[clang] [llvm] [AMDGPU][Clang] Add check of size for __builtin_amdgcn_global_load_lds (PR #93064)
Matt Arsenault via llvm-commits
llvm-commits at lists.llvm.org
Wed May 22 10:10:18 PDT 2024
================
@@ -19040,6 +19040,48 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
CGM.getIntrinsic(Intrinsic::amdgcn_s_sendmsg_rtn, {ResultType});
return Builder.CreateCall(F, {Arg});
}
+ case AMDGPU::BI__builtin_amdgcn_global_load_lds: {
+ SmallVector<Value *, 5> Args;
+ unsigned ICEArguments = 0;
+ ASTContext::GetBuiltinTypeError Error;
+ getContext().GetBuiltinType(BuiltinID, Error, &ICEArguments);
+ assert(Error == ASTContext::GE_None && "Should not codegen an error");
+ Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_global_load_lds);
+ llvm::FunctionType *FTy = F->getFunctionType();
+ for (unsigned i = 0, e = E->getNumArgs(); i != e; ++i) {
+ Value *ArgValue = EmitScalarOrConstFoldImmArg(ICEArguments, i, E);
+ llvm::Type *PTy = FTy->getParamType(i);
+ if (PTy != ArgValue->getType()) {
+ if (auto *PtrTy = dyn_cast<llvm::PointerType>(PTy)) {
+ if (PtrTy->getAddressSpace() !=
+ ArgValue->getType()->getPointerAddressSpace()) {
+ ArgValue = Builder.CreateAddrSpaceCast(
+ ArgValue, llvm::PointerType::get(getLLVMContext(),
+ PtrTy->getAddressSpace()));
+ }
+ }
----------------
arsenm wrote:
> Because the builtin can be used not only in OpenCL, I don't think it would be good to put it in SemaOpenCL.
But your test case is written in OpenCL. You can write a run line for other languages if really needed, but for this you don't really need it
> @yxsamliu Do we have Sema for builtin?
Yes, everything does
https://github.com/llvm/llvm-project/pull/93064
More information about the llvm-commits
mailing list