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