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 <[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 | 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)); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
