https://github.com/jayfoad created https://github.com/llvm/llvm-project/pull/209416
Deprecate these builtins with a warning recommending __builtin_amdgcn_ballot_w32 or __builtin_amdgcn_ballot_w64 instead: __builtin_amdgcn_uicmp __builtin_amdgcn_uicmpl __builtin_amdgcn_sicmp __builtin_amdgcn_sicmpl __builtin_amdgcn_fcmp __builtin_amdgcn_fcmpf Assisted-by: Claude Opus 4.8 >From 615a709c1867bd095f7251cc3d6ab47a1a141ce8 Mon Sep 17 00:00:00 2001 From: Jay Foad <[email protected]> Date: Tue, 14 Jul 2026 09:38:50 +0100 Subject: [PATCH] [Clang][AMDGPU] Deprecate icmp/fcmp builtins in favor of ballot Deprecate these builtins with a warning recommending __builtin_amdgcn_ballot_w32 or __builtin_amdgcn_ballot_w64 instead: __builtin_amdgcn_uicmp __builtin_amdgcn_uicmpl __builtin_amdgcn_sicmp __builtin_amdgcn_sicmpl __builtin_amdgcn_fcmp __builtin_amdgcn_fcmpf Assisted-by: Claude Opus 4.8 --- clang/docs/ReleaseNotes.md | 7 +++ clang/lib/Sema/SemaAMDGPU.cpp | 17 ++++++ .../builtins-amdgcn-cmp-deprecated.cl | 52 +++++++++++++++++++ 3 files changed, 76 insertions(+) create mode 100644 clang/test/SemaOpenCL/builtins-amdgcn-cmp-deprecated.cl diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index 9063e54b3e692..36fbe821b7a08 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -1006,6 +1006,13 @@ latest release, please see the [Clang Web Site](https://clang.llvm.org) or the - The `amdgpu_num_sgpr` and `amdgpu_num_vgpr` function attributes are now deprecated. Using them produces a `-Wdeprecated-declarations` warning. Use `amdgpu_waves_per_eu` instead. +- These builtins have been deprecated in favor of `__builtin_amdgcn_ballot_w32` + or `__builtin_amdgcn_ballot_w64`: + - `__builtin_amdgcn_uicmp` + - `__builtin_amdgcn_uicmpl` + - `__builtin_amdgcn_sicmpl` + - `__builtin_amdgcn_fcmp` + - `__builtin_amdgcn_fcmpf` #### NVPTX Support diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index f1de44e3d2ed7..48230fa262d5c 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -86,6 +86,23 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, return true; } } + case AMDGPU::BI__builtin_amdgcn_uicmp: + case AMDGPU::BI__builtin_amdgcn_uicmpl: + case AMDGPU::BI__builtin_amdgcn_sicmp: + case AMDGPU::BI__builtin_amdgcn_sicmpl: + case AMDGPU::BI__builtin_amdgcn_fcmp: + case AMDGPU::BI__builtin_amdgcn_fcmpf: { + // These builtins are deprecated in favor of + // __builtin_amdgcn_ballot_{w32|w64}. Suggest the replacement matching the + // wavefront size of the calling function. + bool IsWave32 = Builtin::evaluateRequiredTargetFeatures("wavefrontsize32", + CallerFeatureMap); + Diag(TheCall->getBeginLoc(), diag::warn_deprecated_builtin) + << getASTContext().BuiltinInfo.getQuotedName(BuiltinID) + << (IsWave32 ? "__builtin_amdgcn_ballot_w32" + : "__builtin_amdgcn_ballot_w64"); + return false; + } case AMDGPU::BI__builtin_amdgcn_get_fpenv: case AMDGPU::BI__builtin_amdgcn_set_fpenv: return false; diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-cmp-deprecated.cl b/clang/test/SemaOpenCL/builtins-amdgcn-cmp-deprecated.cl new file mode 100644 index 0000000000000..6ceeb4e2c500b --- /dev/null +++ b/clang/test/SemaOpenCL/builtins-amdgcn-cmp-deprecated.cl @@ -0,0 +1,52 @@ +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx900 -verify=expected,wave64 -fsyntax-only %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1010 -target-feature +wavefrontsize64 -verify=expected,wave64 -fsyntax-only %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1010 -target-feature +wavefrontsize32 -verify=expected,wave32 -fsyntax-only %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 -verify=expected,wave32 -fsyntax-only %s + +// Check that -Wno-deprecated-builtins silences the warnings. +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx900 -Wno-deprecated-builtins -verify=silenced -fsyntax-only %s + +// REQUIRES: amdgpu-registered-target + +// silenced-no-diagnostics + +#pragma OPENCL EXTENSION cl_khr_fp64 : enable + +typedef unsigned int uint; +typedef unsigned long ulong; + +void test_sicmp(global ulong* out, int a, int b) { + // wave64-warning@+2 {{builtin '__builtin_amdgcn_sicmp' is deprecated; use __builtin_amdgcn_ballot_w64 instead}} + // wave32-warning@+1 {{builtin '__builtin_amdgcn_sicmp' is deprecated; use __builtin_amdgcn_ballot_w32 instead}} + *out = __builtin_amdgcn_sicmp(a, b, 32); +} + +void test_sicmpl(global ulong* out, long a, long b) { + // wave64-warning@+2 {{builtin '__builtin_amdgcn_sicmpl' is deprecated; use __builtin_amdgcn_ballot_w64 instead}} + // wave32-warning@+1 {{builtin '__builtin_amdgcn_sicmpl' is deprecated; use __builtin_amdgcn_ballot_w32 instead}} + *out = __builtin_amdgcn_sicmpl(a, b, 32); +} + +void test_uicmp(global ulong* out, uint a, uint b) { + // wave64-warning@+2 {{builtin '__builtin_amdgcn_uicmp' is deprecated; use __builtin_amdgcn_ballot_w64 instead}} + // wave32-warning@+1 {{builtin '__builtin_amdgcn_uicmp' is deprecated; use __builtin_amdgcn_ballot_w32 instead}} + *out = __builtin_amdgcn_uicmp(a, b, 32); +} + +void test_uicmpl(global ulong* out, ulong a, ulong b) { + // wave64-warning@+2 {{builtin '__builtin_amdgcn_uicmpl' is deprecated; use __builtin_amdgcn_ballot_w64 instead}} + // wave32-warning@+1 {{builtin '__builtin_amdgcn_uicmpl' is deprecated; use __builtin_amdgcn_ballot_w32 instead}} + *out = __builtin_amdgcn_uicmpl(a, b, 32); +} + +void test_fcmp(global ulong* out, double a, double b) { + // wave64-warning@+2 {{builtin '__builtin_amdgcn_fcmp' is deprecated; use __builtin_amdgcn_ballot_w64 instead}} + // wave32-warning@+1 {{builtin '__builtin_amdgcn_fcmp' is deprecated; use __builtin_amdgcn_ballot_w32 instead}} + *out = __builtin_amdgcn_fcmp(a, b, 1); +} + +void test_fcmpf(global ulong* out, float a, float b) { + // wave64-warning@+2 {{builtin '__builtin_amdgcn_fcmpf' is deprecated; use __builtin_amdgcn_ballot_w64 instead}} + // wave32-warning@+1 {{builtin '__builtin_amdgcn_fcmpf' is deprecated; use __builtin_amdgcn_ballot_w32 instead}} + *out = __builtin_amdgcn_fcmpf(a, b, 1); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
