[clang] 0f24969 - [Clang][AMDGPU] Deprecate icmp/fcmp builtins in favor of ballot (#209416)

[via cfe-commits]

Author: Jay Foad
Date: 2026-07-15T11:07:57Z
New Revision: 0f24969ade9df74b11ff70a2c0b1f49ff49e92fd

URL: https://github.com/llvm/llvm-project/commit/0f24969ade9df74b11ff70a2c0b1f49ff49e92fd
DIFF: https://github.com/llvm/llvm-project/commit/0f24969ade9df74b11ff70a2c0b1f49ff49e92fd.diff

LOG: [Clang][AMDGPU] Deprecate icmp/fcmp builtins in favor of ballot (#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

Added: 
    clang/test/SemaOpenCL/builtins-amdgcn-cmp-deprecated.cl

Modified: 
    clang/docs/ReleaseNotes.md
    clang/lib/Sema/SemaAMDGPU.cpp

Removed: 
    


################################################################################
diff  --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md
index e31871c9c9340..b25edbd2b2e33 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -157,6 +157,14 @@ latest release, please see the [Clang Web Site](https://clang.llvm.org) or the
 
 #### AMDGPU Support
 
+- Deprecated the following builtins 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
 
 #### X86 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 at +2 {{builtin '__builtin_amdgcn_sicmp' is deprecated; use __builtin_amdgcn_ballot_w64 instead}}
+  // wave32-warning at +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 at +2 {{builtin '__builtin_amdgcn_sicmpl' is deprecated; use __builtin_amdgcn_ballot_w64 instead}}
+  // wave32-warning at +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 at +2 {{builtin '__builtin_amdgcn_uicmp' is deprecated; use __builtin_amdgcn_ballot_w64 instead}}
+  // wave32-warning at +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 at +2 {{builtin '__builtin_amdgcn_uicmpl' is deprecated; use __builtin_amdgcn_ballot_w64 instead}}
+  // wave32-warning at +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 at +2 {{builtin '__builtin_amdgcn_fcmp' is deprecated; use __builtin_amdgcn_ballot_w64 instead}}
+  // wave32-warning at +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 at +2 {{builtin '__builtin_amdgcn_fcmpf' is deprecated; use __builtin_amdgcn_ballot_w64 instead}}
+  // wave32-warning at +1 {{builtin '__builtin_amdgcn_fcmpf' is deprecated; use __builtin_amdgcn_ballot_w32 instead}}
+  *out = __builtin_amdgcn_fcmpf(a, b, 1);
+}