[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