[CIR][AMDGPU] Add lowering for amdgcn ds swizzle builtin

Author: skc7

clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp

@@ -91,7 +91,15 @@ mlir::Value CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   case AMDGPU::BI__builtin_amdgcn_div_fmasf: {
     llvm_unreachable("div_fmas_* NYI");
   }
-  case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
+  case AMDGPU::BI__builtin_amdgcn_ds_swizzle: {
+    mlir::Value arg0 = emitScalarExpr(expr->getArg(0));
+    mlir::Value arg1 = emitScalarExpr(expr->getArg(1));
+    return LLVMIntrinsicCallOp::create(
+               builder, getLoc(expr->getExprLoc()),
+               builder.getStringAttr("amdgcn.ds.swizzle"), arg0.getType(),
+               {arg0, arg1})
+        .getResult();
+  }
   case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
   case AMDGPU::BI__builtin_amdgcn_mov_dpp:
   case AMDGPU::BI__builtin_amdgcn_update_dpp: {

clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip

@@ -218,3 +218,13 @@ __device__ void test_wave_reduce_add_u32_iterative_i32(int* out, int in) {
 __device__ void test_wave_reduce_add_u32_dpp_i32(int* out, int in) {
   *out = __builtin_amdgcn_wave_reduce_add_u32(in, 2);
 }
+
+// CIR-LABEL: @_Z19test_ds_swizzle_i32Pii
+// CIR: cir.llvm.intrinsic "amdgcn.ds.swizzle" {{.*}} : (!s32i, !s32i) -> !s32i
+// LLVM: define{{.*}} void @_Z19test_ds_swizzle_i32Pii
+// LLVM: call i32 @llvm.amdgcn.ds.swizzle(i32 %{{.*}}, i32 32)
+// OGCG: define{{.*}} void @_Z19test_ds_swizzle_i32Pii
+// OGCG: call i32 @llvm.amdgcn.ds.swizzle(i32 %{{.*}}, i32 32)
+__device__ void test_ds_swizzle_i32(int* out, int a) {
+  *out = __builtin_amdgcn_ds_swizzle(a, 32);
+}

clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl

@@ -220,3 +220,13 @@ void test_wave_reduce_add_u32_iterative(global int* out, int in) {
 void test_wave_reduce_add_u32_dpp(global int* out, int in) {
   *out = __builtin_amdgcn_wave_reduce_add_u32(in, 2);
 }
+
+// CIR-LABEL: @test_ds_swizzle
+// CIR: cir.llvm.intrinsic "amdgcn.ds.swizzle" {{.*}} : (!s32i, !s32i) -> !s32i
+// LLVM: define{{.*}} void @test_ds_swizzle
+// LLVM: call i32 @llvm.amdgcn.ds.swizzle(i32 %{{.*}}, i32 32)
+// OGCG: define{{.*}} void @test_ds_swizzle
+// OGCG: call i32 @llvm.amdgcn.ds.swizzle(i32 %{{.*}}, i32 32)
+void test_ds_swizzle(global int* out, int a) {
+  *out = __builtin_amdgcn_ds_swizzle(a, 32);
+}