llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang @llvm/pr-subscribers-clangir Author: Ayokunle Amodu (ayokunle321) <details> <summary>Changes</summary> 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. Assisted by: Claude Opus 5 --- Full diff: https://github.com/llvm/llvm-project/pull/223161.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+12) - (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip (+36) ``````````diff 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); +} `````````` </details> https://github.com/llvm/llvm-project/pull/223161 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
