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

[via cfe-commits]

Author: Ayokunle Amodu
Date: 2026-09-08T05:34:26-04:00
New Revision: b3071f2b0db631e8fb0d4b47f5e46d0c8297879a

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

LOG: [CIR][AMDGPU] Add support for AMDGCN class builtins (#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.

Added: 
    

Modified: 
    clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
    clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
    clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 0280d38b93bd3..1c23ea142bf8d 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -436,12 +436,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 e298056ab6d9d..08776b043a900 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
@@ -121,3 +121,11 @@ __device__ void test_frexp_mant_f16(_Float16* out, _Float16 a) {
 __device__ void test_ldexp_f16(_Float16* out, _Float16 a, int b) {
     *out = __builtin_amdgcn_ldexph(a, b);
 }
+
+// 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 7dcd343aabb30..9bbfb6747cbf8 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -315,3 +315,19 @@ __device__ void test_ldexp_f32(float* out, float a, int b) {
 __device__ void test_ldexp_f64(double* out, double a, int b) {
   *out = __builtin_amdgcn_ldexp(a, b);
 }
+
+// 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);
+}