1 //===----- CGCUDANV.cpp - Interface to NVIDIA CUDA Runtime ----------------===// 2 // 3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4 // See https://llvm.org/LICENSE.txt for license information. 5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6 // 7 //===----------------------------------------------------------------------===// 8 // 9 // This provides a class for CUDA code generation targeting the NVIDIA CUDA 10 // runtime library. 11 // 12 //===----------------------------------------------------------------------===// 13 14 #include "CGCUDARuntime.h" 15 #include "CGCXXABI.h" 16 #include "CodeGenFunction.h" 17 #include "CodeGenModule.h" 18 #include "clang/AST/Decl.h" 19 #include "clang/Basic/Cuda.h" 20 #include "clang/CodeGen/CodeGenABITypes.h" 21 #include "clang/CodeGen/ConstantInitBuilder.h" 22 #include "llvm/Frontend/Offloading/Utility.h" 23 #include "llvm/IR/BasicBlock.h" 24 #include "llvm/IR/Constants.h" 25 #include "llvm/IR/DerivedTypes.h" 26 #include "llvm/IR/ReplaceConstant.h" 27 #include "llvm/Support/Format.h" 28 #include "llvm/Support/VirtualFileSystem.h" 29 30 using namespace clang; 31 using namespace CodeGen; 32 33 namespace { 34 constexpr unsigned CudaFatMagic = 0x466243b1; 35 constexpr unsigned HIPFatMagic = 0x48495046; // "HIPF" 36 37 class CGNVCUDARuntime : public CGCUDARuntime { 38 39 private: 40 llvm::IntegerType *IntTy, *SizeTy; 41 llvm::Type *VoidTy; 42 llvm::PointerType *PtrTy; 43 44 /// Convenience reference to LLVM Context 45 llvm::LLVMContext &Context; 46 /// Convenience reference to the current module 47 llvm::Module &TheModule; 48 /// Keeps track of kernel launch stubs and handles emitted in this module 49 struct KernelInfo { 50 llvm::Function *Kernel; // stub function to help launch kernel 51 const Decl *D; 52 }; 53 llvm::SmallVector<KernelInfo, 16> EmittedKernels; 54 // Map a kernel mangled name to a symbol for identifying kernel in host code 55 // For CUDA, the symbol for identifying the kernel is the same as the device 56 // stub function. For HIP, they are different. 57 llvm::DenseMap<StringRef, llvm::GlobalValue *> KernelHandles; 58 // Map a kernel handle to the kernel stub. 59 llvm::DenseMap<llvm::GlobalValue *, llvm::Function *> KernelStubs; 60 struct VarInfo { 61 llvm::GlobalVariable *Var; 62 const VarDecl *D; 63 DeviceVarFlags Flags; 64 }; 65 llvm::SmallVector<VarInfo, 16> DeviceVars; 66 /// Keeps track of variable containing handle of GPU binary. Populated by 67 /// ModuleCtorFunction() and used to create corresponding cleanup calls in 68 /// ModuleDtorFunction() 69 llvm::GlobalVariable *GpuBinaryHandle = nullptr; 70 /// Whether we generate relocatable device code. 71 bool RelocatableDeviceCode; 72 /// Mangle context for device. 73 std::unique_ptr<MangleContext> DeviceMC; 74 /// Some zeros used for GEPs. 75 llvm::Constant *Zeros[2]; 76 77 llvm::FunctionCallee getSetupArgumentFn() const; 78 llvm::FunctionCallee getLaunchFn() const; 79 80 llvm::FunctionType *getRegisterGlobalsFnTy() const; 81 llvm::FunctionType *getCallbackFnTy() const; 82 llvm::FunctionType *getRegisterLinkedBinaryFnTy() const; 83 std::string addPrefixToName(StringRef FuncName) const; 84 std::string addUnderscoredPrefixToName(StringRef FuncName) const; 85 86 /// Creates a function to register all kernel stubs generated in this module. 87 llvm::Function *makeRegisterGlobalsFn(); 88 89 /// Helper function that generates a constant string and returns a pointer to 90 /// the start of the string. The result of this function can be used anywhere 91 /// where the C code specifies const char*. 92 llvm::Constant *makeConstantString(const std::string &Str, 93 const std::string &Name = "") { 94 auto ConstStr = CGM.GetAddrOfConstantCString(Str, Name.c_str()); 95 return llvm::ConstantExpr::getGetElementPtr(ConstStr.getElementType(), 96 ConstStr.getPointer(), Zeros); 97 } 98 99 /// Helper function which generates an initialized constant array from Str, 100 /// and optionally sets section name and alignment. AddNull specifies whether 101 /// the array should nave NUL termination. 102 llvm::Constant *makeConstantArray(StringRef Str, 103 StringRef Name = "", 104 StringRef SectionName = "", 105 unsigned Alignment = 0, 106 bool AddNull = false) { 107 llvm::Constant *Value = 108 llvm::ConstantDataArray::getString(Context, Str, AddNull); 109 auto *GV = new llvm::GlobalVariable( 110 TheModule, Value->getType(), /*isConstant=*/true, 111 llvm::GlobalValue::PrivateLinkage, Value, Name); 112 if (!SectionName.empty()) { 113 GV->setSection(SectionName); 114 // Mark the address as used which make sure that this section isn't 115 // merged and we will really have it in the object file. 116 GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::None); 117 } 118 if (Alignment) 119 GV->setAlignment(llvm::Align(Alignment)); 120 return llvm::ConstantExpr::getGetElementPtr(GV->getValueType(), GV, Zeros); 121 } 122 123 /// Helper function that generates an empty dummy function returning void. 124 llvm::Function *makeDummyFunction(llvm::FunctionType *FnTy) { 125 assert(FnTy->getReturnType()->isVoidTy() && 126 "Can only generate dummy functions returning void!"); 127 llvm::Function *DummyFunc = llvm::Function::Create( 128 FnTy, llvm::GlobalValue::InternalLinkage, "dummy", &TheModule); 129 130 llvm::BasicBlock *DummyBlock = 131 llvm::BasicBlock::Create(Context, "", DummyFunc); 132 CGBuilderTy FuncBuilder(CGM, Context); 133 FuncBuilder.SetInsertPoint(DummyBlock); 134 FuncBuilder.CreateRetVoid(); 135 136 return DummyFunc; 137 } 138 139 void emitDeviceStubBodyLegacy(CodeGenFunction &CGF, FunctionArgList &Args); 140 void emitDeviceStubBodyNew(CodeGenFunction &CGF, FunctionArgList &Args); 141 std::string getDeviceSideName(const NamedDecl *ND) override; 142 143 void registerDeviceVar(const VarDecl *VD, llvm::GlobalVariable &Var, 144 bool Extern, bool Constant) { 145 DeviceVars.push_back({&Var, 146 VD, 147 {DeviceVarFlags::Variable, Extern, Constant, 148 VD->hasAttr<HIPManagedAttr>(), 149 /*Normalized*/ false, 0}}); 150 } 151 void registerDeviceSurf(const VarDecl *VD, llvm::GlobalVariable &Var, 152 bool Extern, int Type) { 153 DeviceVars.push_back({&Var, 154 VD, 155 {DeviceVarFlags::Surface, Extern, /*Constant*/ false, 156 /*Managed*/ false, 157 /*Normalized*/ false, Type}}); 158 } 159 void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var, 160 bool Extern, int Type, bool Normalized) { 161 DeviceVars.push_back({&Var, 162 VD, 163 {DeviceVarFlags::Texture, Extern, /*Constant*/ false, 164 /*Managed*/ false, Normalized, Type}}); 165 } 166 167 /// Creates module constructor function 168 llvm::Function *makeModuleCtorFunction(); 169 /// Creates module destructor function 170 llvm::Function *makeModuleDtorFunction(); 171 /// Transform managed variables for device compilation. 172 void transformManagedVars(); 173 /// Create offloading entries to register globals in RDC mode. 174 void createOffloadingEntries(); 175 176 public: 177 CGNVCUDARuntime(CodeGenModule &CGM); 178 179 llvm::GlobalValue *getKernelHandle(llvm::Function *F, GlobalDecl GD) override; 180 llvm::Function *getKernelStub(llvm::GlobalValue *Handle) override { 181 auto Loc = KernelStubs.find(Handle); 182 assert(Loc != KernelStubs.end()); 183 return Loc->second; 184 } 185 void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) override; 186 void handleVarRegistration(const VarDecl *VD, 187 llvm::GlobalVariable &Var) override; 188 void 189 internalizeDeviceSideVar(const VarDecl *D, 190 llvm::GlobalValue::LinkageTypes &Linkage) override; 191 192 llvm::Function *finalizeModule() override; 193 }; 194 195 } // end anonymous namespace 196 197 std::string CGNVCUDARuntime::addPrefixToName(StringRef FuncName) const { 198 if (CGM.getLangOpts().HIP) 199 return ((Twine("hip") + Twine(FuncName)).str()); 200 return ((Twine("cuda") + Twine(FuncName)).str()); 201 } 202 std::string 203 CGNVCUDARuntime::addUnderscoredPrefixToName(StringRef FuncName) const { 204 if (CGM.getLangOpts().HIP) 205 return ((Twine("__hip") + Twine(FuncName)).str()); 206 return ((Twine("__cuda") + Twine(FuncName)).str()); 207 } 208 209 static std::unique_ptr<MangleContext> InitDeviceMC(CodeGenModule &CGM) { 210 // If the host and device have different C++ ABIs, mark it as the device 211 // mangle context so that the mangling needs to retrieve the additional 212 // device lambda mangling number instead of the regular host one. 213 if (CGM.getContext().getAuxTargetInfo() && 214 CGM.getContext().getTargetInfo().getCXXABI().isMicrosoft() && 215 CGM.getContext().getAuxTargetInfo()->getCXXABI().isItaniumFamily()) { 216 return std::unique_ptr<MangleContext>( 217 CGM.getContext().createDeviceMangleContext( 218 *CGM.getContext().getAuxTargetInfo())); 219 } 220 221 return std::unique_ptr<MangleContext>(CGM.getContext().createMangleContext( 222 CGM.getContext().getAuxTargetInfo())); 223 } 224 225 CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM) 226 : CGCUDARuntime(CGM), Context(CGM.getLLVMContext()), 227 TheModule(CGM.getModule()), 228 RelocatableDeviceCode(CGM.getLangOpts().GPURelocatableDeviceCode), 229 DeviceMC(InitDeviceMC(CGM)) { 230 IntTy = CGM.IntTy; 231 SizeTy = CGM.SizeTy; 232 VoidTy = CGM.VoidTy; 233 Zeros[0] = llvm::ConstantInt::get(SizeTy, 0); 234 Zeros[1] = Zeros[0]; 235 PtrTy = CGM.UnqualPtrTy; 236 } 237 238 llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const { 239 // cudaError_t cudaSetupArgument(void *, size_t, size_t) 240 llvm::Type *Params[] = {PtrTy, SizeTy, SizeTy}; 241 return CGM.CreateRuntimeFunction( 242 llvm::FunctionType::get(IntTy, Params, false), 243 addPrefixToName("SetupArgument")); 244 } 245 246 llvm::FunctionCallee CGNVCUDARuntime::getLaunchFn() const { 247 if (CGM.getLangOpts().HIP) { 248 // hipError_t hipLaunchByPtr(char *); 249 return CGM.CreateRuntimeFunction( 250 llvm::FunctionType::get(IntTy, PtrTy, false), "hipLaunchByPtr"); 251 } 252 // cudaError_t cudaLaunch(char *); 253 return CGM.CreateRuntimeFunction(llvm::FunctionType::get(IntTy, PtrTy, false), 254 "cudaLaunch"); 255 } 256 257 llvm::FunctionType *CGNVCUDARuntime::getRegisterGlobalsFnTy() const { 258 return llvm::FunctionType::get(VoidTy, PtrTy, false); 259 } 260 261 llvm::FunctionType *CGNVCUDARuntime::getCallbackFnTy() const { 262 return llvm::FunctionType::get(VoidTy, PtrTy, false); 263 } 264 265 llvm::FunctionType *CGNVCUDARuntime::getRegisterLinkedBinaryFnTy() const { 266 llvm::Type *Params[] = {llvm::PointerType::getUnqual(Context), PtrTy, PtrTy, 267 llvm::PointerType::getUnqual(Context)}; 268 return llvm::FunctionType::get(VoidTy, Params, false); 269 } 270 271 std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) { 272 GlobalDecl GD; 273 // D could be either a kernel or a variable. 274 if (auto *FD = dyn_cast<FunctionDecl>(ND)) 275 GD = GlobalDecl(FD, KernelReferenceKind::Kernel); 276 else 277 GD = GlobalDecl(ND); 278 std::string DeviceSideName; 279 MangleContext *MC; 280 if (CGM.getLangOpts().CUDAIsDevice) 281 MC = &CGM.getCXXABI().getMangleContext(); 282 else 283 MC = DeviceMC.get(); 284 if (MC->shouldMangleDeclName(ND)) { 285 SmallString<256> Buffer; 286 llvm::raw_svector_ostream Out(Buffer); 287 MC->mangleName(GD, Out); 288 DeviceSideName = std::string(Out.str()); 289 } else 290 DeviceSideName = std::string(ND->getIdentifier()->getName()); 291 292 // Make unique name for device side static file-scope variable for HIP. 293 if (CGM.getContext().shouldExternalize(ND) && 294 CGM.getLangOpts().GPURelocatableDeviceCode) { 295 SmallString<256> Buffer; 296 llvm::raw_svector_ostream Out(Buffer); 297 Out << DeviceSideName; 298 CGM.printPostfixForExternalizedDecl(Out, ND); 299 DeviceSideName = std::string(Out.str()); 300 } 301 return DeviceSideName; 302 } 303 304 void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF, 305 FunctionArgList &Args) { 306 EmittedKernels.push_back({CGF.CurFn, CGF.CurFuncDecl}); 307 if (auto *GV = 308 dyn_cast<llvm::GlobalVariable>(KernelHandles[CGF.CurFn->getName()])) { 309 GV->setLinkage(CGF.CurFn->getLinkage()); 310 GV->setInitializer(CGF.CurFn); 311 } 312 if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(), 313 CudaFeature::CUDA_USES_NEW_LAUNCH) || 314 (CGF.getLangOpts().HIP && CGF.getLangOpts().HIPUseNewLaunchAPI)) 315 emitDeviceStubBodyNew(CGF, Args); 316 else 317 emitDeviceStubBodyLegacy(CGF, Args); 318 } 319 320 // CUDA 9.0+ uses new way to launch kernels. Parameters are packed in a local 321 // array and kernels are launched using cudaLaunchKernel(). 322 void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF, 323 FunctionArgList &Args) { 324 // Build the shadow stack entry at the very start of the function. 325 326 // Calculate amount of space we will need for all arguments. If we have no 327 // args, allocate a single pointer so we still have a valid pointer to the 328 // argument array that we can pass to runtime, even if it will be unused. 329 Address KernelArgs = CGF.CreateTempAlloca( 330 PtrTy, CharUnits::fromQuantity(16), "kernel_args", 331 llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size()))); 332 // Store pointers to the arguments in a locally allocated launch_args. 333 for (unsigned i = 0; i < Args.size(); ++i) { 334 llvm::Value* VarPtr = CGF.GetAddrOfLocalVar(Args[i]).getPointer(); 335 llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, PtrTy); 336 CGF.Builder.CreateDefaultAlignedStore( 337 VoidVarPtr, 338 CGF.Builder.CreateConstGEP1_32(PtrTy, KernelArgs.getPointer(), i)); 339 } 340 341 llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end"); 342 343 // Lookup cudaLaunchKernel/hipLaunchKernel function. 344 // HIP kernel launching API name depends on -fgpu-default-stream option. For 345 // the default value 'legacy', it is hipLaunchKernel. For 'per-thread', 346 // it is hipLaunchKernel_spt. 347 // cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim, 348 // void **args, size_t sharedMem, 349 // cudaStream_t stream); 350 // hipError_t hipLaunchKernel[_spt](const void *func, dim3 gridDim, 351 // dim3 blockDim, void **args, 352 // size_t sharedMem, hipStream_t stream); 353 TranslationUnitDecl *TUDecl = CGM.getContext().getTranslationUnitDecl(); 354 DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl); 355 std::string KernelLaunchAPI = "LaunchKernel"; 356 if (CGF.getLangOpts().GPUDefaultStream == 357 LangOptions::GPUDefaultStreamKind::PerThread) { 358 if (CGF.getLangOpts().HIP) 359 KernelLaunchAPI = KernelLaunchAPI + "_spt"; 360 else if (CGF.getLangOpts().CUDA) 361 KernelLaunchAPI = KernelLaunchAPI + "_ptsz"; 362 } 363 auto LaunchKernelName = addPrefixToName(KernelLaunchAPI); 364 IdentifierInfo &cudaLaunchKernelII = 365 CGM.getContext().Idents.get(LaunchKernelName); 366 FunctionDecl *cudaLaunchKernelFD = nullptr; 367 for (auto *Result : DC->lookup(&cudaLaunchKernelII)) { 368 if (FunctionDecl *FD = dyn_cast<FunctionDecl>(Result)) 369 cudaLaunchKernelFD = FD; 370 } 371 372 if (cudaLaunchKernelFD == nullptr) { 373 CGM.Error(CGF.CurFuncDecl->getLocation(), 374 "Can't find declaration for " + LaunchKernelName); 375 return; 376 } 377 // Create temporary dim3 grid_dim, block_dim. 378 ParmVarDecl *GridDimParam = cudaLaunchKernelFD->getParamDecl(1); 379 QualType Dim3Ty = GridDimParam->getType(); 380 Address GridDim = 381 CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "grid_dim"); 382 Address BlockDim = 383 CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "block_dim"); 384 Address ShmemSize = 385 CGF.CreateTempAlloca(SizeTy, CGM.getSizeAlign(), "shmem_size"); 386 Address Stream = CGF.CreateTempAlloca(PtrTy, CGM.getPointerAlign(), "stream"); 387 llvm::FunctionCallee cudaPopConfigFn = CGM.CreateRuntimeFunction( 388 llvm::FunctionType::get(IntTy, 389 {/*gridDim=*/GridDim.getType(), 390 /*blockDim=*/BlockDim.getType(), 391 /*ShmemSize=*/ShmemSize.getType(), 392 /*Stream=*/Stream.getType()}, 393 /*isVarArg=*/false), 394 addUnderscoredPrefixToName("PopCallConfiguration")); 395 396 CGF.EmitRuntimeCallOrInvoke(cudaPopConfigFn, 397 {GridDim.getPointer(), BlockDim.getPointer(), 398 ShmemSize.getPointer(), Stream.getPointer()}); 399 400 // Emit the call to cudaLaunch 401 llvm::Value *Kernel = 402 CGF.Builder.CreatePointerCast(KernelHandles[CGF.CurFn->getName()], PtrTy); 403 CallArgList LaunchKernelArgs; 404 LaunchKernelArgs.add(RValue::get(Kernel), 405 cudaLaunchKernelFD->getParamDecl(0)->getType()); 406 LaunchKernelArgs.add(RValue::getAggregate(GridDim), Dim3Ty); 407 LaunchKernelArgs.add(RValue::getAggregate(BlockDim), Dim3Ty); 408 LaunchKernelArgs.add(RValue::get(KernelArgs.getPointer()), 409 cudaLaunchKernelFD->getParamDecl(3)->getType()); 410 LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(ShmemSize)), 411 cudaLaunchKernelFD->getParamDecl(4)->getType()); 412 LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(Stream)), 413 cudaLaunchKernelFD->getParamDecl(5)->getType()); 414 415 QualType QT = cudaLaunchKernelFD->getType(); 416 QualType CQT = QT.getCanonicalType(); 417 llvm::Type *Ty = CGM.getTypes().ConvertType(CQT); 418 llvm::FunctionType *FTy = cast<llvm::FunctionType>(Ty); 419 420 const CGFunctionInfo &FI = 421 CGM.getTypes().arrangeFunctionDeclaration(cudaLaunchKernelFD); 422 llvm::FunctionCallee cudaLaunchKernelFn = 423 CGM.CreateRuntimeFunction(FTy, LaunchKernelName); 424 CGF.EmitCall(FI, CGCallee::forDirect(cudaLaunchKernelFn), ReturnValueSlot(), 425 LaunchKernelArgs); 426 CGF.EmitBranch(EndBlock); 427 428 CGF.EmitBlock(EndBlock); 429 } 430 431 void CGNVCUDARuntime::emitDeviceStubBodyLegacy(CodeGenFunction &CGF, 432 FunctionArgList &Args) { 433 // Emit a call to cudaSetupArgument for each arg in Args. 434 llvm::FunctionCallee cudaSetupArgFn = getSetupArgumentFn(); 435 llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end"); 436 CharUnits Offset = CharUnits::Zero(); 437 for (const VarDecl *A : Args) { 438 auto TInfo = CGM.getContext().getTypeInfoInChars(A->getType()); 439 Offset = Offset.alignTo(TInfo.Align); 440 llvm::Value *Args[] = { 441 CGF.Builder.CreatePointerCast(CGF.GetAddrOfLocalVar(A).getPointer(), 442 PtrTy), 443 llvm::ConstantInt::get(SizeTy, TInfo.Width.getQuantity()), 444 llvm::ConstantInt::get(SizeTy, Offset.getQuantity()), 445 }; 446 llvm::CallBase *CB = CGF.EmitRuntimeCallOrInvoke(cudaSetupArgFn, Args); 447 llvm::Constant *Zero = llvm::ConstantInt::get(IntTy, 0); 448 llvm::Value *CBZero = CGF.Builder.CreateICmpEQ(CB, Zero); 449 llvm::BasicBlock *NextBlock = CGF.createBasicBlock("setup.next"); 450 CGF.Builder.CreateCondBr(CBZero, NextBlock, EndBlock); 451 CGF.EmitBlock(NextBlock); 452 Offset += TInfo.Width; 453 } 454 455 // Emit the call to cudaLaunch 456 llvm::FunctionCallee cudaLaunchFn = getLaunchFn(); 457 llvm::Value *Arg = 458 CGF.Builder.CreatePointerCast(KernelHandles[CGF.CurFn->getName()], PtrTy); 459 CGF.EmitRuntimeCallOrInvoke(cudaLaunchFn, Arg); 460 CGF.EmitBranch(EndBlock); 461 462 CGF.EmitBlock(EndBlock); 463 } 464 465 // Replace the original variable Var with the address loaded from variable 466 // ManagedVar populated by HIP runtime. 467 static void replaceManagedVar(llvm::GlobalVariable *Var, 468 llvm::GlobalVariable *ManagedVar) { 469 SmallVector<SmallVector<llvm::User *, 8>, 8> WorkList; 470 for (auto &&VarUse : Var->uses()) { 471 WorkList.push_back({VarUse.getUser()}); 472 } 473 while (!WorkList.empty()) { 474 auto &&WorkItem = WorkList.pop_back_val(); 475 auto *U = WorkItem.back(); 476 if (isa<llvm::ConstantExpr>(U)) { 477 for (auto &&UU : U->uses()) { 478 WorkItem.push_back(UU.getUser()); 479 WorkList.push_back(WorkItem); 480 WorkItem.pop_back(); 481 } 482 continue; 483 } 484 if (auto *I = dyn_cast<llvm::Instruction>(U)) { 485 llvm::Value *OldV = Var; 486 llvm::Instruction *NewV = 487 new llvm::LoadInst(Var->getType(), ManagedVar, "ld.managed", false, 488 llvm::Align(Var->getAlignment()), I); 489 WorkItem.pop_back(); 490 // Replace constant expressions directly or indirectly using the managed 491 // variable with instructions. 492 for (auto &&Op : WorkItem) { 493 auto *CE = cast<llvm::ConstantExpr>(Op); 494 auto *NewInst = CE->getAsInstruction(I); 495 NewInst->replaceUsesOfWith(OldV, NewV); 496 OldV = CE; 497 NewV = NewInst; 498 } 499 I->replaceUsesOfWith(OldV, NewV); 500 } else { 501 llvm_unreachable("Invalid use of managed variable"); 502 } 503 } 504 } 505 506 /// Creates a function that sets up state on the host side for CUDA objects that 507 /// have a presence on both the host and device sides. Specifically, registers 508 /// the host side of kernel functions and device global variables with the CUDA 509 /// runtime. 510 /// \code 511 /// void __cuda_register_globals(void** GpuBinaryHandle) { 512 /// __cudaRegisterFunction(GpuBinaryHandle,Kernel0,...); 513 /// ... 514 /// __cudaRegisterFunction(GpuBinaryHandle,KernelM,...); 515 /// __cudaRegisterVar(GpuBinaryHandle, GlobalVar0, ...); 516 /// ... 517 /// __cudaRegisterVar(GpuBinaryHandle, GlobalVarN, ...); 518 /// } 519 /// \endcode 520 llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() { 521 // No need to register anything 522 if (EmittedKernels.empty() && DeviceVars.empty()) 523 return nullptr; 524 525 llvm::Function *RegisterKernelsFunc = llvm::Function::Create( 526 getRegisterGlobalsFnTy(), llvm::GlobalValue::InternalLinkage, 527 addUnderscoredPrefixToName("_register_globals"), &TheModule); 528 llvm::BasicBlock *EntryBB = 529 llvm::BasicBlock::Create(Context, "entry", RegisterKernelsFunc); 530 CGBuilderTy Builder(CGM, Context); 531 Builder.SetInsertPoint(EntryBB); 532 533 // void __cudaRegisterFunction(void **, const char *, char *, const char *, 534 // int, uint3*, uint3*, dim3*, dim3*, int*) 535 llvm::Type *RegisterFuncParams[] = { 536 PtrTy, PtrTy, PtrTy, PtrTy, IntTy, 537 PtrTy, PtrTy, PtrTy, PtrTy, llvm::PointerType::getUnqual(Context)}; 538 llvm::FunctionCallee RegisterFunc = CGM.CreateRuntimeFunction( 539 llvm::FunctionType::get(IntTy, RegisterFuncParams, false), 540 addUnderscoredPrefixToName("RegisterFunction")); 541 542 // Extract GpuBinaryHandle passed as the first argument passed to 543 // __cuda_register_globals() and generate __cudaRegisterFunction() call for 544 // each emitted kernel. 545 llvm::Argument &GpuBinaryHandlePtr = *RegisterKernelsFunc->arg_begin(); 546 for (auto &&I : EmittedKernels) { 547 llvm::Constant *KernelName = 548 makeConstantString(getDeviceSideName(cast<NamedDecl>(I.D))); 549 llvm::Constant *NullPtr = llvm::ConstantPointerNull::get(PtrTy); 550 llvm::Value *Args[] = { 551 &GpuBinaryHandlePtr, 552 KernelHandles[I.Kernel->getName()], 553 KernelName, 554 KernelName, 555 llvm::ConstantInt::get(IntTy, -1), 556 NullPtr, 557 NullPtr, 558 NullPtr, 559 NullPtr, 560 llvm::ConstantPointerNull::get(llvm::PointerType::getUnqual(Context))}; 561 Builder.CreateCall(RegisterFunc, Args); 562 } 563 564 llvm::Type *VarSizeTy = IntTy; 565 // For HIP or CUDA 9.0+, device variable size is type of `size_t`. 566 if (CGM.getLangOpts().HIP || 567 ToCudaVersion(CGM.getTarget().getSDKVersion()) >= CudaVersion::CUDA_90) 568 VarSizeTy = SizeTy; 569 570 // void __cudaRegisterVar(void **, char *, char *, const char *, 571 // int, int, int, int) 572 llvm::Type *RegisterVarParams[] = {PtrTy, PtrTy, PtrTy, PtrTy, 573 IntTy, VarSizeTy, IntTy, IntTy}; 574 llvm::FunctionCallee RegisterVar = CGM.CreateRuntimeFunction( 575 llvm::FunctionType::get(VoidTy, RegisterVarParams, false), 576 addUnderscoredPrefixToName("RegisterVar")); 577 // void __hipRegisterManagedVar(void **, char *, char *, const char *, 578 // size_t, unsigned) 579 llvm::Type *RegisterManagedVarParams[] = {PtrTy, PtrTy, PtrTy, 580 PtrTy, VarSizeTy, IntTy}; 581 llvm::FunctionCallee RegisterManagedVar = CGM.CreateRuntimeFunction( 582 llvm::FunctionType::get(VoidTy, RegisterManagedVarParams, false), 583 addUnderscoredPrefixToName("RegisterManagedVar")); 584 // void __cudaRegisterSurface(void **, const struct surfaceReference *, 585 // const void **, const char *, int, int); 586 llvm::FunctionCallee RegisterSurf = CGM.CreateRuntimeFunction( 587 llvm::FunctionType::get( 588 VoidTy, {PtrTy, PtrTy, PtrTy, PtrTy, IntTy, IntTy}, false), 589 addUnderscoredPrefixToName("RegisterSurface")); 590 // void __cudaRegisterTexture(void **, const struct textureReference *, 591 // const void **, const char *, int, int, int) 592 llvm::FunctionCallee RegisterTex = CGM.CreateRuntimeFunction( 593 llvm::FunctionType::get( 594 VoidTy, {PtrTy, PtrTy, PtrTy, PtrTy, IntTy, IntTy, IntTy}, false), 595 addUnderscoredPrefixToName("RegisterTexture")); 596 for (auto &&Info : DeviceVars) { 597 llvm::GlobalVariable *Var = Info.Var; 598 assert((!Var->isDeclaration() || Info.Flags.isManaged()) && 599 "External variables should not show up here, except HIP managed " 600 "variables"); 601 llvm::Constant *VarName = makeConstantString(getDeviceSideName(Info.D)); 602 switch (Info.Flags.getKind()) { 603 case DeviceVarFlags::Variable: { 604 uint64_t VarSize = 605 CGM.getDataLayout().getTypeAllocSize(Var->getValueType()); 606 if (Info.Flags.isManaged()) { 607 auto *ManagedVar = new llvm::GlobalVariable( 608 CGM.getModule(), Var->getType(), 609 /*isConstant=*/false, Var->getLinkage(), 610 /*Init=*/Var->isDeclaration() 611 ? nullptr 612 : llvm::ConstantPointerNull::get(Var->getType()), 613 /*Name=*/"", /*InsertBefore=*/nullptr, 614 llvm::GlobalVariable::NotThreadLocal); 615 ManagedVar->setDSOLocal(Var->isDSOLocal()); 616 ManagedVar->setVisibility(Var->getVisibility()); 617 ManagedVar->setExternallyInitialized(true); 618 ManagedVar->takeName(Var); 619 Var->setName(Twine(ManagedVar->getName() + ".managed")); 620 replaceManagedVar(Var, ManagedVar); 621 llvm::Value *Args[] = { 622 &GpuBinaryHandlePtr, 623 ManagedVar, 624 Var, 625 VarName, 626 llvm::ConstantInt::get(VarSizeTy, VarSize), 627 llvm::ConstantInt::get(IntTy, Var->getAlignment())}; 628 if (!Var->isDeclaration()) 629 Builder.CreateCall(RegisterManagedVar, Args); 630 } else { 631 llvm::Value *Args[] = { 632 &GpuBinaryHandlePtr, 633 Var, 634 VarName, 635 VarName, 636 llvm::ConstantInt::get(IntTy, Info.Flags.isExtern()), 637 llvm::ConstantInt::get(VarSizeTy, VarSize), 638 llvm::ConstantInt::get(IntTy, Info.Flags.isConstant()), 639 llvm::ConstantInt::get(IntTy, 0)}; 640 Builder.CreateCall(RegisterVar, Args); 641 } 642 break; 643 } 644 case DeviceVarFlags::Surface: 645 Builder.CreateCall( 646 RegisterSurf, 647 {&GpuBinaryHandlePtr, Var, VarName, VarName, 648 llvm::ConstantInt::get(IntTy, Info.Flags.getSurfTexType()), 649 llvm::ConstantInt::get(IntTy, Info.Flags.isExtern())}); 650 break; 651 case DeviceVarFlags::Texture: 652 Builder.CreateCall( 653 RegisterTex, 654 {&GpuBinaryHandlePtr, Var, VarName, VarName, 655 llvm::ConstantInt::get(IntTy, Info.Flags.getSurfTexType()), 656 llvm::ConstantInt::get(IntTy, Info.Flags.isNormalized()), 657 llvm::ConstantInt::get(IntTy, Info.Flags.isExtern())}); 658 break; 659 } 660 } 661 662 Builder.CreateRetVoid(); 663 return RegisterKernelsFunc; 664 } 665 666 /// Creates a global constructor function for the module: 667 /// 668 /// For CUDA: 669 /// \code 670 /// void __cuda_module_ctor() { 671 /// Handle = __cudaRegisterFatBinary(GpuBinaryBlob); 672 /// __cuda_register_globals(Handle); 673 /// } 674 /// \endcode 675 /// 676 /// For HIP: 677 /// \code 678 /// void __hip_module_ctor() { 679 /// if (__hip_gpubin_handle == 0) { 680 /// __hip_gpubin_handle = __hipRegisterFatBinary(GpuBinaryBlob); 681 /// __hip_register_globals(__hip_gpubin_handle); 682 /// } 683 /// } 684 /// \endcode 685 llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() { 686 bool IsHIP = CGM.getLangOpts().HIP; 687 bool IsCUDA = CGM.getLangOpts().CUDA; 688 // No need to generate ctors/dtors if there is no GPU binary. 689 StringRef CudaGpuBinaryFileName = CGM.getCodeGenOpts().CudaGpuBinaryFileName; 690 if (CudaGpuBinaryFileName.empty() && !IsHIP) 691 return nullptr; 692 if ((IsHIP || (IsCUDA && !RelocatableDeviceCode)) && EmittedKernels.empty() && 693 DeviceVars.empty()) 694 return nullptr; 695 696 // void __{cuda|hip}_register_globals(void* handle); 697 llvm::Function *RegisterGlobalsFunc = makeRegisterGlobalsFn(); 698 // We always need a function to pass in as callback. Create a dummy 699 // implementation if we don't need to register anything. 700 if (RelocatableDeviceCode && !RegisterGlobalsFunc) 701 RegisterGlobalsFunc = makeDummyFunction(getRegisterGlobalsFnTy()); 702 703 // void ** __{cuda|hip}RegisterFatBinary(void *); 704 llvm::FunctionCallee RegisterFatbinFunc = CGM.CreateRuntimeFunction( 705 llvm::FunctionType::get(PtrTy, PtrTy, false), 706 addUnderscoredPrefixToName("RegisterFatBinary")); 707 // struct { int magic, int version, void * gpu_binary, void * dont_care }; 708 llvm::StructType *FatbinWrapperTy = 709 llvm::StructType::get(IntTy, IntTy, PtrTy, PtrTy); 710 711 // Register GPU binary with the CUDA runtime, store returned handle in a 712 // global variable and save a reference in GpuBinaryHandle to be cleaned up 713 // in destructor on exit. Then associate all known kernels with the GPU binary 714 // handle so CUDA runtime can figure out what to call on the GPU side. 715 std::unique_ptr<llvm::MemoryBuffer> CudaGpuBinary = nullptr; 716 if (!CudaGpuBinaryFileName.empty()) { 717 auto VFS = CGM.getFileSystem(); 718 auto CudaGpuBinaryOrErr = 719 VFS->getBufferForFile(CudaGpuBinaryFileName, -1, false); 720 if (std::error_code EC = CudaGpuBinaryOrErr.getError()) { 721 CGM.getDiags().Report(diag::err_cannot_open_file) 722 << CudaGpuBinaryFileName << EC.message(); 723 return nullptr; 724 } 725 CudaGpuBinary = std::move(CudaGpuBinaryOrErr.get()); 726 } 727 728 llvm::Function *ModuleCtorFunc = llvm::Function::Create( 729 llvm::FunctionType::get(VoidTy, false), 730 llvm::GlobalValue::InternalLinkage, 731 addUnderscoredPrefixToName("_module_ctor"), &TheModule); 732 llvm::BasicBlock *CtorEntryBB = 733 llvm::BasicBlock::Create(Context, "entry", ModuleCtorFunc); 734 CGBuilderTy CtorBuilder(CGM, Context); 735 736 CtorBuilder.SetInsertPoint(CtorEntryBB); 737 738 const char *FatbinConstantName; 739 const char *FatbinSectionName; 740 const char *ModuleIDSectionName; 741 StringRef ModuleIDPrefix; 742 llvm::Constant *FatBinStr; 743 unsigned FatMagic; 744 if (IsHIP) { 745 FatbinConstantName = ".hip_fatbin"; 746 FatbinSectionName = ".hipFatBinSegment"; 747 748 ModuleIDSectionName = "__hip_module_id"; 749 ModuleIDPrefix = "__hip_"; 750 751 if (CudaGpuBinary) { 752 // If fatbin is available from early finalization, create a string 753 // literal containing the fat binary loaded from the given file. 754 const unsigned HIPCodeObjectAlign = 4096; 755 FatBinStr = makeConstantArray(std::string(CudaGpuBinary->getBuffer()), "", 756 FatbinConstantName, HIPCodeObjectAlign); 757 } else { 758 // If fatbin is not available, create an external symbol 759 // __hip_fatbin in section .hip_fatbin. The external symbol is supposed 760 // to contain the fat binary but will be populated somewhere else, 761 // e.g. by lld through link script. 762 FatBinStr = new llvm::GlobalVariable( 763 CGM.getModule(), CGM.Int8Ty, 764 /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage, nullptr, 765 "__hip_fatbin", nullptr, 766 llvm::GlobalVariable::NotThreadLocal); 767 cast<llvm::GlobalVariable>(FatBinStr)->setSection(FatbinConstantName); 768 } 769 770 FatMagic = HIPFatMagic; 771 } else { 772 if (RelocatableDeviceCode) 773 FatbinConstantName = CGM.getTriple().isMacOSX() 774 ? "__NV_CUDA,__nv_relfatbin" 775 : "__nv_relfatbin"; 776 else 777 FatbinConstantName = 778 CGM.getTriple().isMacOSX() ? "__NV_CUDA,__nv_fatbin" : ".nv_fatbin"; 779 // NVIDIA's cuobjdump looks for fatbins in this section. 780 FatbinSectionName = 781 CGM.getTriple().isMacOSX() ? "__NV_CUDA,__fatbin" : ".nvFatBinSegment"; 782 783 ModuleIDSectionName = CGM.getTriple().isMacOSX() 784 ? "__NV_CUDA,__nv_module_id" 785 : "__nv_module_id"; 786 ModuleIDPrefix = "__nv_"; 787 788 // For CUDA, create a string literal containing the fat binary loaded from 789 // the given file. 790 FatBinStr = makeConstantArray(std::string(CudaGpuBinary->getBuffer()), "", 791 FatbinConstantName, 8); 792 FatMagic = CudaFatMagic; 793 } 794 795 // Create initialized wrapper structure that points to the loaded GPU binary 796 ConstantInitBuilder Builder(CGM); 797 auto Values = Builder.beginStruct(FatbinWrapperTy); 798 // Fatbin wrapper magic. 799 Values.addInt(IntTy, FatMagic); 800 // Fatbin version. 801 Values.addInt(IntTy, 1); 802 // Data. 803 Values.add(FatBinStr); 804 // Unused in fatbin v1. 805 Values.add(llvm::ConstantPointerNull::get(PtrTy)); 806 llvm::GlobalVariable *FatbinWrapper = Values.finishAndCreateGlobal( 807 addUnderscoredPrefixToName("_fatbin_wrapper"), CGM.getPointerAlign(), 808 /*constant*/ true); 809 FatbinWrapper->setSection(FatbinSectionName); 810 811 // There is only one HIP fat binary per linked module, however there are 812 // multiple constructor functions. Make sure the fat binary is registered 813 // only once. The constructor functions are executed by the dynamic loader 814 // before the program gains control. The dynamic loader cannot execute the 815 // constructor functions concurrently since doing that would not guarantee 816 // thread safety of the loaded program. Therefore we can assume sequential 817 // execution of constructor functions here. 818 if (IsHIP) { 819 auto Linkage = CudaGpuBinary ? llvm::GlobalValue::InternalLinkage : 820 llvm::GlobalValue::LinkOnceAnyLinkage; 821 llvm::BasicBlock *IfBlock = 822 llvm::BasicBlock::Create(Context, "if", ModuleCtorFunc); 823 llvm::BasicBlock *ExitBlock = 824 llvm::BasicBlock::Create(Context, "exit", ModuleCtorFunc); 825 // The name, size, and initialization pattern of this variable is part 826 // of HIP ABI. 827 GpuBinaryHandle = new llvm::GlobalVariable( 828 TheModule, PtrTy, /*isConstant=*/false, Linkage, 829 /*Initializer=*/llvm::ConstantPointerNull::get(PtrTy), 830 "__hip_gpubin_handle"); 831 if (Linkage == llvm::GlobalValue::LinkOnceAnyLinkage) 832 GpuBinaryHandle->setComdat( 833 CGM.getModule().getOrInsertComdat(GpuBinaryHandle->getName())); 834 GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign()); 835 // Prevent the weak symbol in different shared libraries being merged. 836 if (Linkage != llvm::GlobalValue::InternalLinkage) 837 GpuBinaryHandle->setVisibility(llvm::GlobalValue::HiddenVisibility); 838 Address GpuBinaryAddr( 839 GpuBinaryHandle, PtrTy, 840 CharUnits::fromQuantity(GpuBinaryHandle->getAlignment())); 841 { 842 auto *HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr); 843 llvm::Constant *Zero = 844 llvm::Constant::getNullValue(HandleValue->getType()); 845 llvm::Value *EQZero = CtorBuilder.CreateICmpEQ(HandleValue, Zero); 846 CtorBuilder.CreateCondBr(EQZero, IfBlock, ExitBlock); 847 } 848 { 849 CtorBuilder.SetInsertPoint(IfBlock); 850 // GpuBinaryHandle = __hipRegisterFatBinary(&FatbinWrapper); 851 llvm::CallInst *RegisterFatbinCall = 852 CtorBuilder.CreateCall(RegisterFatbinFunc, FatbinWrapper); 853 CtorBuilder.CreateStore(RegisterFatbinCall, GpuBinaryAddr); 854 CtorBuilder.CreateBr(ExitBlock); 855 } 856 { 857 CtorBuilder.SetInsertPoint(ExitBlock); 858 // Call __hip_register_globals(GpuBinaryHandle); 859 if (RegisterGlobalsFunc) { 860 auto *HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr); 861 CtorBuilder.CreateCall(RegisterGlobalsFunc, HandleValue); 862 } 863 } 864 } else if (!RelocatableDeviceCode) { 865 // Register binary with CUDA runtime. This is substantially different in 866 // default mode vs. separate compilation! 867 // GpuBinaryHandle = __cudaRegisterFatBinary(&FatbinWrapper); 868 llvm::CallInst *RegisterFatbinCall = 869 CtorBuilder.CreateCall(RegisterFatbinFunc, FatbinWrapper); 870 GpuBinaryHandle = new llvm::GlobalVariable( 871 TheModule, PtrTy, false, llvm::GlobalValue::InternalLinkage, 872 llvm::ConstantPointerNull::get(PtrTy), "__cuda_gpubin_handle"); 873 GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign()); 874 CtorBuilder.CreateAlignedStore(RegisterFatbinCall, GpuBinaryHandle, 875 CGM.getPointerAlign()); 876 877 // Call __cuda_register_globals(GpuBinaryHandle); 878 if (RegisterGlobalsFunc) 879 CtorBuilder.CreateCall(RegisterGlobalsFunc, RegisterFatbinCall); 880 881 // Call __cudaRegisterFatBinaryEnd(Handle) if this CUDA version needs it. 882 if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(), 883 CudaFeature::CUDA_USES_FATBIN_REGISTER_END)) { 884 // void __cudaRegisterFatBinaryEnd(void **); 885 llvm::FunctionCallee RegisterFatbinEndFunc = CGM.CreateRuntimeFunction( 886 llvm::FunctionType::get(VoidTy, PtrTy, false), 887 "__cudaRegisterFatBinaryEnd"); 888 CtorBuilder.CreateCall(RegisterFatbinEndFunc, RegisterFatbinCall); 889 } 890 } else { 891 // Generate a unique module ID. 892 SmallString<64> ModuleID; 893 llvm::raw_svector_ostream OS(ModuleID); 894 OS << ModuleIDPrefix << llvm::format("%" PRIx64, FatbinWrapper->getGUID()); 895 llvm::Constant *ModuleIDConstant = makeConstantArray( 896 std::string(ModuleID), "", ModuleIDSectionName, 32, /*AddNull=*/true); 897 898 // Create an alias for the FatbinWrapper that nvcc will look for. 899 llvm::GlobalAlias::create(llvm::GlobalValue::ExternalLinkage, 900 Twine("__fatbinwrap") + ModuleID, FatbinWrapper); 901 902 // void __cudaRegisterLinkedBinary%ModuleID%(void (*)(void *), void *, 903 // void *, void (*)(void **)) 904 SmallString<128> RegisterLinkedBinaryName("__cudaRegisterLinkedBinary"); 905 RegisterLinkedBinaryName += ModuleID; 906 llvm::FunctionCallee RegisterLinkedBinaryFunc = CGM.CreateRuntimeFunction( 907 getRegisterLinkedBinaryFnTy(), RegisterLinkedBinaryName); 908 909 assert(RegisterGlobalsFunc && "Expecting at least dummy function!"); 910 llvm::Value *Args[] = {RegisterGlobalsFunc, FatbinWrapper, ModuleIDConstant, 911 makeDummyFunction(getCallbackFnTy())}; 912 CtorBuilder.CreateCall(RegisterLinkedBinaryFunc, Args); 913 } 914 915 // Create destructor and register it with atexit() the way NVCC does it. Doing 916 // it during regular destructor phase worked in CUDA before 9.2 but results in 917 // double-free in 9.2. 918 if (llvm::Function *CleanupFn = makeModuleDtorFunction()) { 919 // extern "C" int atexit(void (*f)(void)); 920 llvm::FunctionType *AtExitTy = 921 llvm::FunctionType::get(IntTy, CleanupFn->getType(), false); 922 llvm::FunctionCallee AtExitFunc = 923 CGM.CreateRuntimeFunction(AtExitTy, "atexit", llvm::AttributeList(), 924 /*Local=*/true); 925 CtorBuilder.CreateCall(AtExitFunc, CleanupFn); 926 } 927 928 CtorBuilder.CreateRetVoid(); 929 return ModuleCtorFunc; 930 } 931 932 /// Creates a global destructor function that unregisters the GPU code blob 933 /// registered by constructor. 934 /// 935 /// For CUDA: 936 /// \code 937 /// void __cuda_module_dtor() { 938 /// __cudaUnregisterFatBinary(Handle); 939 /// } 940 /// \endcode 941 /// 942 /// For HIP: 943 /// \code 944 /// void __hip_module_dtor() { 945 /// if (__hip_gpubin_handle) { 946 /// __hipUnregisterFatBinary(__hip_gpubin_handle); 947 /// __hip_gpubin_handle = 0; 948 /// } 949 /// } 950 /// \endcode 951 llvm::Function *CGNVCUDARuntime::makeModuleDtorFunction() { 952 // No need for destructor if we don't have a handle to unregister. 953 if (!GpuBinaryHandle) 954 return nullptr; 955 956 // void __cudaUnregisterFatBinary(void ** handle); 957 llvm::FunctionCallee UnregisterFatbinFunc = CGM.CreateRuntimeFunction( 958 llvm::FunctionType::get(VoidTy, PtrTy, false), 959 addUnderscoredPrefixToName("UnregisterFatBinary")); 960 961 llvm::Function *ModuleDtorFunc = llvm::Function::Create( 962 llvm::FunctionType::get(VoidTy, false), 963 llvm::GlobalValue::InternalLinkage, 964 addUnderscoredPrefixToName("_module_dtor"), &TheModule); 965 966 llvm::BasicBlock *DtorEntryBB = 967 llvm::BasicBlock::Create(Context, "entry", ModuleDtorFunc); 968 CGBuilderTy DtorBuilder(CGM, Context); 969 DtorBuilder.SetInsertPoint(DtorEntryBB); 970 971 Address GpuBinaryAddr( 972 GpuBinaryHandle, GpuBinaryHandle->getValueType(), 973 CharUnits::fromQuantity(GpuBinaryHandle->getAlignment())); 974 auto *HandleValue = DtorBuilder.CreateLoad(GpuBinaryAddr); 975 // There is only one HIP fat binary per linked module, however there are 976 // multiple destructor functions. Make sure the fat binary is unregistered 977 // only once. 978 if (CGM.getLangOpts().HIP) { 979 llvm::BasicBlock *IfBlock = 980 llvm::BasicBlock::Create(Context, "if", ModuleDtorFunc); 981 llvm::BasicBlock *ExitBlock = 982 llvm::BasicBlock::Create(Context, "exit", ModuleDtorFunc); 983 llvm::Constant *Zero = llvm::Constant::getNullValue(HandleValue->getType()); 984 llvm::Value *NEZero = DtorBuilder.CreateICmpNE(HandleValue, Zero); 985 DtorBuilder.CreateCondBr(NEZero, IfBlock, ExitBlock); 986 987 DtorBuilder.SetInsertPoint(IfBlock); 988 DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue); 989 DtorBuilder.CreateStore(Zero, GpuBinaryAddr); 990 DtorBuilder.CreateBr(ExitBlock); 991 992 DtorBuilder.SetInsertPoint(ExitBlock); 993 } else { 994 DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue); 995 } 996 DtorBuilder.CreateRetVoid(); 997 return ModuleDtorFunc; 998 } 999 1000 CGCUDARuntime *CodeGen::CreateNVCUDARuntime(CodeGenModule &CGM) { 1001 return new CGNVCUDARuntime(CGM); 1002 } 1003 1004 void CGNVCUDARuntime::internalizeDeviceSideVar( 1005 const VarDecl *D, llvm::GlobalValue::LinkageTypes &Linkage) { 1006 // For -fno-gpu-rdc, host-side shadows of external declarations of device-side 1007 // global variables become internal definitions. These have to be internal in 1008 // order to prevent name conflicts with global host variables with the same 1009 // name in a different TUs. 1010 // 1011 // For -fgpu-rdc, the shadow variables should not be internalized because 1012 // they may be accessed by different TU. 1013 if (CGM.getLangOpts().GPURelocatableDeviceCode) 1014 return; 1015 1016 // __shared__ variables are odd. Shadows do get created, but 1017 // they are not registered with the CUDA runtime, so they 1018 // can't really be used to access their device-side 1019 // counterparts. It's not clear yet whether it's nvcc's bug or 1020 // a feature, but we've got to do the same for compatibility. 1021 if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() || 1022 D->hasAttr<CUDASharedAttr>() || 1023 D->getType()->isCUDADeviceBuiltinSurfaceType() || 1024 D->getType()->isCUDADeviceBuiltinTextureType()) { 1025 Linkage = llvm::GlobalValue::InternalLinkage; 1026 } 1027 } 1028 1029 void CGNVCUDARuntime::handleVarRegistration(const VarDecl *D, 1030 llvm::GlobalVariable &GV) { 1031 if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) { 1032 // Shadow variables and their properties must be registered with CUDA 1033 // runtime. Skip Extern global variables, which will be registered in 1034 // the TU where they are defined. 1035 // 1036 // Don't register a C++17 inline variable. The local symbol can be 1037 // discarded and referencing a discarded local symbol from outside the 1038 // comdat (__cuda_register_globals) is disallowed by the ELF spec. 1039 // 1040 // HIP managed variables need to be always recorded in device and host 1041 // compilations for transformation. 1042 // 1043 // HIP managed variables and variables in CUDADeviceVarODRUsedByHost are 1044 // added to llvm.compiler-used, therefore they are safe to be registered. 1045 if ((!D->hasExternalStorage() && !D->isInline()) || 1046 CGM.getContext().CUDADeviceVarODRUsedByHost.contains(D) || 1047 D->hasAttr<HIPManagedAttr>()) { 1048 registerDeviceVar(D, GV, !D->hasDefinition(), 1049 D->hasAttr<CUDAConstantAttr>()); 1050 } 1051 } else if (D->getType()->isCUDADeviceBuiltinSurfaceType() || 1052 D->getType()->isCUDADeviceBuiltinTextureType()) { 1053 // Builtin surfaces and textures and their template arguments are 1054 // also registered with CUDA runtime. 1055 const auto *TD = cast<ClassTemplateSpecializationDecl>( 1056 D->getType()->castAs<RecordType>()->getDecl()); 1057 const TemplateArgumentList &Args = TD->getTemplateArgs(); 1058 if (TD->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>()) { 1059 assert(Args.size() == 2 && 1060 "Unexpected number of template arguments of CUDA device " 1061 "builtin surface type."); 1062 auto SurfType = Args[1].getAsIntegral(); 1063 if (!D->hasExternalStorage()) 1064 registerDeviceSurf(D, GV, !D->hasDefinition(), SurfType.getSExtValue()); 1065 } else { 1066 assert(Args.size() == 3 && 1067 "Unexpected number of template arguments of CUDA device " 1068 "builtin texture type."); 1069 auto TexType = Args[1].getAsIntegral(); 1070 auto Normalized = Args[2].getAsIntegral(); 1071 if (!D->hasExternalStorage()) 1072 registerDeviceTex(D, GV, !D->hasDefinition(), TexType.getSExtValue(), 1073 Normalized.getZExtValue()); 1074 } 1075 } 1076 } 1077 1078 // Transform managed variables to pointers to managed variables in device code. 1079 // Each use of the original managed variable is replaced by a load from the 1080 // transformed managed variable. The transformed managed variable contains 1081 // the address of managed memory which will be allocated by the runtime. 1082 void CGNVCUDARuntime::transformManagedVars() { 1083 for (auto &&Info : DeviceVars) { 1084 llvm::GlobalVariable *Var = Info.Var; 1085 if (Info.Flags.getKind() == DeviceVarFlags::Variable && 1086 Info.Flags.isManaged()) { 1087 auto *ManagedVar = new llvm::GlobalVariable( 1088 CGM.getModule(), Var->getType(), 1089 /*isConstant=*/false, Var->getLinkage(), 1090 /*Init=*/Var->isDeclaration() 1091 ? nullptr 1092 : llvm::ConstantPointerNull::get(Var->getType()), 1093 /*Name=*/"", /*InsertBefore=*/nullptr, 1094 llvm::GlobalVariable::NotThreadLocal, 1095 CGM.getContext().getTargetAddressSpace(LangAS::cuda_device)); 1096 ManagedVar->setDSOLocal(Var->isDSOLocal()); 1097 ManagedVar->setVisibility(Var->getVisibility()); 1098 ManagedVar->setExternallyInitialized(true); 1099 replaceManagedVar(Var, ManagedVar); 1100 ManagedVar->takeName(Var); 1101 Var->setName(Twine(ManagedVar->getName()) + ".managed"); 1102 // Keep managed variables even if they are not used in device code since 1103 // they need to be allocated by the runtime. 1104 if (!Var->isDeclaration()) { 1105 assert(!ManagedVar->isDeclaration()); 1106 CGM.addCompilerUsedGlobal(Var); 1107 CGM.addCompilerUsedGlobal(ManagedVar); 1108 } 1109 } 1110 } 1111 } 1112 1113 // Creates offloading entries for all the kernels and globals that must be 1114 // registered. The linker will provide a pointer to this section so we can 1115 // register the symbols with the linked device image. 1116 void CGNVCUDARuntime::createOffloadingEntries() { 1117 StringRef Section = CGM.getLangOpts().HIP ? "hip_offloading_entries" 1118 : "cuda_offloading_entries"; 1119 llvm::Module &M = CGM.getModule(); 1120 for (KernelInfo &I : EmittedKernels) 1121 llvm::offloading::emitOffloadingEntry( 1122 M, KernelHandles[I.Kernel->getName()], 1123 getDeviceSideName(cast<NamedDecl>(I.D)), /*Flags=*/0, /*Data=*/0, 1124 llvm::offloading::OffloadGlobalEntry, Section); 1125 1126 for (VarInfo &I : DeviceVars) { 1127 uint64_t VarSize = 1128 CGM.getDataLayout().getTypeAllocSize(I.Var->getValueType()); 1129 int32_t Flags = 1130 (I.Flags.isExtern() 1131 ? static_cast<int32_t>(llvm::offloading::OffloadGlobalExtern) 1132 : 0) | 1133 (I.Flags.isConstant() 1134 ? static_cast<int32_t>(llvm::offloading::OffloadGlobalConstant) 1135 : 0) | 1136 (I.Flags.isNormalized() 1137 ? static_cast<int32_t>(llvm::offloading::OffloadGlobalNormalized) 1138 : 0); 1139 if (I.Flags.getKind() == DeviceVarFlags::Variable) { 1140 llvm::offloading::emitOffloadingEntry( 1141 M, I.Var, getDeviceSideName(I.D), VarSize, 1142 (I.Flags.isManaged() ? llvm::offloading::OffloadGlobalManagedEntry 1143 : llvm::offloading::OffloadGlobalEntry) | 1144 Flags, 1145 /*Data=*/0, Section); 1146 } else if (I.Flags.getKind() == DeviceVarFlags::Surface) { 1147 llvm::offloading::emitOffloadingEntry( 1148 M, I.Var, getDeviceSideName(I.D), VarSize, 1149 llvm::offloading::OffloadGlobalSurfaceEntry | Flags, 1150 I.Flags.getSurfTexType(), Section); 1151 } else if (I.Flags.getKind() == DeviceVarFlags::Texture) { 1152 llvm::offloading::emitOffloadingEntry( 1153 M, I.Var, getDeviceSideName(I.D), VarSize, 1154 llvm::offloading::OffloadGlobalTextureEntry | Flags, 1155 I.Flags.getSurfTexType(), Section); 1156 } 1157 } 1158 } 1159 1160 // Returns module constructor to be added. 1161 llvm::Function *CGNVCUDARuntime::finalizeModule() { 1162 if (CGM.getLangOpts().CUDAIsDevice) { 1163 transformManagedVars(); 1164 1165 // Mark ODR-used device variables as compiler used to prevent it from being 1166 // eliminated by optimization. This is necessary for device variables 1167 // ODR-used by host functions. Sema correctly marks them as ODR-used no 1168 // matter whether they are ODR-used by device or host functions. 1169 // 1170 // We do not need to do this if the variable has used attribute since it 1171 // has already been added. 1172 // 1173 // Static device variables have been externalized at this point, therefore 1174 // variables with LLVM private or internal linkage need not be added. 1175 for (auto &&Info : DeviceVars) { 1176 auto Kind = Info.Flags.getKind(); 1177 if (!Info.Var->isDeclaration() && 1178 !llvm::GlobalValue::isLocalLinkage(Info.Var->getLinkage()) && 1179 (Kind == DeviceVarFlags::Variable || 1180 Kind == DeviceVarFlags::Surface || 1181 Kind == DeviceVarFlags::Texture) && 1182 Info.D->isUsed() && !Info.D->hasAttr<UsedAttr>()) { 1183 CGM.addCompilerUsedGlobal(Info.Var); 1184 } 1185 } 1186 return nullptr; 1187 } 1188 if (CGM.getLangOpts().OffloadingNewDriver && RelocatableDeviceCode) 1189 createOffloadingEntries(); 1190 else 1191 return makeModuleCtorFunction(); 1192 1193 return nullptr; 1194 } 1195 1196 llvm::GlobalValue *CGNVCUDARuntime::getKernelHandle(llvm::Function *F, 1197 GlobalDecl GD) { 1198 auto Loc = KernelHandles.find(F->getName()); 1199 if (Loc != KernelHandles.end()) { 1200 auto OldHandle = Loc->second; 1201 if (KernelStubs[OldHandle] == F) 1202 return OldHandle; 1203 1204 // We've found the function name, but F itself has changed, so we need to 1205 // update the references. 1206 if (CGM.getLangOpts().HIP) { 1207 // For HIP compilation the handle itself does not change, so we only need 1208 // to update the Stub value. 1209 KernelStubs[OldHandle] = F; 1210 return OldHandle; 1211 } 1212 // For non-HIP compilation, erase the old Stub and fall-through to creating 1213 // new entries. 1214 KernelStubs.erase(OldHandle); 1215 } 1216 1217 if (!CGM.getLangOpts().HIP) { 1218 KernelHandles[F->getName()] = F; 1219 KernelStubs[F] = F; 1220 return F; 1221 } 1222 1223 auto *Var = new llvm::GlobalVariable( 1224 TheModule, F->getType(), /*isConstant=*/true, F->getLinkage(), 1225 /*Initializer=*/nullptr, 1226 CGM.getMangledName( 1227 GD.getWithKernelReferenceKind(KernelReferenceKind::Kernel))); 1228 Var->setAlignment(CGM.getPointerAlign().getAsAlign()); 1229 Var->setDSOLocal(F->isDSOLocal()); 1230 Var->setVisibility(F->getVisibility()); 1231 auto *FD = cast<FunctionDecl>(GD.getDecl()); 1232 auto *FT = FD->getPrimaryTemplate(); 1233 if (!FT || FT->isThisDeclarationADefinition()) 1234 CGM.maybeSetTrivialComdat(*FD, *Var); 1235 KernelHandles[F->getName()] = Var; 1236 KernelStubs[Var] = F; 1237 return Var; 1238 } 1239