llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang @llvm/pr-subscribers-backend-amdgpu Author: Chaitanya (skc7) <details> <summary>Changes</summary> Upstreaming clangIR PR: https://github.com/llvm/clangir/pull/2053 This PR adds support for lowering of "_builtin_amdgcn_readlane" and "_builtin_amdgcn_readfirstlane" amdgpu builtins to clangIR. --- Full diff: https://github.com/llvm/llvm-project/pull/197153.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+13-2) - (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip (+16) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 04ab1c29b0d63..7e6e0f1a06046 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -184,8 +184,19 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, getContext().BuiltinInfo.getName(builtinId)); return mlir::Value{}; } - case AMDGPU::BI__builtin_amdgcn_readlane: - case AMDGPU::BI__builtin_amdgcn_readfirstlane: + case AMDGPU::BI__builtin_amdgcn_readlane: { + mlir::Value src0 = emitScalarExpr(expr->getArg(0)); + mlir::Value src1 = emitScalarExpr(expr->getArg(1)); + return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), + "amdgcn.readlane", src0.getType(), + mlir::ValueRange{src0, src1}); + } + case AMDGPU::BI__builtin_amdgcn_readfirstlane: { + mlir::Value src0 = emitScalarExpr(expr->getArg(0)); + return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), + "amdgcn.readfirstlane", src0.getType(), + mlir::ValueRange{src0}); + } case AMDGPU::BI__builtin_amdgcn_wave_shuffle: { cgm.errorNYI(expr->getSourceRange(), std::string("unimplemented AMDGPU builtin call: ") + diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 4a61fde7aa90c..b3c5ac59679a3 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -71,3 +71,19 @@ __device__ void test_div_fmas_f64(double* out, double a, double b, double c, int __device__ void test_ds_swizzle(int* out, int a) { *out = __builtin_amdgcn_ds_swizzle(a, 32); } + +// CIR-LABEL: @_Z13test_readlanePiii +// CIR: cir.call_llvm_intrinsic "amdgcn.readlane" {{.*}} : (!s32i, !s32i) -> !s32i +// LLVM: define{{.*}} void @_Z13test_readlanePiii +// LLVM: call{{.*}} i32 @llvm.amdgcn.readlane.i32(i32 %{{.*}}, i32 %{{.*}}) +__device__ void test_readlane(int* out, int a, int b) { + *out = __builtin_amdgcn_readlane(a, b); +} + +// CIR-LABEL: @_Z18test_readfirstlanePii +// CIR: cir.call_llvm_intrinsic "amdgcn.readfirstlane" {{.*}} : (!s32i) -> !s32i +// LLVM: define{{.*}} void @_Z18test_readfirstlanePii +// LLVM: call{{.*}} i32 @llvm.amdgcn.readfirstlane.i32(i32 %{{.*}}) +__device__ void test_readfirstlane(int* out, int a) { + *out = __builtin_amdgcn_readfirstlane(a); +} `````````` </details> https://github.com/llvm/llvm-project/pull/197153 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
