================
@@ -206,10 +206,74 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
case AMDGPU::BI__builtin_amdgcn_mov_dpp:
case AMDGPU::BI__builtin_amdgcn_update_dpp: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ mlir::Location loc = getLoc(expr->getExprLoc());
+ unsigned iceArguments = 0;
+ ASTContext::GetBuiltinTypeError error;
+ getContext().GetBuiltinType(builtinId, error, &iceArguments);
+ assert(error == ASTContext::GE_None && "Should not codegen an error");
+ assert(expr->getNumArgs() == 5 || expr->getNumArgs() == 6 ||
+ expr->getNumArgs() == 2);
+
+ mlir::Type dataTy = convertType(expr->getArg(0)->getType());
+ unsigned size = cgm.getDataLayout().getTypeSizeInBits(dataTy);
+ cir::IntType intTy = builder.getUIntNTy(std::max(size, 32u));
+
+ bool isMovDpp8 = builtinId == AMDGPU::BI__builtin_amdgcn_mov_dpp8;
+ bool isMovDpp = builtinId == AMDGPU::BI__builtin_amdgcn_mov_dpp;
+ bool isUpdateDpp = builtinId == AMDGPU::BI__builtin_amdgcn_update_dpp;
+ llvm::StringRef intrinsicName =
+ isMovDpp8 ? "amdgcn.mov.dpp8" : "amdgcn.update.dpp";
+
+ // Fixed parameter types of the target LLVM intrinsics, following the
+ // "old"/"data" operands which share the overloaded integer type.
+ llvm::SmallVector<mlir::Type, 4> fixedTailTypes;
+ mlir::Type ui32 = builder.getUInt32Ty();
+ if (isMovDpp8)
+ fixedTailTypes = {ui32};
+ else
+ fixedTailTypes = {ui32, ui32, ui32, builder.getUIntNTy(1)};
----------------
steffenlarsen wrote:
Agreed. I have a patch that will make it able to find the builtin signature,
but I think it is more suited as a follow-up.
https://github.com/llvm/llvm-project/pull/226469
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
