Author: Ayokunle Amodu Date: 2026-09-09T14:56:29-04:00 New Revision: 2c9afb46c32ccf19fe048688e273cad4c7b3dfc6
URL: https://github.com/llvm/llvm-project/commit/2c9afb46c32ccf19fe048688e273cad4c7b3dfc6 DIFF: https://github.com/llvm/llvm-project/commit/2c9afb46c32ccf19fe048688e273cad4c7b3dfc6.diff LOG: [CIR][AMDGPU] Add support for AMDGCN bitfield extract builtins (#222186) Adds codegen for the following AMDGCN bitfield extract builtins: * __builtin_amdgcn_ubfe (unsigned) * __builtin_amdgcn_sbfe (signed) These are lowered to the corresponding `llvm.amdgcn.ubfe` and `llvm.amdgcn.sbfe` intrinsics. Added: Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp 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 94d5a835d60fa..c31ac8fa946c9 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -346,18 +346,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, getContext().BuiltinInfo.getName(builtinId)); return mlir::Value{}; } - case AMDGPU::BI__builtin_amdgcn_ubfe: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } - case AMDGPU::BI__builtin_amdgcn_sbfe: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case AMDGPU::BI__builtin_amdgcn_ubfe: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.ubfe").getValue(); + case AMDGPU::BI__builtin_amdgcn_sbfe: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.sbfe").getValue(); case AMDGPU::BI__builtin_amdgcn_ballot_w32: case AMDGPU::BI__builtin_amdgcn_ballot_w64: return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ballot", diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 9bbfb6747cbf8..9b687e98d1195 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -331,3 +331,21 @@ __device__ void test_class_f32(bool* out, float a, int b) { __device__ void test_class_f64(bool* out, double a, int b) { *out = __builtin_amdgcn_class(a, b); } + +// CIR-LABEL: @_Z9test_ubfePjjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.ubfe" {{.*}} : (!u32i, !u32i, !u32i) -> !u32i +// LLVM: define{{.*}} void @_Z9test_ubfePjjjj +// LLVM: call{{.*}} i32 @llvm.amdgcn.ubfe.i32(i32 %{{.+}}, i32 %{{.+}}, i32 %{{.+}}) +__device__ void test_ubfe(unsigned int* out, unsigned int src0, unsigned int src1, + unsigned int src2) { + *out = __builtin_amdgcn_ubfe(src0, src1, src2); +} + +// CIR-LABEL: @_Z9test_sbfePjjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.sbfe" {{.*}} : (!u32i, !u32i, !u32i) -> !u32i +// LLVM: define{{.*}} void @_Z9test_sbfePjjjj +// LLVM: call{{.*}} i32 @llvm.amdgcn.sbfe.i32(i32 %{{.+}}, i32 %{{.+}}, i32 %{{.+}}) +__device__ void test_sbfe(unsigned int* out, unsigned int src0, unsigned int src1, + unsigned int src2) { + *out = __builtin_amdgcn_sbfe(src0, src1, src2); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
