https://github.com/ayokunle321 created https://github.com/llvm/llvm-project/pull/223161
Adds codegen for the following AMDGCN permlane builtins: - __builtin_amdgcn_permlane_bcast - __builtin_amdgcn_permlane_up - __builtin_amdgcn_permlane_down - __builtin_amdgcn_permlane_xor These are lowered to the corresponding `llvm.amdgcn.permlane.*` intrinsics. >From 0084b73f160f33034249d337fa2e9f87814077bf Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 12 Sep 2026 19:04:58 +0200 Subject: [PATCH] add permlance builtins --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 12 +++++++ .../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 36 +++++++++++++++++++ 2 files changed, 48 insertions(+) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 1c23ea142bf8d..00920cdebeec5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -225,6 +225,18 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_permlane64: return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.permlane64") .getValue(); + case AMDGPU::BI__builtin_amdgcn_permlane_bcast: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.bcast") + .getValue(); + case AMDGPU::BI__builtin_amdgcn_permlane_up: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.up") + .getValue(); + case AMDGPU::BI__builtin_amdgcn_permlane_down: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.down") + .getValue(); + case AMDGPU::BI__builtin_amdgcn_permlane_xor: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.xor") + .getValue(); case AMDGPU::BI__builtin_amdgcn_readlane: return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane") .getValue(); diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip index b6e46ba180c34..72ec665a782b3 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip @@ -96,3 +96,39 @@ __device__ void test_tanh_f16(_Float16* out, _Float16 a) { __device__ void test_tanh_bf16(__bf16* out, __bf16 a) { *out = __builtin_amdgcn_tanh_bf16(a); } + +// CIR-LABEL: @_Z19test_permlane_bcastPjjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.permlane.bcast" {{.*}} : (!s32i, !s32i, !s32i) -> !s32i +// LLVM: define{{.*}} void @_Z19test_permlane_bcastPjjjj +// LLVM: call{{.*}} i32 @llvm.amdgcn.permlane.bcast.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}) +__device__ void test_permlane_bcast(unsigned int* out, unsigned int src0, + unsigned int src1, unsigned int src2) { + *out = __builtin_amdgcn_permlane_bcast(src0, src1, src2); +} + +// CIR-LABEL: @_Z16test_permlane_upPjjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.permlane.up" {{.*}} : (!s32i, !s32i, !s32i) -> !s32i +// LLVM: define{{.*}} void @_Z16test_permlane_upPjjjj +// LLVM: call{{.*}} i32 @llvm.amdgcn.permlane.up.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}) +__device__ void test_permlane_up(unsigned int* out, unsigned int src0, + unsigned int src1, unsigned int src2) { + *out = __builtin_amdgcn_permlane_up(src0, src1, src2); +} + +// CIR-LABEL: @_Z18test_permlane_downPjjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.permlane.down" {{.*}} : (!s32i, !s32i, !s32i) -> !s32i +// LLVM: define{{.*}} void @_Z18test_permlane_downPjjjj +// LLVM: call{{.*}} i32 @llvm.amdgcn.permlane.down.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}) +__device__ void test_permlane_down(unsigned int* out, unsigned int src0, + unsigned int src1, unsigned int src2) { + *out = __builtin_amdgcn_permlane_down(src0, src1, src2); +} + +// CIR-LABEL: @_Z17test_permlane_xorPjjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.permlane.xor" {{.*}} : (!s32i, !s32i, !s32i) -> !s32i +// LLVM: define{{.*}} void @_Z17test_permlane_xorPjjjj +// LLVM: call{{.*}} i32 @llvm.amdgcn.permlane.xor.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}) +__device__ void test_permlane_xor(unsigned int* out, unsigned int src0, + unsigned int src1, unsigned int src2) { + *out = __builtin_amdgcn_permlane_xor(src0, src1, src2); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
