https://github.com/shiltian created https://github.com/llvm/llvm-project/pull/176599
We should also not allow `-wavefrontsize32` and `-wavefrontsize64` to be specified at the same time. >From 727f8f9f05aa1f8d4bdfd397e63778a166cb8c01 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> 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 | 13 ++++++---- .../builtins-amdgcn-error-wave32.cl | 2 +- llvm/lib/TargetParser/TargetParser.cpp | 25 +++++++++++++------ 3 files changed, 27 insertions(+), 13 deletions(-) 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)); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
