[clang] [CIR][AMDGPU] Add support for AMDGCN permlane builtins (PR #223161)

[via cfe-commits]

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