[clang] [llvm] [Clang][AMDGPU] Handle `wavefrontsize32` and `wavefrontsize64` features more robustly (PR #176599)

[via cfe-commits]

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