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

Andy Kaylor via cfe-commits cfe-commits at lists.llvm.org
Wed May 6 17:04:54 PDT 2026


================
@@ -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();
----------------
andykaylor wrote:

```suggestion
    return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), 
                                       "amdgcn.ds.swizzle", src0.getType(),
                                       {src0, src1});
```

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


More information about the cfe-commits mailing list