Author: Ayokunle Amodu Date: 2026-09-09T14:56:54-04:00 New Revision: 3cab00bd5674d374e702807a57a3db432a351086
URL: https://github.com/llvm/llvm-project/commit/3cab00bd5674d374e702807a57a3db432a351086 DIFF: https://github.com/llvm/llvm-project/commit/3cab00bd5674d374e702807a57a3db432a351086.diff LOG: [CIR][AMDGPU] Add support for AMDGCN fract builtins (#222185) Adds codegen for the following AMDGCN fract builtins: * __builtin_amdgcn_fract (double) * __builtin_amdgcn_fractf (float) * __builtin_amdgcn_fracth (half) These are lowered to the corresponding `llvm.amdgcn.fract` 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 c31ac8fa946c9..3494d1363d70f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -334,12 +334,8 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, } case AMDGPU::BI__builtin_amdgcn_fract: case AMDGPU::BI__builtin_amdgcn_fractf: - case AMDGPU::BI__builtin_amdgcn_fracth: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case AMDGPU::BI__builtin_amdgcn_fracth: + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.fract").getValue(); case AMDGPU::BI__builtin_amdgcn_lerp: { cgm.errorNYI(expr->getSourceRange(), std::string("unimplemented AMDGPU builtin call: ") + diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip index 08776b043a900..afc4c74572489 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip @@ -122,6 +122,14 @@ __device__ void test_ldexp_f16(_Float16* out, _Float16 a, int b) { *out = __builtin_amdgcn_ldexph(a, b); } +// CIR-LABEL: @_Z14test_fract_f16PDF16_DF16_ +// CIR: cir.call_llvm_intrinsic "amdgcn.fract" {{.*}} : (!cir.f16) -> !cir.f16 +// LLVM: define{{.*}} void @_Z14test_fract_f16PDF16_DF16_ +// LLVM: call{{.*}} half @llvm.amdgcn.fract.f16(half %{{.*}}) +__device__ void test_fract_f16(_Float16* out, _Float16 a) { + *out = __builtin_amdgcn_fracth(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 diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 9b687e98d1195..713bc10700cdd 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -224,6 +224,22 @@ __device__ void test_frexp_mant_f64(double* out, double a) { *out = __builtin_amdgcn_frexp_mant(a); } +// CIR-LABEL: @_Z14test_fract_f32Pff +// CIR: cir.call_llvm_intrinsic "amdgcn.fract" {{.*}} : (!cir.float) -> !cir.float +// LLVM: define{{.*}} void @_Z14test_fract_f32Pff +// LLVM: call{{.*}} float @llvm.amdgcn.fract.f32(float %{{.+}}) +__device__ void test_fract_f32(float* out, float a) { + *out = __builtin_amdgcn_fractf(a); +} + +// CIR-LABEL: @_Z14test_fract_f64Pdd +// CIR: cir.call_llvm_intrinsic "amdgcn.fract" {{.*}} : (!cir.double) -> !cir.double +// LLVM: define{{.*}} void @_Z14test_fract_f64Pdd +// LLVM: call{{.*}} double @llvm.amdgcn.fract.f64(double %{{.+}}) +__device__ void test_fract_f64(double* out, double a) { + *out = __builtin_amdgcn_fract(a); +} + // CIR-LABEL: @_Z14test_sicmp_i32Pmii // CIR: [[CMP:%.+]] = cir.cmp eq %{{.+}}, %{{.+}} : !s32i // CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[CMP]] : (!cir.bool) -> !u64i _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
