================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
