[clang] [CIR][AMDGPU] Add support for AMDGCN class builtins (PR #213496)

[Ayokunle Amodu via cfe-commits]

https://github.com/ayokunle321 created https://github.com/llvm/llvm-project/pull/213496

Adds codegen for the following AMDGCN class builtins:

- __builtin_amdgcn_class (double)
- __builtin_amdgcn_classf (float)
- __builtin_amdgcn_classh (half)

These are lowered to the corresponding `llvm.amdgcn.class` intrinsics.

>From 207237679573c15febde956dff7c9a481a17e56f Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sun, 2 Aug 2026 00:56:23 +0000
Subject: [PATCH] add amdgcn class builtins

---
 clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp    | 10 ++++------
 .../CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip    |  8 ++++++++
 clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip    | 16 ++++++++++++++++
 3 files changed, 28 insertions(+), 6 deletions(-)

diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index b9ae5948ba013..be4075fa9f93b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -384,12 +384,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   }
   case AMDGPU::BI__builtin_amdgcn_class:
   case AMDGPU::BI__builtin_amdgcn_classf:
-  case AMDGPU::BI__builtin_amdgcn_classh: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
-  }
+  case AMDGPU::BI__builtin_amdgcn_classh:
+    return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.class",
+                                               convertType(expr->getType()))
+        .getValue();
   case AMDGPU::BI__builtin_amdgcn_fmed3f:
   case AMDGPU::BI__builtin_amdgcn_fmed3h: {
     cgm.errorNYI(expr->getSourceRange(),
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
index 11a97c9befa9d..113bc15e5e291 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
@@ -111,3 +111,11 @@ __device__ void test_cos_f16(_Float16* out, _Float16 a) {
 __device__ void test_frexp_mant_f16(_Float16* out, _Float16 a) {
   *out = __builtin_amdgcn_frexp_manth(a);
 }
+
+// CIR-LABEL: @_Z14test_class_f16PbDF16_i
+// CIR: cir.call_llvm_intrinsic "amdgcn.class" {{.*}} : (!cir.f16, !s32i) -> !cir.bool
+// LLVM: define{{.*}} void @_Z14test_class_f16PbDF16_i
+// LLVM: call{{.*}} i1 @llvm.amdgcn.class.f16(half %{{.*}}, i32 %{{.*}})
+__device__ void test_class_f16(bool* out, _Float16 a, int b) {
+  *out = __builtin_amdgcn_classh(a, b);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index 4a129a7e534ac..f590727cb38d1 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -223,3 +223,19 @@ __device__ void test_frexp_mant_f32(float* out, float a) {
 __device__ void test_frexp_mant_f64(double* out, double a) {
   *out = __builtin_amdgcn_frexp_mant(a);
 }
+
+// CIR-LABEL: @_Z14test_class_f32Pbfi
+// CIR: cir.call_llvm_intrinsic "amdgcn.class" {{.*}} : (!cir.float, !s32i) -> !cir.bool
+// LLVM: define{{.*}} void @_Z14test_class_f32Pbfi
+// LLVM: call{{.*}} i1 @llvm.amdgcn.class.f32(float %{{.*}}, i32 %{{.*}})
+__device__ void test_class_f32(bool* out, float a, int b) {
+  *out = __builtin_amdgcn_classf(a, b);
+}
+
+// CIR-LABEL: @_Z14test_class_f64Pbdi
+// CIR: cir.call_llvm_intrinsic "amdgcn.class" {{.*}} : (!cir.double, !s32i) -> !cir.bool
+// LLVM: define{{.*}} void @_Z14test_class_f64Pbdi
+// LLVM: call{{.*}} i1 @llvm.amdgcn.class.f64(double %{{.*}}, i32 %{{.*}})
+__device__ void test_class_f64(bool* out, double a, int b) {
+  *out = __builtin_amdgcn_class(a, b);
+}