https://github.com/skc7 created https://github.com/llvm/llvm-project/pull/197153
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. >From 9bb1933385708f02c2a2fb3c92c8cec38f02e979 Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Tue, 12 May 2026 16:40:14 +0530 Subject: [PATCH] [CIR][AMDGPU] Add lowering for amdgcn readlane readfirstlane builtins --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 15 +++++++++++++-- clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 16 ++++++++++++++++ 2 files changed, 29 insertions(+), 2 deletions(-) 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); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
