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); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
