================
@@ -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();
}
----------------
steffenlarsen wrote:
Even if it is a little repetitive, it would probably be about as short an
implementation to just call `emitBuiltinWithOneOverloadedType` for each case
and inline the intrinsic name.
https://github.com/llvm/llvm-project/pull/233756
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
