================
@@ -824,17 +824,32 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:
case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:
case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
- case AMDGPU::BI__builtin_amdgcn_load_to_lds: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ llvm::StringRef intrinsicName;
+ switch (builtinId) {
+ case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:
+ intrinsicName = "amdgcn.cluster.load.b32";
+ break;
+ case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:
+ intrinsicName = "amdgcn.cluster.load.b64";
+ break;
+ case AMDGPU::BI__builtin_amdgcn_cluster_load_b128:
+ intrinsicName = "amdgcn.cluster.load.b128";
+ break;
+ }
+ return emitBuiltinWithOneOverloadedType<3>(expr, intrinsicName,
+ convertType(expr->getType()))
+ .getValue();
}
----------------
aobolensk wrote:
Not at all, that became even better, diff is +8 -15, thanks for the suggestion!
https://github.com/llvm/llvm-project/pull/233756
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
