[clang] [flang] [llvm] [Clang][AMDGPU] Handle `wavefrontsize32` and `wavefrontsize64` features more robustly (PR #176599)
Shilei Tian via cfe-commits
cfe-commits at lists.llvm.org
Sat Jan 17 15:18:31 PST 2026
https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/176599
>From 032eb06ad6caa94a68230cd5876a69bd657ddb13 Mon Sep 17 00:00:00 2001
From: Shilei Tian <i at tianshilei.me>
Date: Sat, 17 Jan 2026 17:30:07 -0500
Subject: [PATCH] [Clang][AMDGPU] Handle `wavefrontsize32` and
`wavefrontsize64` features more robustly
We should also not allow `-wavefrontsize32` and `-wavefrontsize64` to be specified at the same time.
---
.../CodeGenOpenCL/amdgpu-features-illegal.cl | 21 ++++++----
.../builtins-amdgcn-error-wave32.cl | 4 +-
.../builtins-amdgcn-wave32-func-attr.cl | 4 +-
.../Driver/target-cpu-features-invalid.f90 | 8 +++-
llvm/lib/TargetParser/TargetParser.cpp | 39 ++++++++++++++-----
5 files changed, 54 insertions(+), 22 deletions(-)
diff --git a/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl b/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl
index 04de5dca3f6c0..f3857178d928e 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-features-illegal.cl
@@ -1,10 +1,17 @@
-// 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-WAVE32
+// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx900 -target-feature -wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck %s --check-prefix=GFX9-WAVE64
+// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx1250 -target-feature +wavefrontsize64 -o /dev/null %s 2>&1 | FileCheck %s --check-prefix=GFX1250-WAVE64
+// RUN: not %clang_cc1 -triple amdgcn -target-cpu gfx1250 -target-feature -wavefrontsize32 -o /dev/null %s 2>&1 | FileCheck %s --check-prefix=GFX1250-WAVE32
-// CHECK: 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
+// 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-WAVE32: error: option '+wavefrontsize32' cannot be specified on this target
+// GFX9-WAVE64: error: option '-wavefrontsize64' cannot be specified on this target
+// GFX1250-WAVE64: error: option '+wavefrontsize64' cannot be specified on this target
+// GFX1250-WAVE32: error: option '-wavefrontsize32' cannot be specified on this target
kernel void test() {}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl
index 949ff07f76154..34c4ae42391d3 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-wave32.cl
@@ -12,7 +12,7 @@ void test_ballot_wave32(global uint* out, int a, int b) {
*out = __builtin_amdgcn_ballot_w32(a == b); // expected-error {{'__builtin_amdgcn_ballot_w32' needs target feature wavefrontsize32}}
}
-__attribute__((target("wavefrontsize32"))) // gfx9-error@*:* {{option 'wavefrontsize32' cannot be specified on this target}}
+__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/clang/test/SemaOpenCL/builtins-amdgcn-wave32-func-attr.cl b/clang/test/SemaOpenCL/builtins-amdgcn-wave32-func-attr.cl
index 55c890368525a..6ef872e6add3a 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-wave32-func-attr.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-wave32-func-attr.cl
@@ -11,8 +11,8 @@
typedef unsigned int uint;
-// GFX9: error: option 'wavefrontsize32' cannot be specified on this target
-__attribute__((target("wavefrontsize32"))) // gfx9-error@*:* {{option 'wavefrontsize32' cannot be specified on this target}}
+// GFX9: error: option '+wavefrontsize32' cannot be specified on this target
+__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);
}
diff --git a/flang/test/Driver/target-cpu-features-invalid.f90 b/flang/test/Driver/target-cpu-features-invalid.f90
index 288da8d57e81d..7e72bf1684c55 100644
--- a/flang/test/Driver/target-cpu-features-invalid.f90
+++ b/flang/test/Driver/target-cpu-features-invalid.f90
@@ -9,8 +9,12 @@
! RUN: -o /dev/null -S %s 2>&1 | FileCheck %s -check-prefix=CHECK-INVALID-FEATURE
! RUN: not %flang_fc1 -triple amdgcn-amd-amdhsa -target-feature +wavefrontsize32 \
-! RUN: -target-feature +wavefrontsize64 -o /dev/null -S %s 2>&1 | FileCheck %s -check-prefix=CHECK-INVALID-WAVEFRONT
+! RUN: -target-feature +wavefrontsize64 -o /dev/null -S %s 2>&1 | FileCheck %s -check-prefix=CHECK-INVALID-WAVEFRONT-PLUS
+
+! RUN: not %flang_fc1 -triple amdgcn-amd-amdhsa -target-feature -wavefrontsize32 \
+! RUN: -target-feature -wavefrontsize64 -o /dev/null -S %s 2>&1 | FileCheck %s -check-prefix=CHECK-INVALID-WAVEFRONT-MINUS
! CHECK-INVALID-CPU: 'supercpu' is not a recognized processor for this target (ignoring processor)
! CHECK-INVALID-FEATURE: '+superspeed' is not a recognized feature for this target (ignoring feature)
-! CHECK-INVALID-WAVEFRONT: 'wavefrontsize32' and 'wavefrontsize64' are mutually exclusive
+! CHECK-INVALID-WAVEFRONT-PLUS: '+wavefrontsize32' and '+wavefrontsize64' are mutually exclusive
+! CHECK-INVALID-WAVEFRONT-MINUS: '-wavefrontsize32' and '-wavefrontsize64' are mutually exclusive
diff --git a/llvm/lib/TargetParser/TargetParser.cpp b/llvm/lib/TargetParser/TargetParser.cpp
index b34295b6a1cb1..ee5e46b2ff8a9 100644
--- a/llvm/lib/TargetParser/TargetParser.cpp
+++ b/llvm/lib/TargetParser/TargetParser.cpp
@@ -373,21 +373,42 @@ 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"};
+ if (DisableWave32 && DisableWave64)
return {AMDGPU::INVALID_FEATURE_COMBINATION,
- "'wavefrontsize32' and 'wavefrontsize64' are mutually exclusive"};
+ "'-wavefrontsize32' and '-wavefrontsize64' are mutually exclusive"};
- if (HaveWave32 && !IsNullGPU && TargetHasWave64)
- return {AMDGPU::UNSUPPORTED_TARGET_FEATURE, "wavefrontsize32"};
+ if (!IsNullGPU && TargetHasWave64) {
+ if (EnableWave32)
+ return {AMDGPU::UNSUPPORTED_TARGET_FEATURE, "+wavefrontsize32"};
+ if (DisableWave64)
+ return {AMDGPU::UNSUPPORTED_TARGET_FEATURE, "-wavefrontsize64"};
+ }
- if (HaveWave64 && !IsNullGPU && TargetHasWave32)
- return {AMDGPU::UNSUPPORTED_TARGET_FEATURE, "wavefrontsize64"};
+ if (!IsNullGPU && TargetHasWave32) {
+ if (EnableWave64)
+ return {AMDGPU::UNSUPPORTED_TARGET_FEATURE, "+wavefrontsize64"};
+ if (DisableWave32)
+ return {AMDGPU::UNSUPPORTED_TARGET_FEATURE, "-wavefrontsize32"};
+ }
// 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));
More information about the cfe-commits
mailing list