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.getContext().getCUIDHash(), nullptr,
844 llvm::GlobalVariable::NotThreadLocal);
845 cast<llvm::GlobalVariable>(FatBinStr)->setSection(FatbinConstantName);
846 }
847
848 FatMagic = HIPFatMagic;
849 } else {
850 if (RelocatableDeviceCode)
851 FatbinConstantName = CGM.getTriple().isMacOSX()
852 ? "__NV_CUDA,__nv_relfatbin"
853 : "__nv_relfatbin";
854 else
855 FatbinConstantName =
856 CGM.getTriple().isMacOSX() ? "__NV_CUDA,__nv_fatbin" : ".nv_fatbin";
857 // NVIDIA's cuobjdump looks for fatbins in this section.
858 FatbinSectionName =
859 CGM.getTriple().isMacOSX() ? "__NV_CUDA,__fatbin" : ".nvFatBinSegment";
860
861 ModuleIDSectionName = CGM.getTriple().isMacOSX()
862 ? "__NV_CUDA,__nv_module_id"
863 : "__nv_module_id";
864 ModuleIDPrefix = "__nv_";
865
866 // For CUDA, create a string literal containing the fat binary loaded from
867 // the given file.
868 FatBinStr = makeConstantArray(std::string(CudaGpuBinary->getBuffer()), "",
869 FatbinConstantName, 8);
870 FatMagic = CudaFatMagic;
871 }
872
873 // Create initialized wrapper structure that points to the loaded GPU binary
874 ConstantInitBuilder Builder(CGM);
875 auto Values = Builder.beginStruct(FatbinWrapperTy);
876 // Fatbin wrapper magic.
877 Values.addInt(IntTy, FatMagic);
878 // Fatbin version.
879 Values.addInt(IntTy, 1);
880 // Data.
881 Values.add(FatBinStr);
882 // Unused in fatbin v1.
883 Values.add(llvm::ConstantPointerNull::get(PtrTy));
884 llvm::GlobalVariable *FatbinWrapper = Values.finishAndCreateGlobal(
885 addUnderscoredPrefixToName("_fatbin_wrapper"), CGM.getPointerAlign(),
886 /*constant*/ true);
887 FatbinWrapper->setSection(FatbinSectionName);
888
889 // There is only one HIP fat binary per linked module, however there are
890 // multiple constructor functions. Make sure the fat binary is registered
891 // only once. The constructor functions are executed by the dynamic loader
892 // before the program gains control. The dynamic loader cannot execute the
893 // constructor functions concurrently since doing that would not guarantee
894 // thread safety of the loaded program. Therefore we can assume sequential
895 // execution of constructor functions here.
896 if (IsHIP) {
897 auto Linkage = CudaGpuBinary ? llvm::GlobalValue::InternalLinkage
898 : llvm::GlobalValue::ExternalLinkage;
899 llvm::BasicBlock *IfBlock =
900 llvm::BasicBlock::Create(Context, "if", ModuleCtorFunc);
901 llvm::BasicBlock *ExitBlock =
902 llvm::BasicBlock::Create(Context, "exit", ModuleCtorFunc);
903 // The name, size, and initialization pattern of this variable is part
904 // of HIP ABI.
905 GpuBinaryHandle = new llvm::GlobalVariable(
906 TheModule, PtrTy, /*isConstant=*/false, Linkage,
907 /*Initializer=*/
908 CudaGpuBinary ? llvm::ConstantPointerNull::get(PtrTy) : nullptr,
909 CudaGpuBinary
910 ? "__hip_gpubin_handle"
911 : "__hip_gpubin_handle_" + CGM.getContext().getCUIDHash());
912 GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign());
913 // Prevent the weak symbol in different shared libraries being merged.
914 if (Linkage != llvm::GlobalValue::InternalLinkage)
915 GpuBinaryHandle->setVisibility(llvm::GlobalValue::HiddenVisibility);
916 Address GpuBinaryAddr(
917 GpuBinaryHandle, PtrTy,
918 CharUnits::fromQuantity(GpuBinaryHandle->getAlignment()));
919 {
920 auto *HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr);
921 llvm::Constant *Zero =
922 llvm::Constant::getNullValue(HandleValue->getType());
923 llvm::Value *EQZero = CtorBuilder.CreateICmpEQ(HandleValue, Zero);
924 CtorBuilder.CreateCondBr(EQZero, IfBlock, ExitBlock);
925 }
926 {
927 CtorBuilder.SetInsertPoint(IfBlock);
928 // GpuBinaryHandle = __hipRegisterFatBinary(&FatbinWrapper);
929 llvm::CallInst *RegisterFatbinCall =
930 CtorBuilder.CreateCall(RegisterFatbinFunc, FatbinWrapper);
931 CtorBuilder.CreateStore(RegisterFatbinCall, GpuBinaryAddr);
932 CtorBuilder.CreateBr(ExitBlock);
933 }
934 {
935 CtorBuilder.SetInsertPoint(ExitBlock);
936 // Call __hip_register_globals(GpuBinaryHandle);
937 if (RegisterGlobalsFunc) {
938 auto *HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr);
939 CtorBuilder.CreateCall(RegisterGlobalsFunc, HandleValue);
940 }
941 }
942 } else if (!RelocatableDeviceCode) {
943 // Register binary with CUDA runtime. This is substantially different in
944 // default mode vs. separate compilation!
945 // GpuBinaryHandle = __cudaRegisterFatBinary(&FatbinWrapper);
946 llvm::CallInst *RegisterFatbinCall =
947 CtorBuilder.CreateCall(RegisterFatbinFunc, FatbinWrapper);
948 GpuBinaryHandle = new llvm::GlobalVariable(
949 TheModule, PtrTy, false, llvm::GlobalValue::InternalLinkage,
950 llvm::ConstantPointerNull::get(PtrTy), "__cuda_gpubin_handle");
951 GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign());
952 CtorBuilder.CreateAlignedStore(RegisterFatbinCall, GpuBinaryHandle,
953 CGM.getPointerAlign());
954
955 // Call __cuda_register_globals(GpuBinaryHandle);
956 if (RegisterGlobalsFunc)
957 CtorBuilder.CreateCall(RegisterGlobalsFunc, RegisterFatbinCall);
958
959 // Call __cudaRegisterFatBinaryEnd(Handle) if this CUDA version needs it.
961 CudaFeature::CUDA_USES_FATBIN_REGISTER_END)) {
962 // void __cudaRegisterFatBinaryEnd(void **);
963 llvm::FunctionCallee RegisterFatbinEndFunc = CGM.CreateRuntimeFunction(
964 llvm::FunctionType::get(VoidTy, PtrTy, false),
965 "__cudaRegisterFatBinaryEnd");
966 CtorBuilder.CreateCall(RegisterFatbinEndFunc, RegisterFatbinCall);
967 }
968 } else {
969 // Generate a unique module ID.
970 SmallString<64> ModuleID;
971 llvm::raw_svector_ostream OS(ModuleID);
972 OS << ModuleIDPrefix << llvm::format("%" PRIx64, FatbinWrapper->getGUID());
973 llvm::Constant *ModuleIDConstant = makeConstantArray(
974 std::string(ModuleID), "", ModuleIDSectionName, 32, /*AddNull=*/true);
975
976 // Create an alias for the FatbinWrapper that nvcc will look for.
977 llvm::GlobalAlias::create(llvm::GlobalValue::ExternalLinkage,
978 Twine("__fatbinwrap") + ModuleID, FatbinWrapper);
979
980 // void __cudaRegisterLinkedBinary%ModuleID%(void (*)(void *), void *,
981 // void *, void (*)(void **))
982 SmallString<128> RegisterLinkedBinaryName("__cudaRegisterLinkedBinary");
983 RegisterLinkedBinaryName += ModuleID;
984 llvm::FunctionCallee RegisterLinkedBinaryFunc = CGM.CreateRuntimeFunction(
985 getRegisterLinkedBinaryFnTy(), RegisterLinkedBinaryName);
986
987 assert(RegisterGlobalsFunc && "Expecting at least dummy function!");
988 llvm::Value *Args[] = {RegisterGlobalsFunc, FatbinWrapper, ModuleIDConstant,
989 makeDummyFunction(getCallbackFnTy())};
990 CtorBuilder.CreateCall(RegisterLinkedBinaryFunc, Args);
991 }
992
993 // Create destructor and register it with atexit() the way NVCC does it. Doing
994 // it during regular destructor phase worked in CUDA before 9.2 but results in
995 // double-free in 9.2.
996 if (llvm::Function *CleanupFn = makeModuleDtorFunction()) {
997 // extern "C" int atexit(void (*f)(void));
998 llvm::FunctionType *AtExitTy =
999 llvm::FunctionType::get(IntTy, CleanupFn->getType(), false);
1000 llvm::FunctionCallee AtExitFunc =
1001 CGM.CreateRuntimeFunction(AtExitTy, "atexit", llvm::AttributeList(),
1002 /*Local=*/true);
1003 CtorBuilder.CreateCall(AtExitFunc, CleanupFn);
1004 }
1005
1006 CtorBuilder.CreateRetVoid();
1007 return ModuleCtorFunc;
1008}
1009
1010/// Creates a global destructor function that unregisters the GPU code blob
1011/// registered by constructor.
1012///
1013/// For CUDA:
1014/// \code
1015/// void __cuda_module_dtor() {
1016/// __cudaUnregisterFatBinary(Handle);
1017/// }
1018/// \endcode
1019///
1020/// For HIP:
1021/// \code
1022/// void __hip_module_dtor() {
1023/// if (__hip_gpubin_handle) {
1024/// __hipUnregisterFatBinary(__hip_gpubin_handle);
1025/// __hip_gpubin_handle = 0;
1026/// }
1027/// }
1028/// \endcode
1029llvm::Function *CGNVCUDARuntime::makeModuleDtorFunction() {
1030 // No need for destructor if we don't have a handle to unregister.
1031 if (!GpuBinaryHandle)
1032 return nullptr;
1033
1034 // void __cudaUnregisterFatBinary(void ** handle);
1035 llvm::FunctionCallee UnregisterFatbinFunc = CGM.CreateRuntimeFunction(
1036 llvm::FunctionType::get(VoidTy, PtrTy, false),
1037 addUnderscoredPrefixToName("UnregisterFatBinary"));
1038
1039 llvm::Function *ModuleDtorFunc = llvm::Function::Create(
1040 llvm::FunctionType::get(VoidTy, false),
1041 llvm::GlobalValue::InternalLinkage,
1042 addUnderscoredPrefixToName("_module_dtor"), &TheModule);
1043
1044 llvm::BasicBlock *DtorEntryBB =
1045 llvm::BasicBlock::Create(Context, "entry", ModuleDtorFunc);
1046 CGBuilderTy DtorBuilder(CGM, Context);
1047 DtorBuilder.SetInsertPoint(DtorEntryBB);
1048
1049 Address GpuBinaryAddr(
1050 GpuBinaryHandle, GpuBinaryHandle->getValueType(),
1051 CharUnits::fromQuantity(GpuBinaryHandle->getAlignment()));
1052 auto *HandleValue = DtorBuilder.CreateLoad(GpuBinaryAddr);
1053 // There is only one HIP fat binary per linked module, however there are
1054 // multiple destructor functions. Make sure the fat binary is unregistered
1055 // only once.
1056 if (CGM.getLangOpts().HIP) {
1057 llvm::BasicBlock *IfBlock =
1058 llvm::BasicBlock::Create(Context, "if", ModuleDtorFunc);
1059 llvm::BasicBlock *ExitBlock =
1060 llvm::BasicBlock::Create(Context, "exit", ModuleDtorFunc);
1061 llvm::Constant *Zero = llvm::Constant::getNullValue(HandleValue->getType());
1062 llvm::Value *NEZero = DtorBuilder.CreateICmpNE(HandleValue, Zero);
1063 DtorBuilder.CreateCondBr(NEZero, IfBlock, ExitBlock);
1064
1065 DtorBuilder.SetInsertPoint(IfBlock);
1066 DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue);
1067 DtorBuilder.CreateStore(Zero, GpuBinaryAddr);
1068 DtorBuilder.CreateBr(ExitBlock);
1069
1070 DtorBuilder.SetInsertPoint(ExitBlock);
1071 } else {
1072 DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue);
1073 }
1074 DtorBuilder.CreateRetVoid();
1075 return ModuleDtorFunc;
1076}
1077
1079 return new CGNVCUDARuntime(CGM);
1080}
1081
1082void CGNVCUDARuntime::internalizeDeviceSideVar(
1083 const VarDecl *D, llvm::GlobalValue::LinkageTypes &Linkage) {
1084 // For -fno-gpu-rdc, host-side shadows of external declarations of device-side
1085 // global variables become internal definitions. These have to be internal in
1086 // order to prevent name conflicts with global host variables with the same
1087 // name in a different TUs.
1088 //
1089 // For -fgpu-rdc, the shadow variables should not be internalized because
1090 // they may be accessed by different TU.
1091 if (CGM.getLangOpts().GPURelocatableDeviceCode)
1092 return;
1093
1094 // __shared__ variables are odd. Shadows do get created, but
1095 // they are not registered with the CUDA runtime, so they
1096 // can't really be used to access their device-side
1097 // counterparts. It's not clear yet whether it's nvcc's bug or
1098 // a feature, but we've got to do the same for compatibility.
1099 if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
1100 D->hasAttr<CUDASharedAttr>() ||
1101 D->getType()->isCUDADeviceBuiltinSurfaceType() ||
1102 D->getType()->isCUDADeviceBuiltinTextureType()) {
1103 Linkage = llvm::GlobalValue::InternalLinkage;
1104 }
1105}
1106
1107void CGNVCUDARuntime::handleVarRegistration(const VarDecl *D,
1108 llvm::GlobalVariable &GV) {
1109 if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
1110 // Shadow variables and their properties must be registered with CUDA
1111 // runtime. Skip Extern global variables, which will be registered in
1112 // the TU where they are defined.
1113 //
1114 // Don't register a C++17 inline variable. The local symbol can be
1115 // discarded and referencing a discarded local symbol from outside the
1116 // comdat (__cuda_register_globals) is disallowed by the ELF spec.
1117 //
1118 // HIP managed variables need to be always recorded in device and host
1119 // compilations for transformation.
1120 //
1121 // HIP managed variables and variables in CUDADeviceVarODRUsedByHost are
1122 // added to llvm.compiler-used, therefore they are safe to be registered.
1123 if ((!D->hasExternalStorage() && !D->isInline()) ||
1124 CGM.getContext().CUDADeviceVarODRUsedByHost.contains(D) ||
1125 D->hasAttr<HIPManagedAttr>()) {
1126 registerDeviceVar(D, GV, !D->hasDefinition(),
1127 D->hasAttr<CUDAConstantAttr>());
1128 }
1129 } else if (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
1130 D->getType()->isCUDADeviceBuiltinTextureType()) {
1131 // Builtin surfaces and textures and their template arguments are
1132 // also registered with CUDA runtime.
1133 const auto *TD = cast<ClassTemplateSpecializationDecl>(
1134 D->getType()->castAs<RecordType>()->getDecl());
1135 const TemplateArgumentList &Args = TD->getTemplateArgs();
1136 if (TD->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>()) {
1137 assert(Args.size() == 2 &&
1138 "Unexpected number of template arguments of CUDA device "
1139 "builtin surface type.");
1140 auto SurfType = Args[1].getAsIntegral();
1141 if (!D->hasExternalStorage())
1142 registerDeviceSurf(D, GV, !D->hasDefinition(), SurfType.getSExtValue());
1143 } else {
1144 assert(Args.size() == 3 &&
1145 "Unexpected number of template arguments of CUDA device "
1146 "builtin texture type.");
1147 auto TexType = Args[1].getAsIntegral();
1148 auto Normalized = Args[2].getAsIntegral();
1149 if (!D->hasExternalStorage())
1150 registerDeviceTex(D, GV, !D->hasDefinition(), TexType.getSExtValue(),
1151 Normalized.getZExtValue());
1152 }
1153 }
1154}
1155
1156// Transform managed variables to pointers to managed variables in device code.
1157// Each use of the original managed variable is replaced by a load from the
1158// transformed managed variable. The transformed managed variable contains
1159// the address of managed memory which will be allocated by the runtime.
1160void CGNVCUDARuntime::transformManagedVars() {
1161 for (auto &&Info : DeviceVars) {
1162 llvm::GlobalVariable *Var = Info.Var;
1163 if (Info.Flags.getKind() == DeviceVarFlags::Variable &&
1164 Info.Flags.isManaged()) {
1165 auto *ManagedVar = new llvm::GlobalVariable(
1166 CGM.getModule(), Var->getType(),
1167 /*isConstant=*/false, Var->getLinkage(),
1168 /*Init=*/Var->isDeclaration()
1169 ? nullptr
1170 : llvm::ConstantPointerNull::get(Var->getType()),
1171 /*Name=*/"", /*InsertBefore=*/nullptr,
1172 llvm::GlobalVariable::NotThreadLocal,
1173 CGM.getContext().getTargetAddressSpace(CGM.getLangOpts().CUDAIsDevice
1174 ? LangAS::cuda_device
1175 : LangAS::Default));
1176 ManagedVar->setDSOLocal(Var->isDSOLocal());
1177 ManagedVar->setVisibility(Var->getVisibility());
1178 ManagedVar->setExternallyInitialized(true);
1179 replaceManagedVar(Var, ManagedVar);
1180 ManagedVar->takeName(Var);
1181 Var->setName(Twine(ManagedVar->getName()) + ".managed");
1182 // Keep managed variables even if they are not used in device code since
1183 // they need to be allocated by the runtime.
1184 if (CGM.getLangOpts().CUDAIsDevice && !Var->isDeclaration()) {
1185 assert(!ManagedVar->isDeclaration());
1186 CGM.addCompilerUsedGlobal(Var);
1187 CGM.addCompilerUsedGlobal(ManagedVar);
1188 }
1189 }
1190 }
1191}
1192
1193// Creates offloading entries for all the kernels and globals that must be
1194// registered. The linker will provide a pointer to this section so we can
1195// register the symbols with the linked device image.
1196void CGNVCUDARuntime::createOffloadingEntries() {
1198 StringRef Section = (SectionPrefix + "_offloading_entries").toStringRef(Out);
1199
1200 llvm::Module &M = CGM.getModule();
1201 for (KernelInfo &I : EmittedKernels)
1202 llvm::offloading::emitOffloadingEntry(
1203 M, KernelHandles[I.Kernel->getName()],
1204 getDeviceSideName(cast<NamedDecl>(I.D)), /*Flags=*/0, /*Data=*/0,
1205 llvm::offloading::OffloadGlobalEntry, Section);
1206
1207 for (VarInfo &I : DeviceVars) {
1208 uint64_t VarSize =
1209 CGM.getDataLayout().getTypeAllocSize(I.Var->getValueType());
1210 int32_t Flags =
1211 (I.Flags.isExtern()
1212 ? static_cast<int32_t>(llvm::offloading::OffloadGlobalExtern)
1213 : 0) |
1214 (I.Flags.isConstant()
1215 ? static_cast<int32_t>(llvm::offloading::OffloadGlobalConstant)
1216 : 0) |
1217 (I.Flags.isNormalized()
1218 ? static_cast<int32_t>(llvm::offloading::OffloadGlobalNormalized)
1219 : 0);
1220 if (I.Flags.getKind() == DeviceVarFlags::Variable) {
1221 llvm::offloading::emitOffloadingEntry(
1222 M, I.Var, getDeviceSideName(I.D), VarSize,
1223 (I.Flags.isManaged() ? llvm::offloading::OffloadGlobalManagedEntry
1224 : llvm::offloading::OffloadGlobalEntry) |
1225 Flags,
1226 /*Data=*/0, Section);
1227 } else if (I.Flags.getKind() == DeviceVarFlags::Surface) {
1228 llvm::offloading::emitOffloadingEntry(
1229 M, I.Var, getDeviceSideName(I.D), VarSize,
1230 llvm::offloading::OffloadGlobalSurfaceEntry | Flags,
1231 I.Flags.getSurfTexType(), Section);
1232 } else if (I.Flags.getKind() == DeviceVarFlags::Texture) {
1233 llvm::offloading::emitOffloadingEntry(
1234 M, I.Var, getDeviceSideName(I.D), VarSize,
1235 llvm::offloading::OffloadGlobalTextureEntry | Flags,
1236 I.Flags.getSurfTexType(), Section);
1237 }
1238 }
1239}
1240
1241// Returns module constructor to be added.
1242llvm::Function *CGNVCUDARuntime::finalizeModule() {
1243 transformManagedVars();
1244 if (CGM.getLangOpts().CUDAIsDevice) {
1245 // Mark ODR-used device variables as compiler used to prevent it from being
1246 // eliminated by optimization. This is necessary for device variables
1247 // ODR-used by host functions. Sema correctly marks them as ODR-used no
1248 // matter whether they are ODR-used by device or host functions.
1249 //
1250 // We do not need to do this if the variable has used attribute since it
1251 // has already been added.
1252 //
1253 // Static device variables have been externalized at this point, therefore
1254 // variables with LLVM private or internal linkage need not be added.
1255 for (auto &&Info : DeviceVars) {
1256 auto Kind = Info.Flags.getKind();
1257 if (!Info.Var->isDeclaration() &&
1258 !llvm::GlobalValue::isLocalLinkage(Info.Var->getLinkage()) &&
1259 (Kind == DeviceVarFlags::Variable ||
1260 Kind == DeviceVarFlags::Surface ||
1261 Kind == DeviceVarFlags::Texture) &&
1262 Info.D->isUsed() && !Info.D->hasAttr<UsedAttr>()) {
1263 CGM.addCompilerUsedGlobal(Info.Var);
1264 }
1265 }
1266 return nullptr;
1267 }
1268 if (CGM.getLangOpts().OffloadViaLLVM ||
1269 (CGM.getLangOpts().OffloadingNewDriver && RelocatableDeviceCode))
1270 createOffloadingEntries();
1271 else
1272 return makeModuleCtorFunction();
1273
1274 return nullptr;
1275}
1276
1277llvm::GlobalValue *CGNVCUDARuntime::getKernelHandle(llvm::Function *F,
1278 GlobalDecl GD) {
1279 auto Loc = KernelHandles.find(F->getName());
1280 if (Loc != KernelHandles.end()) {
1281 auto OldHandle = Loc->second;
1282 if (KernelStubs[OldHandle] == F)
1283 return OldHandle;
1284
1285 // We've found the function name, but F itself has changed, so we need to
1286 // update the references.
1287 if (CGM.getLangOpts().HIP) {
1288 // For HIP compilation the handle itself does not change, so we only need
1289 // to update the Stub value.
1290 KernelStubs[OldHandle] = F;
1291 return OldHandle;
1292 }
1293 // For non-HIP compilation, erase the old Stub and fall-through to creating
1294 // new entries.
1295 KernelStubs.erase(OldHandle);
1296 }
1297
1298 if (!CGM.getLangOpts().HIP) {
1299 KernelHandles[F->getName()] = F;
1300 KernelStubs[F] = F;
1301 return F;
1302 }
1303
1304 auto *Var = new llvm::GlobalVariable(
1305 TheModule, F->getType(), /*isConstant=*/true, F->getLinkage(),
1306 /*Initializer=*/nullptr,
1307 CGM.getMangledName(
1308 GD.getWithKernelReferenceKind(KernelReferenceKind::Kernel)));
1309 Var->setAlignment(CGM.getPointerAlign().getAsAlign());
1310 Var->setDSOLocal(F->isDSOLocal());
1311 Var->setVisibility(F->getVisibility());
1312 auto *FD = cast<FunctionDecl>(GD.getDecl());
1313 auto *FT = FD->getPrimaryTemplate();
1314 if (!FT || FT->isThisDeclarationADefinition())
1315 CGM.maybeSetTrivialComdat(*FD, *Var);
1316 KernelHandles[F->getName()] = Var;
1317 KernelStubs[Var] = F;
1318 return Var;
1319}
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:758
TranslationUnitDecl * getTranslationUnitDecl() const
Definition: ASTContext.h:1101
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:660
const TargetInfo * getAuxTargetInfo() const
Definition: ASTContext.h:780
llvm::DenseSet< const VarDecl * > CUDADeviceVarODRUsedByHost
Keep track of CUDA/HIP device-side variables ODR-used by host code.
Definition: ASTContext.h:1195
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:779
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:135
llvm::StoreInst * CreateAlignedStore(llvm::Value *Val, llvm::Value *Addr, CharUnits Align, bool IsVolatile=false)
Definition: CGBuilder.h:142
llvm::StoreInst * CreateDefaultAlignedStore(llvm::Value *Val, llvm::Value *Addr, bool IsVolatile=false)
Definition: CGBuilder.h:150
Address CreateStructGEP(Address Addr, unsigned Index, const llvm::Twine &Name="")
Definition: CGBuilder.h:218
llvm::LoadInst * CreateLoad(Address Addr, const llvm::Twine &Name="")
Definition: CGBuilder.h:107
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:298
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...
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,...
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,...
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:463
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:368
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:372
DeclContext - This is used only as base class of specific decl types that can act as declaration cont...
Definition: DeclBase.h:1436
lookup_result lookup(DeclarationName Name) const
lookup - Find the declarations (if any) with the given Name in this context.
Definition: DeclBase.cpp:1852
Decl - This represents one declaration (or definition), e.g.
Definition: DeclBase.h:86
SourceLocation getLocation() const
Definition: DeclBase.h:446
bool hasAttr() const
Definition: DeclBase.h:584
DiagnosticBuilder Report(SourceLocation Loc, unsigned DiagID)
Issue the message to the client.
Definition: Diagnostic.h:1547
Represents a function declaration or definition.
Definition: Decl.h:1932
const ParmVarDecl * getParamDecl(unsigned i) const
Definition: Decl.h:2669
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.
GPUDefaultStreamKind GPUDefaultStream
The default stream kind used for HIP kernel launching.
Definition: LangOptions.h:556
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:249
IdentifierInfo * getIdentifier() const
Get the identifier that names this declaration, if there is one.
Definition: Decl.h:270
Represents a parameter to a function.
Definition: Decl.h:1722
A (possibly-)qualified type.
Definition: Type.h:941
QualType getCanonicalType() const
Definition: Type.h:7802
A helper class that allows the use of isa/cast/dyncast to detect TagType objects of structs/unions/cl...
Definition: Type.h:5965
RecordDecl * getDecl() const
Definition: Type.h:5975
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:1327
const llvm::VersionTuple & getSDKVersion() const
Definition: TargetInfo.h:1795
A template argument list.
Definition: DeclTemplate.h:244
unsigned size() const
Retrieve the number of template arguments in this template argument list.
Definition: DeclTemplate.h:280
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:1829
QualType getType() const
Definition: Decl.h:678
QualType getType() const
Definition: Value.cpp:234
Represents a variable declaration or definition.
Definition: Decl.h:879
CGCUDARuntime * CreateNVCUDARuntime(CodeGenModule &CGM)
Creates an instance of a CUDA runtime class.
Definition: CGCUDANV.cpp:1078
bool Zero(InterpState &S, CodePtr OpPC)
Definition: Interp.h:2237
@ VFS
Remove unused -ivfsoverlay arguments.
The JSON file list parser is used to communicate input to InstallAPI.
CudaVersion ToCudaVersion(llvm::VersionTuple)
Definition: Cuda.cpp:67
bool CudaFeatureEnabled(llvm::VersionTuple, CudaFeature)
Definition: Cuda.cpp:251
Linkage
Describes the different kinds of linkage (C++ [basic.link], C99 6.2.2) that an entity may have.
Definition: Linkage.h:24
unsigned long uint64_t
llvm::IntegerType * Int8Ty
i8, i16, i32, and i64
llvm::IntegerType * IntTy
int