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

[Jay Foad via cfe-commits]

https://github.com/jayfoad updated https://github.com/llvm/llvm-project/pull/209416

>From 615a709c1867bd095f7251cc3d6ab47a1a141ce8 Mon Sep 17 00:00:00 2001
From: Jay Foad <jay.foad at amd.com>
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 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);
+}