[clang] [CIR][AMDGPU] Add lowering for amdgcn ds swizzle builtin. (PR #196011)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clangir

Author: Chaitanya (skc7)

<details>
<summary>Changes</summary>

Upstreaming clangIR PR: https://github.com/llvm/clangir/pull/2052

This PR adds support for lowering of _builtin_amdgcn_ds_swizzle* amdgpu builtin to clangIR.


---
Full diff: https://github.com/llvm/llvm-project/pull/196011.diff


2 Files Affected:

- (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+10-1) 
- (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip (+8) 


``````````diff
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 929cdf8e88789..7e722d4e654d3 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -161,7 +161,16 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
                              .getResult();
     return result;
   }
-  case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
+  case AMDGPU::BI__builtin_amdgcn_ds_swizzle: {
+    mlir::Value src0 = emitScalarExpr(expr->getArg(0));
+    mlir::Value src1 = emitScalarExpr(expr->getArg(1));
+    mlir::Value result = cir::LLVMIntrinsicCallOp::create(
+                             builder, getLoc(expr->getExprLoc()),
+                             builder.getStringAttr("amdgcn.ds.swizzle"),
+                             src0.getType(), {src0, src1})
+                             .getResult();
+    return result;
+  }
   case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
   case AMDGPU::BI__builtin_amdgcn_mov_dpp:
   case AMDGPU::BI__builtin_amdgcn_update_dpp: {
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index d374479e6182e..4a61fde7aa90c 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -63,3 +63,11 @@ __device__ void test_div_fmas_f32(double* out, float a, float b, float c, int d)
 __device__ void test_div_fmas_f64(double* out, double a, double b, double c, int d) {
   *out = __builtin_amdgcn_div_fmas(a, b, c, d);
 }
+
+// CIR-LABEL: @_Z15test_ds_swizzlePii
+// CIR: cir.call_llvm_intrinsic "amdgcn.ds.swizzle" {{.*}} : (!s32i, !s32i) -> !s32i
+// LLVM: define{{.*}} void @_Z15test_ds_swizzlePii
+// LLVM: call{{.*}} i32 @llvm.amdgcn.ds.swizzle(i32 %{{.*}}, i32 32)
+__device__ void test_ds_swizzle(int* out, int a) {
+  *out = __builtin_amdgcn_ds_swizzle(a, 32);
+}

``````````

</details>


https://github.com/llvm/llvm-project/pull/196011