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