llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Shilei Tian (shiltian) <details> <summary>Changes</summary> We should also not allow `-wavefrontsize32` and `-wavefrontsize64` to be specified at the same time. --- Full diff: https://github.com/llvm/llvm-project/pull/176599.diff 3 Files Affected: - (modified) clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl (+8-5) - (modified) clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl (+1-1) - (modified) llvm/lib/TargetParser/TargetParser.cpp (+18-7) ``````````diff diff --git a/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl b/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl index 04de5dca3f6c0..d49d5f54fd6fc 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl @@ -1,9 +1,12 @@ -// RUN: not %clang_cc1 -triple amdgcn -target-feature +wavefrontsize32 -target-feature +wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck %s -// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx1103 -target-feature +wavefrontsize32 -target-feature +wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck %s -// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx900 -target-feature +wavefrontsize32 -o /dev/null %s 2>&1 | FileCheck %s --check-prefix=GFX9 -// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx1250 -target-feature +wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck %s --check-prefix=GFX1250 +// RUN: not %clang_cc1 -triple amdgcn -target-feature +wavefrontsize32 -target-feature +wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck --check-prefix=INVALID-PLUS %s +// RUN: not %clang_cc1 -triple amdgcn -target-feature -wavefrontsize32 -target-feature -wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck --check-prefix=INVALID-MINUS %s +// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx1103 -target-feature +wavefrontsize32 -target-feature +wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck --check-prefix=INVALID-PLUS %s +// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx1103 -target-feature -wavefrontsize32 -target-feature -wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck --check-prefix=INVALID-MINUS %s +// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx900 -target-feature +wavefrontsize32 -target-feature -wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck %s --check-prefix=GFX9 +// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx1250 -target-feature -wavefrontsize32 -target-feature +wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck %s --check-prefix=GFX1250 -// CHECK: error: invalid feature combination: 'wavefrontsize32' and 'wavefrontsize64' are mutually exclusive +// INVALID-PLUS: error: invalid feature combination: '+wavefrontsize32' and '+wavefrontsize64' are mutually exclusive +// INVALID-MINUS: error: invalid feature combination: '-wavefrontsize32' and '-wavefrontsize64' are mutually exclusive // GFX9: error: option 'wavefrontsize32' cannot be specified on this target // GFX1250: error: option 'wavefrontsize64' cannot be specified on this target diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl index 949ff07f76154..aea75871eba45 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl @@ -14,5 +14,5 @@ void test_ballot_wave32(global uint* out, int a, int b) { __attribute__((target("wavefrontsize32"))) // gfx9-error@*:* {{option 'wavefrontsize32' cannot be specified on this target}} void test_ballot_wave32_target_attr(global uint* out, int a, int b) { - *out = __builtin_amdgcn_ballot_w32(a == b); // wavefront64-error@*:* {{invalid feature combination: 'wavefrontsize32' and 'wavefrontsize64' are mutually exclusive}} + *out = __builtin_amdgcn_ballot_w32(a == b); // wavefront64-error@*:* {{invalid feature combination: '+wavefrontsize32' and '+wavefrontsize64' are mutually exclusive}} } diff --git a/llvm/lib/TargetParser/TargetParser.cpp b/llvm/lib/TargetParser/TargetParser.cpp index b34295b6a1cb1..c13b0689140c1 100644 --- a/llvm/lib/TargetParser/TargetParser.cpp +++ b/llvm/lib/TargetParser/TargetParser.cpp @@ -373,21 +373,32 @@ insertWaveSizeFeature(StringRef GPU, const Triple &T, const bool IsNullGPU = GPU.empty(); const bool TargetHasWave32 = DefaultFeatures.count("wavefrontsize32"); const bool TargetHasWave64 = DefaultFeatures.count("wavefrontsize64"); - const bool HaveWave32 = Features.count("wavefrontsize32"); - const bool HaveWave64 = Features.count("wavefrontsize64"); - if (HaveWave32 && HaveWave64) + auto Wave32Itr = Features.find("wavefrontsize32"); + auto Wave64Itr = Features.find("wavefrontsize64"); + const bool EnableWave32 = + Wave32Itr != Features.end() && Wave32Itr->getValue(); + const bool EnableWave64 = + Wave64Itr != Features.end() && Wave64Itr->getValue(); + const bool DisableWave32 = + Wave32Itr != Features.end() && !Wave32Itr->getValue(); + const bool DisableWave64 = + Wave64Itr != Features.end() && !Wave64Itr->getValue(); + if (EnableWave32 && EnableWave64) return {AMDGPU::INVALID_FEATURE_COMBINATION, - "'wavefrontsize32' and 'wavefrontsize64' are mutually exclusive"}; + "'+wavefrontsize32' and '+wavefrontsize64' are mutually exclusive"}; + if (DisableWave32 && DisableWave64) + return {AMDGPU::INVALID_FEATURE_COMBINATION, + "'-wavefrontsize32' and '-wavefrontsize64' are mutually exclusive"}; - if (HaveWave32 && !IsNullGPU && TargetHasWave64) + if (EnableWave32 && !IsNullGPU && TargetHasWave64) return {AMDGPU::UNSUPPORTED_TARGET_FEATURE, "wavefrontsize32"}; - if (HaveWave64 && !IsNullGPU && TargetHasWave32) + if (EnableWave64 && !IsNullGPU && TargetHasWave32) return {AMDGPU::UNSUPPORTED_TARGET_FEATURE, "wavefrontsize64"}; // Don't assume any wavesize with an unknown subtarget. // Default to wave32 if target supports both. - if (!IsNullGPU && !HaveWave32 && !HaveWave64 && !TargetHasWave32 && + if (!IsNullGPU && !EnableWave32 && !EnableWave64 && !TargetHasWave32 && !TargetHasWave64) Features.insert(std::make_pair("wavefrontsize32", true)); `````````` </details> https://github.com/llvm/llvm-project/pull/176599 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
