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"
21#include "clang/CodeGen/CodeGenABITypes.h"
22#include "clang/CodeGen/ConstantInitBuilder.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/GlobalValue.h"
29#include "llvm/IR/ReplaceConstant.h"
30#include "llvm/ProfileData/InstrProf.h"
31#include "llvm/Support/Format.h"
32#include "llvm/Support/MD5.h"
33#include "llvm/Support/VirtualFileSystem.h"
34#include "llvm/Transforms/Utils/ModuleUtils.h"
35
36using namespace clang;
37using namespace CodeGen;
38
39namespace {
40constexpr unsigned CudaFatMagic = 0x466243b1;
41constexpr unsigned HIPFatMagic = 0x48495046; // "HIPF"
42
43class CGNVCUDARuntime : public CGCUDARuntime {
44
45 /// The prefix used for function calls and section names (CUDA, HIP, LLVM)
46 StringRef Prefix;
47
48private:
49 llvm::IntegerType *IntTy, *SizeTy;
50 llvm::Type *VoidTy;
51 llvm::PointerType *PtrTy;
52
53 /// Convenience reference to LLVM Context
54 llvm::LLVMContext &Context;
55 /// Convenience reference to the current module
56 llvm::Module &TheModule;
57 /// Keeps track of kernel launch stubs and handles emitted in this module
58 struct KernelInfo {
59 llvm::Function *Kernel; // stub function to help launch kernel
60 const Decl *D;
61 };
62 llvm::SmallVector<KernelInfo, 16> EmittedKernels;
63 // Map a kernel mangled name to a symbol for identifying kernel in host code
64 // For CUDA, the symbol for identifying the kernel is the same as the device
65 // stub function. For HIP, they are different.
66 llvm::DenseMap<StringRef, llvm::GlobalValue *> KernelHandles;
67 // Map a kernel handle to the kernel stub.
68 llvm::DenseMap<llvm::GlobalValue *, llvm::Function *> KernelStubs;
69 struct VarInfo {
70 llvm::GlobalVariable *Var;
71 const VarDecl *D;
72 DeviceVarFlags Flags;
73 };
74 llvm::SmallVector<VarInfo, 16> DeviceVars;
75 /// Keeps track of variable containing handle of GPU binary. Populated by
76 /// ModuleCtorFunction() and used to create corresponding cleanup calls in
77 /// ModuleDtorFunction()
78 llvm::GlobalVariable *GpuBinaryHandle = nullptr;
79 /// Host-side shadow for the per-TU __llvm_profile_sections_<CUID> global,
80 /// emitted only for HIP host compiles when PGO is on. Registered via
81 /// __hipRegisterVar (non-RDC) or an offloading entry (RDC) so the runtime
82 /// can locate the device-side table by name.
83 llvm::GlobalVariable *OffloadProfShadow = nullptr;
84 struct OffloadProfSectionShadowInfo {
85 llvm::GlobalVariable *Shadow;
86 std::string DeviceName;
87 };
88 llvm::SmallVector<OffloadProfSectionShadowInfo, 16> OffloadProfSectionShadows;
89 /// Whether we generate relocatable device code.
90 bool RelocatableDeviceCode;
91 /// Mangle context for device.
92 std::unique_ptr<MangleContext> DeviceMC;
93
94 llvm::FunctionCallee getSetupArgumentFn() const;
95 llvm::FunctionCallee getLaunchFn() const;
96
97 llvm::FunctionType *getRegisterGlobalsFnTy() const;
98 llvm::FunctionType *getCallbackFnTy() const;
99 llvm::FunctionType *getRegisterLinkedBinaryFnTy() const;
100 std::string addPrefixToName(StringRef FuncName) const;
101 std::string addUnderscoredPrefixToName(StringRef FuncName) const;
102
103 /// Creates a function to register all kernel stubs generated in this module.
104 llvm::Function *makeRegisterGlobalsFn();
105
106 /// Helper function that generates a constant string and returns a pointer to
107 /// the start of the string. The result of this function can be used anywhere
108 /// where the C code specifies const char*.
109 llvm::Constant *makeConstantString(const std::string &Str,
110 const std::string &Name = "") {
111 return CGM.GetAddrOfConstantCString(Str, GlobalName: Name).getPointer();
112 }
113
114 /// Helper function which generates an initialized constant array from Str,
115 /// and optionally sets section name and alignment. AddNull specifies whether
116 /// the array should nave NUL termination.
117 llvm::Constant *makeConstantArray(StringRef Str,
118 StringRef Name = "",
119 StringRef SectionName = "",
120 unsigned Alignment = 0,
121 bool AddNull = false) {
122 llvm::Constant *Value =
123 llvm::ConstantDataArray::getString(Context, Initializer: Str, AddNull);
124 auto *GV = new llvm::GlobalVariable(
125 TheModule, Value->getType(), /*isConstant=*/true,
126 llvm::GlobalValue::PrivateLinkage, Value, Name);
127 if (!SectionName.empty()) {
128 GV->setSection(SectionName);
129 // Mark the address as used which make sure that this section isn't
130 // merged and we will really have it in the object file.
131 GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::None);
132 }
133 if (Alignment)
134 GV->setAlignment(llvm::Align(Alignment));
135 return GV;
136 }
137
138 /// Helper function that generates an empty dummy function returning void.
139 llvm::Function *makeDummyFunction(llvm::FunctionType *FnTy) {
140 assert(FnTy->getReturnType()->isVoidTy() &&
141 "Can only generate dummy functions returning void!");
142 llvm::Function *DummyFunc = llvm::Function::Create(
143 Ty: FnTy, Linkage: llvm::GlobalValue::InternalLinkage, N: "dummy", M: &TheModule);
144
145 llvm::BasicBlock *DummyBlock =
146 llvm::BasicBlock::Create(Context, Name: "", Parent: DummyFunc);
147 CGBuilderTy FuncBuilder(CGM, Context);
148 FuncBuilder.SetInsertPoint(DummyBlock);
149 FuncBuilder.CreateRetVoid();
150
151 return DummyFunc;
152 }
153
154 Address prepareKernelArgs(CodeGenFunction &CGF, FunctionArgList &Args);
155 Address prepareKernelArgsLLVMOffload(CodeGenFunction &CGF,
156 FunctionArgList &Args);
157 void emitDeviceStubBodyLegacy(CodeGenFunction &CGF, FunctionArgList &Args);
158 void emitDeviceStubBodyNew(CodeGenFunction &CGF, FunctionArgList &Args);
159 std::string getDeviceSideName(const NamedDecl *ND) override;
160
161 void registerDeviceVar(const VarDecl *VD, llvm::GlobalVariable &Var,
162 bool Extern, bool Constant) {
163 DeviceVars.push_back(Elt: {.Var: &Var,
164 .D: VD,
165 .Flags: {DeviceVarFlags::Variable, Extern, Constant,
166 VD->hasAttr<HIPManagedAttr>(),
167 /*Normalized*/ false, 0}});
168 }
169 void registerDeviceSurf(const VarDecl *VD, llvm::GlobalVariable &Var,
170 bool Extern, int Type) {
171 DeviceVars.push_back(Elt: {.Var: &Var,
172 .D: VD,
173 .Flags: {DeviceVarFlags::Surface, Extern, /*Constant*/ false,
174 /*Managed*/ false,
175 /*Normalized*/ false, Type}});
176 }
177 void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var,
178 bool Extern, int Type, bool Normalized) {
179 DeviceVars.push_back(Elt: {.Var: &Var,
180 .D: VD,
181 .Flags: {DeviceVarFlags::Texture, Extern, /*Constant*/ false,
182 /*Managed*/ false, Normalized, Type}});
183 }
184
185 /// Creates module constructor function
186 llvm::Function *makeModuleCtorFunction();
187 /// Creates module destructor function
188 llvm::Function *makeModuleDtorFunction();
189 /// Transform managed variables for device compilation.
190 void transformManagedVars();
191 /// Create offloading entries to register globals in RDC mode.
192 void createOffloadingEntries();
193 /// For HIP+PGO, emit the per-TU __llvm_profile_sections_<CUID> global.
194 /// On the device side, InstrProfiling emits the populated section-bounds
195 /// table only when the TU has real profile data. On the host side it is a
196 /// placeholder void* shadow stored in
197 /// OffloadProfShadow, registered later by makeRegisterGlobalsFn (non-RDC)
198 /// or createOffloadingEntries (RDC) so the runtime can locate the
199 /// device-side table by name.
200 void emitOffloadProfilingSections();
201
202public:
203 CGNVCUDARuntime(CodeGenModule &CGM);
204
205 llvm::GlobalValue *getKernelHandle(llvm::Function *F, GlobalDecl GD) override;
206 llvm::Function *getKernelStub(llvm::GlobalValue *Handle) override {
207 auto Loc = KernelStubs.find(Val: Handle);
208 assert(Loc != KernelStubs.end());
209 return Loc->second;
210 }
211 void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) override;
212 void handleVarRegistration(const VarDecl *VD,
213 llvm::GlobalVariable &Var) override;
214 void
215 internalizeDeviceSideVar(const VarDecl *D,
216 llvm::GlobalValue::LinkageTypes &Linkage) override;
217
218 llvm::Function *finalizeModule() override;
219};
220
221} // end anonymous namespace
222
223std::string CGNVCUDARuntime::addPrefixToName(StringRef FuncName) const {
224 return (Prefix + FuncName).str();
225}
226std::string
227CGNVCUDARuntime::addUnderscoredPrefixToName(StringRef FuncName) const {
228 return ("__" + Prefix + FuncName).str();
229}
230
231static std::unique_ptr<MangleContext> InitDeviceMC(CodeGenModule &CGM) {
232 // If the host and device have different C++ ABIs, mark it as the device
233 // mangle context so that the mangling needs to retrieve the additional
234 // device lambda mangling number instead of the regular host one.
235 if (CGM.getContext().getAuxTargetInfo() &&
236 CGM.getContext().getTargetInfo().getCXXABI().isMicrosoft() &&
237 CGM.getContext().getAuxTargetInfo()->getCXXABI().isItaniumFamily()) {
238 return std::unique_ptr<MangleContext>(
239 CGM.getContext().createDeviceMangleContext(
240 T: *CGM.getContext().getAuxTargetInfo()));
241 }
242
243 return std::unique_ptr<MangleContext>(CGM.getContext().createMangleContext(
244 T: CGM.getContext().getAuxTargetInfo()));
245}
246
247CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
248 : CGCUDARuntime(CGM), Context(CGM.getLLVMContext()),
249 TheModule(CGM.getModule()),
250 RelocatableDeviceCode(CGM.getLangOpts().GPURelocatableDeviceCode),
251 DeviceMC(InitDeviceMC(CGM)) {
252 IntTy = CGM.IntTy;
253 SizeTy = CGM.SizeTy;
254 VoidTy = CGM.VoidTy;
255 PtrTy = CGM.DefaultPtrTy;
256
257 if (CGM.getLangOpts().OffloadViaLLVM)
258 Prefix = "llvm";
259 else if (CGM.getLangOpts().HIP)
260 Prefix = "hip";
261 else
262 Prefix = "cuda";
263}
264
265llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const {
266 // cudaError_t cudaSetupArgument(void *, size_t, size_t)
267 llvm::Type *Params[] = {PtrTy, SizeTy, SizeTy};
268 return CGM.CreateRuntimeFunction(
269 Ty: llvm::FunctionType::get(Result: IntTy, Params, isVarArg: false),
270 Name: addPrefixToName(FuncName: "SetupArgument"));
271}
272
273llvm::FunctionCallee CGNVCUDARuntime::getLaunchFn() const {
274 if (CGM.getLangOpts().HIP) {
275 // hipError_t hipLaunchByPtr(char *);
276 return CGM.CreateRuntimeFunction(
277 Ty: llvm::FunctionType::get(Result: IntTy, Params: PtrTy, isVarArg: false), Name: "hipLaunchByPtr");
278 }
279 // cudaError_t cudaLaunch(char *);
280 return CGM.CreateRuntimeFunction(Ty: llvm::FunctionType::get(Result: IntTy, Params: PtrTy, isVarArg: false),
281 Name: "cudaLaunch");
282}
283
284llvm::FunctionType *CGNVCUDARuntime::getRegisterGlobalsFnTy() const {
285 return llvm::FunctionType::get(Result: VoidTy, Params: PtrTy, isVarArg: false);
286}
287
288llvm::FunctionType *CGNVCUDARuntime::getCallbackFnTy() const {
289 return llvm::FunctionType::get(Result: VoidTy, Params: PtrTy, isVarArg: false);
290}
291
292llvm::FunctionType *CGNVCUDARuntime::getRegisterLinkedBinaryFnTy() const {
293 llvm::Type *Params[] = {llvm::PointerType::getUnqual(C&: Context), PtrTy, PtrTy,
294 llvm::PointerType::getUnqual(C&: Context)};
295 return llvm::FunctionType::get(Result: VoidTy, Params, isVarArg: false);
296}
297
298std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) {
299 GlobalDecl GD;
300 // D could be either a kernel or a variable.
301 if (auto *FD = dyn_cast<FunctionDecl>(Val: ND))
302 GD = GlobalDecl(FD, KernelReferenceKind::Kernel);
303 else
304 GD = GlobalDecl(ND);
305 std::string DeviceSideName;
306 MangleContext *MC;
307 if (CGM.getLangOpts().CUDAIsDevice)
308 MC = &CGM.getCXXABI().getMangleContext();
309 else
310 MC = DeviceMC.get();
311 if (MC->shouldMangleDeclName(D: ND)) {
312 SmallString<256> Buffer;
313 llvm::raw_svector_ostream Out(Buffer);
314 MC->mangleName(GD, Out);
315 DeviceSideName = std::string(Out.str());
316 } else
317 DeviceSideName = std::string(ND->getIdentifier()->getName());
318
319 // Make unique name for device side static file-scope variable for HIP.
320 if (CGM.getContext().shouldExternalize(D: ND) &&
321 CGM.getLangOpts().GPURelocatableDeviceCode) {
322 SmallString<256> Buffer;
323 llvm::raw_svector_ostream Out(Buffer);
324 Out << DeviceSideName;
325 CGM.printPostfixForExternalizedDecl(OS&: Out, D: ND);
326 DeviceSideName = std::string(Out.str());
327 }
328 return DeviceSideName;
329}
330
331void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
332 FunctionArgList &Args) {
333 EmittedKernels.push_back(Elt: {.Kernel: CGF.CurFn, .D: CGF.CurFuncDecl});
334 if (auto *GV =
335 dyn_cast<llvm::GlobalVariable>(Val: KernelHandles[CGF.CurFn->getName()])) {
336 GV->setLinkage(CGF.CurFn->getLinkage());
337 GV->setInitializer(CGF.CurFn);
338 CGM.setDSOLocal(GV);
339 }
340 if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
341 CudaFeature::CUDA_USES_NEW_LAUNCH) ||
342 (CGF.getLangOpts().HIP && CGF.getLangOpts().HIPUseNewLaunchAPI) ||
343 (CGF.getLangOpts().OffloadViaLLVM))
344 emitDeviceStubBodyNew(CGF, Args);
345 else
346 emitDeviceStubBodyLegacy(CGF, Args);
347}
348
349/// CUDA passes the arguments with a level of indirection. For example, a
350/// (void*, short, void*) is passed as {void **, short *, void **} to the launch
351/// function. For the LLVM/Offload launch we include the number of arguments and
352/// their size. Thus, we pass {{void **, short*, void **}, 3, {sizeof(void*),
353/// sizeof(short), sizeof(void*)}}.
354Address CGNVCUDARuntime::prepareKernelArgsLLVMOffload(CodeGenFunction &CGF,
355 FunctionArgList &Args) {
356 SmallVector<llvm::Type *> KernelLaunchParamsTypes;
357
358 auto *Int64Ty = CGF.Builder.getInt64Ty();
359 KernelLaunchParamsTypes.push_back(Elt: PtrTy);
360 KernelLaunchParamsTypes.push_back(Elt: Int64Ty);
361 KernelLaunchParamsTypes.push_back(Elt: PtrTy);
362
363 llvm::StructType *KernelLaunchParamsTy =
364 llvm::StructType::create(Elements: KernelLaunchParamsTypes);
365 Address KernelLaunchParams = CGF.CreateTempAllocaWithoutCast(
366 Ty: KernelLaunchParamsTy, align: CharUnits::fromQuantity(Quantity: 16),
367 Name: "kernel_launch_params");
368 Address KernelArgs = CGF.CreateTempAlloca(
369 Ty: PtrTy, UseAddrSpace: LangAS::Default, align: CharUnits::fromQuantity(Quantity: 16), Name: "kernel_args",
370 ArraySize: llvm::ConstantInt::get(Ty: SizeTy, V: std::max<size_t>(a: 1, b: Args.size())));
371 Address KernelArgSizes = CGF.CreateTempAlloca(
372 Ty: SizeTy, UseAddrSpace: LangAS::Default, align: CharUnits::fromQuantity(Quantity: 16), Name: "kernel_arg_sizes",
373 ArraySize: llvm::ConstantInt::get(Ty: SizeTy, V: std::max<size_t>(a: 1, b: Args.size())));
374
375 CGF.Builder.CreateStore(Val: KernelArgs.emitRawPointer(CGF),
376 Addr: CGF.Builder.CreateStructGEP(Addr: KernelLaunchParams, Index: 0));
377 CGF.Builder.CreateStore(Val: llvm::ConstantInt::get(Ty: Int64Ty, V: Args.size()),
378 Addr: CGF.Builder.CreateStructGEP(Addr: KernelLaunchParams, Index: 1));
379 CGF.Builder.CreateStore(Val: KernelArgSizes.emitRawPointer(CGF),
380 Addr: CGF.Builder.CreateStructGEP(Addr: KernelLaunchParams, Index: 2));
381
382 for (unsigned i = 0; i < Args.size(); ++i) {
383 llvm::Value *VarPtr = CGF.GetAddrOfLocalVar(VD: Args[i]).emitRawPointer(CGF);
384 llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(V: VarPtr, DestTy: PtrTy);
385 CGF.Builder.CreateDefaultAlignedStore(
386 Val: VoidVarPtr, Addr: CGF.Builder.CreateConstGEP1_32(
387 Ty: PtrTy, Ptr: KernelArgs.emitRawPointer(CGF), Idx0: i));
388
389 auto ArgSize = CGM.getDataLayout().getTypeAllocSize(
390 Ty: CGM.getTypes().ConvertType(T: Args[i]->getType()));
391 CGF.Builder.CreateDefaultAlignedStore(
392 Val: llvm::ConstantInt::get(Ty: SizeTy, V: ArgSize),
393 Addr: CGF.Builder.CreateConstGEP1_32(Ty: SizeTy,
394 Ptr: KernelArgSizes.emitRawPointer(CGF), Idx0: i));
395 }
396
397 return KernelLaunchParams;
398}
399
400Address CGNVCUDARuntime::prepareKernelArgs(CodeGenFunction &CGF,
401 FunctionArgList &Args) {
402 // Calculate amount of space we will need for all arguments. If we have no
403 // args, allocate a single pointer so we still have a valid pointer to the
404 // argument array that we can pass to runtime, even if it will be unused.
405 Address KernelArgs = CGF.CreateTempAlloca(
406 Ty: PtrTy, UseAddrSpace: LangAS::Default, align: CharUnits::fromQuantity(Quantity: 16), Name: "kernel_args",
407 ArraySize: llvm::ConstantInt::get(Ty: SizeTy, V: std::max<size_t>(a: 1, b: Args.size())));
408 // Store pointers to the arguments in a locally allocated launch_args.
409 for (unsigned i = 0; i < Args.size(); ++i) {
410 llvm::Value *VarPtr = CGF.GetAddrOfLocalVar(VD: Args[i]).emitRawPointer(CGF);
411 llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(V: VarPtr, DestTy: PtrTy);
412 CGF.Builder.CreateDefaultAlignedStore(
413 Val: VoidVarPtr, Addr: CGF.Builder.CreateConstGEP1_32(
414 Ty: PtrTy, Ptr: KernelArgs.emitRawPointer(CGF), Idx0: i));
415 }
416 return KernelArgs;
417}
418
419// CUDA 9.0+ uses new way to launch kernels. Parameters are packed in a local
420// array and kernels are launched using cudaLaunchKernel().
421void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
422 FunctionArgList &Args) {
423 bool UsesLLVMOffloading = CGF.getLangOpts().OffloadViaLLVM;
424 // Build the shadow stack entry at the very start of the function.
425 Address KernelArgs = UsesLLVMOffloading
426 ? prepareKernelArgsLLVMOffload(CGF, Args)
427 : prepareKernelArgs(CGF, Args);
428
429 llvm::BasicBlock *EndBlock = CGF.createBasicBlock(name: "setup.end");
430
431 // Lookup cudaLaunchKernel/hipLaunchKernel function.
432 // HIP kernel launching API name depends on -fgpu-default-stream option. For
433 // the default value 'legacy', it is hipLaunchKernel. For 'per-thread',
434 // it is hipLaunchKernel_spt.
435 // cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
436 // void **args, size_t sharedMem,
437 // cudaStream_t stream);
438 // hipError_t hipLaunchKernel[_spt](const void *func, dim3 gridDim,
439 // dim3 blockDim, void **args,
440 // size_t sharedMem, hipStream_t stream);
441 TranslationUnitDecl *TUDecl = CGM.getContext().getTranslationUnitDecl();
442 DeclContext *DC = TranslationUnitDecl::castToDeclContext(D: TUDecl);
443 std::string KernelLaunchAPI = "LaunchKernel";
444 if (CGF.getLangOpts().GPUDefaultStream ==
445 LangOptions::GPUDefaultStreamKind::PerThread) {
446 if (CGF.getLangOpts().HIP)
447 KernelLaunchAPI = KernelLaunchAPI + "_spt";
448 else if (CGF.getLangOpts().CUDA)
449 KernelLaunchAPI = KernelLaunchAPI + "_ptsz";
450 }
451 auto LaunchKernelName = addPrefixToName(FuncName: KernelLaunchAPI);
452 const IdentifierInfo &cudaLaunchKernelII =
453 CGM.getContext().Idents.get(Name: LaunchKernelName);
454 FunctionDecl *cudaLaunchKernelFD = nullptr;
455 for (auto *Result : DC->lookup(Name: &cudaLaunchKernelII)) {
456 if (FunctionDecl *FD = dyn_cast<FunctionDecl>(Val: Result))
457 cudaLaunchKernelFD = FD;
458 }
459
460 if (cudaLaunchKernelFD == nullptr) {
461 CGM.Error(loc: CGF.CurFuncDecl->getLocation(),
462 error: "Can't find declaration for " + LaunchKernelName);
463 return;
464 }
465 // Create temporary dim3 grid_dim, block_dim.
466 ParmVarDecl *GridDimParam = cudaLaunchKernelFD->getParamDecl(i: 1);
467 QualType Dim3Ty = GridDimParam->getType();
468 Address GridDim = CGF.CreateMemTempWithoutCast(
469 T: Dim3Ty, Align: CharUnits::fromQuantity(Quantity: 8), Name: "grid_dim");
470 Address BlockDim = CGF.CreateMemTempWithoutCast(
471 T: Dim3Ty, Align: CharUnits::fromQuantity(Quantity: 8), Name: "block_dim");
472 Address ShmemSize = CGF.CreateTempAlloca(Ty: SizeTy, UseAddrSpace: LangAS::Default,
473 align: CGM.getSizeAlign(), Name: "shmem_size");
474 Address Stream = CGF.CreateTempAlloca(Ty: PtrTy, UseAddrSpace: LangAS::Default,
475 align: CGM.getPointerAlign(), Name: "stream");
476 llvm::FunctionCallee cudaPopConfigFn = CGM.CreateRuntimeFunction(
477 Ty: llvm::FunctionType::get(Result: IntTy,
478 Params: {/*gridDim=*/GridDim.getType(),
479 /*blockDim=*/BlockDim.getType(),
480 /*ShmemSize=*/ShmemSize.getType(),
481 /*Stream=*/Stream.getType()},
482 /*isVarArg=*/false),
483 Name: addUnderscoredPrefixToName(FuncName: "PopCallConfiguration"));
484
485 CGF.EmitRuntimeCallOrInvoke(callee: cudaPopConfigFn, args: {GridDim.emitRawPointer(CGF),
486 BlockDim.emitRawPointer(CGF),
487 ShmemSize.emitRawPointer(CGF),
488 Stream.emitRawPointer(CGF)});
489
490 // Emit the call to cudaLaunch
491 llvm::Value *Kernel =
492 CGF.Builder.CreatePointerCast(V: KernelHandles[CGF.CurFn->getName()], DestTy: PtrTy);
493 CallArgList LaunchKernelArgs;
494 LaunchKernelArgs.add(rvalue: RValue::get(V: Kernel),
495 type: cudaLaunchKernelFD->getParamDecl(i: 0)->getType());
496 LaunchKernelArgs.add(rvalue: RValue::getAggregate(addr: GridDim), type: Dim3Ty);
497 LaunchKernelArgs.add(rvalue: RValue::getAggregate(addr: BlockDim), type: Dim3Ty);
498 LaunchKernelArgs.add(rvalue: RValue::get(Addr: KernelArgs, CGF),
499 type: cudaLaunchKernelFD->getParamDecl(i: 3)->getType());
500 LaunchKernelArgs.add(rvalue: RValue::get(V: CGF.Builder.CreateLoad(Addr: ShmemSize)),
501 type: cudaLaunchKernelFD->getParamDecl(i: 4)->getType());
502 LaunchKernelArgs.add(rvalue: RValue::get(V: CGF.Builder.CreateLoad(Addr: Stream)),
503 type: cudaLaunchKernelFD->getParamDecl(i: 5)->getType());
504
505 QualType QT = cudaLaunchKernelFD->getType();
506 QualType CQT = QT.getCanonicalType();
507 llvm::Type *Ty = CGM.getTypes().ConvertType(T: CQT);
508 llvm::FunctionType *FTy = cast<llvm::FunctionType>(Val: Ty);
509
510 const CGFunctionInfo &FI =
511 CGM.getTypes().arrangeFunctionDeclaration(GD: cudaLaunchKernelFD);
512 llvm::FunctionCallee cudaLaunchKernelFn =
513 CGM.CreateRuntimeFunction(Ty: FTy, Name: LaunchKernelName);
514 CGF.EmitCall(CallInfo: FI, Callee: CGCallee::forDirect(functionPtr: cudaLaunchKernelFn), ReturnValue: ReturnValueSlot(),
515 Args: LaunchKernelArgs);
516
517 // To prevent CUDA device stub functions from being merged by ICF in MSVC
518 // environment, create an unique global variable for each kernel and write to
519 // the variable in the device stub.
520 if (CGM.getContext().getTargetInfo().getCXXABI().isMicrosoft() &&
521 !CGF.getLangOpts().HIP) {
522 llvm::Function *KernelFunction = llvm::cast<llvm::Function>(Val: Kernel);
523 std::string GlobalVarName = (KernelFunction->getName() + ".id").str();
524
525 llvm::GlobalVariable *HandleVar =
526 CGM.getModule().getNamedGlobal(Name: GlobalVarName);
527 if (!HandleVar) {
528 HandleVar = new llvm::GlobalVariable(
529 CGM.getModule(), CGM.Int8Ty,
530 /*Constant=*/false, KernelFunction->getLinkage(),
531 llvm::ConstantInt::get(Ty: CGM.Int8Ty, V: 0), GlobalVarName);
532 HandleVar->setDSOLocal(KernelFunction->isDSOLocal());
533 HandleVar->setVisibility(KernelFunction->getVisibility());
534 if (KernelFunction->hasComdat())
535 HandleVar->setComdat(CGM.getModule().getOrInsertComdat(Name: GlobalVarName));
536 }
537
538 CGF.Builder.CreateAlignedStore(Val: llvm::ConstantInt::get(Ty: CGM.Int8Ty, V: 1),
539 Addr: HandleVar, Align: CharUnits::One(),
540 /*IsVolatile=*/true);
541 }
542
543 CGF.EmitBranch(Block: EndBlock);
544
545 CGF.EmitBlock(BB: EndBlock);
546}
547
548void CGNVCUDARuntime::emitDeviceStubBodyLegacy(CodeGenFunction &CGF,
549 FunctionArgList &Args) {
550 // Emit a call to cudaSetupArgument for each arg in Args.
551 llvm::FunctionCallee cudaSetupArgFn = getSetupArgumentFn();
552 llvm::BasicBlock *EndBlock = CGF.createBasicBlock(name: "setup.end");
553 CharUnits Offset = CharUnits::Zero();
554 for (const VarDecl *A : Args) {
555 auto TInfo = CGM.getContext().getTypeInfoInChars(T: A->getType());
556 Offset = Offset.alignTo(Align: TInfo.Align);
557 llvm::Value *Args[] = {
558 CGF.Builder.CreatePointerCast(
559 V: CGF.GetAddrOfLocalVar(VD: A).emitRawPointer(CGF), DestTy: PtrTy),
560 llvm::ConstantInt::get(Ty: SizeTy, V: TInfo.Width.getQuantity()),
561 llvm::ConstantInt::get(Ty: SizeTy, V: Offset.getQuantity()),
562 };
563 llvm::CallBase *CB = CGF.EmitRuntimeCallOrInvoke(callee: cudaSetupArgFn, args: Args);
564 llvm::Constant *Zero = llvm::ConstantInt::get(Ty: IntTy, V: 0);
565 llvm::Value *CBZero = CGF.Builder.CreateICmpEQ(LHS: CB, RHS: Zero);
566 llvm::BasicBlock *NextBlock = CGF.createBasicBlock(name: "setup.next");
567 CGF.Builder.CreateCondBr(Cond: CBZero, True: NextBlock, False: EndBlock);
568 CGF.EmitBlock(BB: NextBlock);
569 Offset += TInfo.Width;
570 }
571
572 // Emit the call to cudaLaunch
573 llvm::FunctionCallee cudaLaunchFn = getLaunchFn();
574 llvm::Value *Arg =
575 CGF.Builder.CreatePointerCast(V: KernelHandles[CGF.CurFn->getName()], DestTy: PtrTy);
576 CGF.EmitRuntimeCallOrInvoke(callee: cudaLaunchFn, args: Arg);
577 CGF.EmitBranch(Block: EndBlock);
578
579 CGF.EmitBlock(BB: EndBlock);
580}
581
582// Replace the original variable Var with the address loaded from variable
583// ManagedVar populated by HIP runtime.
584static void replaceManagedVar(llvm::GlobalVariable *Var,
585 llvm::GlobalVariable *ManagedVar) {
586 SmallVector<SmallVector<llvm::User *, 8>, 8> WorkList;
587 for (auto &&VarUse : Var->uses()) {
588 WorkList.push_back(Elt: {VarUse.getUser()});
589 }
590 while (!WorkList.empty()) {
591 auto &&WorkItem = WorkList.pop_back_val();
592 auto *U = WorkItem.back();
593 if (isa<llvm::ConstantExpr>(Val: U)) {
594 for (auto &&UU : U->uses()) {
595 WorkItem.push_back(Elt: UU.getUser());
596 WorkList.push_back(Elt: WorkItem);
597 WorkItem.pop_back();
598 }
599 continue;
600 }
601 if (auto *I = dyn_cast<llvm::Instruction>(Val: U)) {
602 llvm::Value *OldV = Var;
603 llvm::Instruction *NewV =
604 new llvm::LoadInst(Var->getType(), ManagedVar, "ld.managed", false,
605 Var->getAlign().valueOrOne(), I->getIterator());
606 WorkItem.pop_back();
607 // Replace constant expressions directly or indirectly using the managed
608 // variable with instructions.
609 for (auto &&Op : WorkItem) {
610 auto *CE = cast<llvm::ConstantExpr>(Val: Op);
611 auto *NewInst = CE->getAsInstruction();
612 NewInst->insertBefore(BB&: *I->getParent(), InsertPos: I->getIterator());
613 NewInst->replaceUsesOfWith(From: OldV, To: NewV);
614 OldV = CE;
615 NewV = NewInst;
616 }
617 I->replaceUsesOfWith(From: OldV, To: NewV);
618 } else {
619 llvm_unreachable("Invalid use of managed variable");
620 }
621 }
622}
623
624/// Creates a function that sets up state on the host side for CUDA objects that
625/// have a presence on both the host and device sides. Specifically, registers
626/// the host side of kernel functions and device global variables with the CUDA
627/// runtime.
628/// \code
629/// void __cuda_register_globals(void** GpuBinaryHandle) {
630/// __cudaRegisterFunction(GpuBinaryHandle,Kernel0,...);
631/// ...
632/// __cudaRegisterFunction(GpuBinaryHandle,KernelM,...);
633/// __cudaRegisterVar(GpuBinaryHandle, GlobalVar0, ...);
634/// ...
635/// __cudaRegisterVar(GpuBinaryHandle, GlobalVarN, ...);
636/// }
637/// \endcode
638llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
639 // No need to register anything
640 if (EmittedKernels.empty() && DeviceVars.empty())
641 return nullptr;
642
643 llvm::Function *RegisterKernelsFunc = llvm::Function::Create(
644 Ty: getRegisterGlobalsFnTy(), Linkage: llvm::GlobalValue::InternalLinkage,
645 N: addUnderscoredPrefixToName(FuncName: "_register_globals"), M: &TheModule);
646 llvm::BasicBlock *EntryBB =
647 llvm::BasicBlock::Create(Context, Name: "entry", Parent: RegisterKernelsFunc);
648 CGBuilderTy Builder(CGM, Context);
649 Builder.SetInsertPoint(EntryBB);
650
651 // void __cudaRegisterFunction(void **, const char *, char *, const char *,
652 // int, uint3*, uint3*, dim3*, dim3*, int*)
653 llvm::Type *RegisterFuncParams[] = {
654 PtrTy, PtrTy, PtrTy, PtrTy, IntTy,
655 PtrTy, PtrTy, PtrTy, PtrTy, llvm::PointerType::getUnqual(C&: Context)};
656 llvm::FunctionCallee RegisterFunc = CGM.CreateRuntimeFunction(
657 Ty: llvm::FunctionType::get(Result: IntTy, Params: RegisterFuncParams, isVarArg: false),
658 Name: addUnderscoredPrefixToName(FuncName: "RegisterFunction"));
659
660 // Extract GpuBinaryHandle passed as the first argument passed to
661 // __cuda_register_globals() and generate __cudaRegisterFunction() call for
662 // each emitted kernel.
663 llvm::Argument &GpuBinaryHandlePtr = *RegisterKernelsFunc->arg_begin();
664 for (auto &&I : EmittedKernels) {
665 llvm::Constant *KernelName =
666 makeConstantString(Str: getDeviceSideName(ND: cast<NamedDecl>(Val: I.D)));
667 llvm::Constant *NullPtr = llvm::ConstantPointerNull::get(T: PtrTy);
668 llvm::Value *Args[] = {
669 &GpuBinaryHandlePtr,
670 KernelHandles[I.Kernel->getName()],
671 KernelName,
672 KernelName,
673 llvm::ConstantInt::getAllOnesValue(Ty: IntTy),
674 NullPtr,
675 NullPtr,
676 NullPtr,
677 NullPtr,
678 llvm::ConstantPointerNull::get(T: llvm::PointerType::getUnqual(C&: Context))};
679 Builder.CreateCall(Callee: RegisterFunc, Args);
680 }
681
682 llvm::Type *VarSizeTy = IntTy;
683 // For HIP or CUDA 9.0+, device variable size is type of `size_t`.
684 if (CGM.getLangOpts().HIP ||
685 ToCudaVersion(CGM.getTarget().getSDKVersion()) >= CudaVersion::CUDA_90)
686 VarSizeTy = SizeTy;
687
688 // void __cudaRegisterVar(void **, char *, char *, const char *,
689 // int, int, int, int)
690 llvm::Type *RegisterVarParams[] = {PtrTy, PtrTy, PtrTy, PtrTy,
691 IntTy, VarSizeTy, IntTy, IntTy};
692 llvm::FunctionCallee RegisterVar = CGM.CreateRuntimeFunction(
693 Ty: llvm::FunctionType::get(Result: VoidTy, Params: RegisterVarParams, isVarArg: false),
694 Name: addUnderscoredPrefixToName(FuncName: "RegisterVar"));
695 // void __hipRegisterManagedVar(void **, char *, char *, const char *,
696 // size_t, unsigned)
697 llvm::Type *RegisterManagedVarParams[] = {PtrTy, PtrTy, PtrTy,
698 PtrTy, VarSizeTy, IntTy};
699 llvm::FunctionCallee RegisterManagedVar = CGM.CreateRuntimeFunction(
700 Ty: llvm::FunctionType::get(Result: VoidTy, Params: RegisterManagedVarParams, isVarArg: false),
701 Name: addUnderscoredPrefixToName(FuncName: "RegisterManagedVar"));
702 // void __cudaRegisterSurface(void **, const struct surfaceReference *,
703 // const void **, const char *, int, int);
704 llvm::FunctionCallee RegisterSurf = CGM.CreateRuntimeFunction(
705 Ty: llvm::FunctionType::get(
706 Result: VoidTy, Params: {PtrTy, PtrTy, PtrTy, PtrTy, IntTy, IntTy}, isVarArg: false),
707 Name: addUnderscoredPrefixToName(FuncName: "RegisterSurface"));
708 // void __cudaRegisterTexture(void **, const struct textureReference *,
709 // const void **, const char *, int, int, int)
710 llvm::FunctionCallee RegisterTex = CGM.CreateRuntimeFunction(
711 Ty: llvm::FunctionType::get(
712 Result: VoidTy, Params: {PtrTy, PtrTy, PtrTy, PtrTy, IntTy, IntTy, IntTy}, isVarArg: false),
713 Name: addUnderscoredPrefixToName(FuncName: "RegisterTexture"));
714 for (auto &&Info : DeviceVars) {
715 llvm::GlobalVariable *Var = Info.Var;
716 assert((!Var->isDeclaration() || Info.Flags.isManaged()) &&
717 "External variables should not show up here, except HIP managed "
718 "variables");
719 llvm::Constant *VarName = makeConstantString(Str: getDeviceSideName(ND: Info.D));
720 switch (Info.Flags.getKind()) {
721 case DeviceVarFlags::Variable: {
722 uint64_t VarSize =
723 CGM.getDataLayout().getTypeAllocSize(Ty: Var->getValueType());
724 if (Info.Flags.isManaged()) {
725 assert(Var->getName().ends_with(".managed") &&
726 "HIP managed variables not transformed");
727 auto *ManagedVar = CGM.getModule().getNamedGlobal(
728 Name: Var->getName().drop_back(N: StringRef(".managed").size()));
729 llvm::Value *Args[] = {
730 &GpuBinaryHandlePtr,
731 ManagedVar,
732 Var,
733 VarName,
734 llvm::ConstantInt::get(Ty: VarSizeTy, V: VarSize),
735 llvm::ConstantInt::get(Ty: IntTy,
736 V: Var->getAlign().valueOrOne().value())};
737 if (!Var->isDeclaration())
738 Builder.CreateCall(Callee: RegisterManagedVar, Args);
739 } else {
740 llvm::Value *Args[] = {
741 &GpuBinaryHandlePtr,
742 Var,
743 VarName,
744 VarName,
745 llvm::ConstantInt::get(Ty: IntTy, V: Info.Flags.isExtern()),
746 llvm::ConstantInt::get(Ty: VarSizeTy, V: VarSize),
747 llvm::ConstantInt::get(Ty: IntTy, V: Info.Flags.isConstant()),
748 llvm::ConstantInt::get(Ty: IntTy, V: 0)};
749 Builder.CreateCall(Callee: RegisterVar, Args);
750 }
751 break;
752 }
753 case DeviceVarFlags::Surface:
754 Builder.CreateCall(
755 Callee: RegisterSurf,
756 Args: {&GpuBinaryHandlePtr, Var, VarName, VarName,
757 llvm::ConstantInt::get(Ty: IntTy, V: Info.Flags.getSurfTexType()),
758 llvm::ConstantInt::get(Ty: IntTy, V: Info.Flags.isExtern())});
759 break;
760 case DeviceVarFlags::Texture:
761 Builder.CreateCall(
762 Callee: RegisterTex,
763 Args: {&GpuBinaryHandlePtr, Var, VarName, VarName,
764 llvm::ConstantInt::get(Ty: IntTy, V: Info.Flags.getSurfTexType()),
765 llvm::ConstantInt::get(Ty: IntTy, V: Info.Flags.isNormalized()),
766 llvm::ConstantInt::get(Ty: IntTy, V: Info.Flags.isExtern())});
767 break;
768 }
769 }
770
771 // Register the per-TU offload-profiling shadow so the host runtime can
772 // locate the matching device-side __llvm_profile_sections_<CUID>. We
773 // emit both __hipRegisterVar (so the HIP runtime can map the host
774 // shadow to the device symbol) and
775 // __llvm_profile_offload_register_shadow_variable (so the profile
776 // runtime adds the shadow to its drain list).
777 if (OffloadProfShadow) {
778 llvm::Constant *Name =
779 makeConstantString(Str: std::string(OffloadProfShadow->getName()));
780 llvm::Constant *IntZero = llvm::ConstantInt::get(Ty: IntTy, V: 0);
781 llvm::Value *RegisterVarArgs[] = {
782 &GpuBinaryHandlePtr,
783 OffloadProfShadow,
784 Name,
785 Name,
786 IntZero,
787 llvm::ConstantInt::get(Ty: VarSizeTy,
788 V: CGM.getDataLayout().getPointerSize(/*AS=*/0)),
789 IntZero,
790 IntZero};
791 Builder.CreateCall(Callee: RegisterVar, Args: RegisterVarArgs);
792
793 llvm::FunctionCallee RegisterShadow = CGM.CreateRuntimeFunction(
794 Ty: llvm::FunctionType::get(Result: VoidTy, Params: {PtrTy}, isVarArg: false),
795 Name: "__llvm_profile_offload_register_shadow_variable");
796 Builder.CreateCall(Callee: RegisterShadow, Args: {OffloadProfShadow});
797 }
798
799 if (!OffloadProfSectionShadows.empty()) {
800 llvm::FunctionCallee RegisterSectionShadow = CGM.CreateRuntimeFunction(
801 Ty: llvm::FunctionType::get(Result: VoidTy, Params: {PtrTy}, isVarArg: false),
802 Name: "__llvm_profile_offload_register_section_shadow_variable");
803 llvm::Constant *IntZero = llvm::ConstantInt::get(Ty: IntTy, V: 0);
804 for (const auto &Info : OffloadProfSectionShadows) {
805 llvm::Constant *Name = makeConstantString(Str: Info.DeviceName);
806 llvm::Value *RegisterVarArgs[] = {
807 &GpuBinaryHandlePtr,
808 Info.Shadow,
809 Name,
810 Name,
811 IntZero,
812 llvm::ConstantInt::get(Ty: VarSizeTy,
813 V: CGM.getDataLayout().getPointerSize(/*AS=*/0)),
814 IntZero,
815 IntZero};
816 Builder.CreateCall(Callee: RegisterVar, Args: RegisterVarArgs);
817 Builder.CreateCall(Callee: RegisterSectionShadow, Args: {Info.Shadow});
818 }
819 }
820
821 Builder.CreateRetVoid();
822 return RegisterKernelsFunc;
823}
824
825/// Creates a global constructor function for the module:
826///
827/// For CUDA:
828/// \code
829/// void __cuda_module_ctor() {
830/// Handle = __cudaRegisterFatBinary(GpuBinaryBlob);
831/// __cuda_register_globals(Handle);
832/// }
833/// \endcode
834///
835/// For HIP:
836/// \code
837/// void __hip_module_ctor() {
838/// if (__hip_gpubin_handle == 0) {
839/// __hip_gpubin_handle = __hipRegisterFatBinary(GpuBinaryBlob);
840/// __hip_register_globals(__hip_gpubin_handle);
841/// }
842/// }
843/// \endcode
844llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
845 bool IsHIP = CGM.getLangOpts().HIP;
846 bool IsCUDA = CGM.getLangOpts().CUDA;
847 // No need to generate ctors/dtors if there is no GPU binary.
848 StringRef CudaGpuBinaryFileName =
849 CGM.getCodeGenOpts().OffloadBinaryToEmbedFile;
850 if (CudaGpuBinaryFileName.empty() && !IsHIP)
851 return nullptr;
852 if ((IsHIP || (IsCUDA && !RelocatableDeviceCode)) && EmittedKernels.empty() &&
853 DeviceVars.empty())
854 return nullptr;
855
856 // void __{cuda|hip}_register_globals(void* handle);
857 llvm::Function *RegisterGlobalsFunc = makeRegisterGlobalsFn();
858 // We always need a function to pass in as callback. Create a dummy
859 // implementation if we don't need to register anything.
860 if (RelocatableDeviceCode && !RegisterGlobalsFunc)
861 RegisterGlobalsFunc = makeDummyFunction(FnTy: getRegisterGlobalsFnTy());
862
863 // void ** __{cuda|hip}RegisterFatBinary(void *);
864 llvm::FunctionCallee RegisterFatbinFunc = CGM.CreateRuntimeFunction(
865 Ty: llvm::FunctionType::get(Result: PtrTy, Params: PtrTy, isVarArg: false),
866 Name: addUnderscoredPrefixToName(FuncName: "RegisterFatBinary"));
867 // struct { int magic, int version, void * gpu_binary, void * dont_care };
868 llvm::StructType *FatbinWrapperTy =
869 llvm::StructType::get(elt1: IntTy, elts: IntTy, elts: PtrTy, elts: PtrTy);
870
871 // Register GPU binary with the CUDA runtime, store returned handle in a
872 // global variable and save a reference in GpuBinaryHandle to be cleaned up
873 // in destructor on exit. Then associate all known kernels with the GPU binary
874 // handle so CUDA runtime can figure out what to call on the GPU side.
875 std::unique_ptr<llvm::MemoryBuffer> CudaGpuBinary = nullptr;
876 if (!CudaGpuBinaryFileName.empty()) {
877 auto VFS = CGM.getFileSystem();
878 auto CudaGpuBinaryOrErr =
879 VFS->getBufferForFile(Name: CudaGpuBinaryFileName, FileSize: -1, RequiresNullTerminator: false);
880 if (std::error_code EC = CudaGpuBinaryOrErr.getError()) {
881 CGM.getDiags().Report(DiagID: diag::err_cannot_open_file)
882 << CudaGpuBinaryFileName << EC.message();
883 return nullptr;
884 }
885 CudaGpuBinary = std::move(CudaGpuBinaryOrErr.get());
886 }
887
888 llvm::Function *ModuleCtorFunc = llvm::Function::Create(
889 Ty: llvm::FunctionType::get(Result: VoidTy, isVarArg: false),
890 Linkage: llvm::GlobalValue::InternalLinkage,
891 N: addUnderscoredPrefixToName(FuncName: "_module_ctor"), M: &TheModule);
892 llvm::BasicBlock *CtorEntryBB =
893 llvm::BasicBlock::Create(Context, Name: "entry", Parent: ModuleCtorFunc);
894 CGBuilderTy CtorBuilder(CGM, Context);
895
896 CtorBuilder.SetInsertPoint(CtorEntryBB);
897
898 const char *FatbinConstantName;
899 const char *FatbinSectionName;
900 const char *ModuleIDSectionName;
901 StringRef ModuleIDPrefix;
902 llvm::Constant *FatBinStr;
903 unsigned FatMagic;
904 if (IsHIP) {
905 // On macOS (Mach-O), section names must be in "segment,section" format.
906 FatbinConstantName =
907 CGM.getTriple().isMacOSX() ? "__HIP,__hip_fatbin" : ".hip_fatbin";
908 FatbinSectionName =
909 CGM.getTriple().isMacOSX() ? "__HIP,__fatbin" : ".hipFatBinSegment";
910
911 ModuleIDSectionName =
912 CGM.getTriple().isMacOSX() ? "__HIP,__module_id" : "__hip_module_id";
913 ModuleIDPrefix = "__hip_";
914
915 if (CudaGpuBinary) {
916 // If fatbin is available from early finalization, create a string
917 // literal containing the fat binary loaded from the given file.
918 const unsigned HIPCodeObjectAlign = 4096;
919 FatBinStr = makeConstantArray(Str: std::string(CudaGpuBinary->getBuffer()), Name: "",
920 SectionName: FatbinConstantName, Alignment: HIPCodeObjectAlign);
921 } else {
922 // If fatbin is not available, create an external symbol
923 // __hip_fatbin in section .hip_fatbin. The external symbol is supposed
924 // to contain the fat binary but will be populated somewhere else,
925 // e.g. by lld through link script.
926 FatBinStr = new llvm::GlobalVariable(
927 CGM.getModule(), CGM.Int8Ty,
928 /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage, nullptr,
929 "__hip_fatbin" + (CGM.getLangOpts().CUID.empty()
930 ? ""
931 : "_" + CGM.getContext().getCUIDHash()),
932 nullptr, llvm::GlobalVariable::NotThreadLocal);
933 cast<llvm::GlobalVariable>(Val: FatBinStr)->setSection(FatbinConstantName);
934 }
935
936 FatMagic = HIPFatMagic;
937 } else {
938 if (RelocatableDeviceCode)
939 FatbinConstantName = CGM.getTriple().isMacOSX()
940 ? "__NV_CUDA,__nv_relfatbin"
941 : "__nv_relfatbin";
942 else
943 FatbinConstantName =
944 CGM.getTriple().isMacOSX() ? "__NV_CUDA,__nv_fatbin" : ".nv_fatbin";
945 // NVIDIA's cuobjdump looks for fatbins in this section.
946 FatbinSectionName =
947 CGM.getTriple().isMacOSX() ? "__NV_CUDA,__fatbin" : ".nvFatBinSegment";
948
949 ModuleIDSectionName = CGM.getTriple().isMacOSX()
950 ? "__NV_CUDA,__nv_module_id"
951 : "__nv_module_id";
952 ModuleIDPrefix = "__nv_";
953
954 // For CUDA, create a string literal containing the fat binary loaded from
955 // the given file.
956 FatBinStr = makeConstantArray(Str: std::string(CudaGpuBinary->getBuffer()), Name: "",
957 SectionName: FatbinConstantName, Alignment: 8);
958 FatMagic = CudaFatMagic;
959 }
960
961 // Create initialized wrapper structure that points to the loaded GPU binary
962 ConstantInitBuilder Builder(CGM);
963 auto Values = Builder.beginStruct(structTy: FatbinWrapperTy);
964 // Fatbin wrapper magic.
965 Values.addInt(intTy: IntTy, value: FatMagic);
966 // Fatbin version.
967 Values.addInt(intTy: IntTy, value: 1);
968 // Data.
969 Values.add(value: FatBinStr);
970 // Unused in fatbin v1.
971 Values.add(value: llvm::ConstantPointerNull::get(T: PtrTy));
972 llvm::GlobalVariable *FatbinWrapper = Values.finishAndCreateGlobal(
973 args: addUnderscoredPrefixToName(FuncName: "_fatbin_wrapper"), args: CGM.getPointerAlign(),
974 /*constant*/ args: true);
975 FatbinWrapper->setSection(FatbinSectionName);
976 CGM.getSanitizerMetadata()->disableSanitizerForGlobal(GV: FatbinWrapper);
977
978 // There is only one HIP fat binary per linked module, however there are
979 // multiple constructor functions. Make sure the fat binary is registered
980 // only once. The constructor functions are executed by the dynamic loader
981 // before the program gains control. The dynamic loader cannot execute the
982 // constructor functions concurrently since doing that would not guarantee
983 // thread safety of the loaded program. Therefore we can assume sequential
984 // execution of constructor functions here.
985 if (IsHIP) {
986 auto Linkage = RelocatableDeviceCode ? llvm::GlobalValue::ExternalLinkage
987 : llvm::GlobalValue::InternalLinkage;
988 llvm::BasicBlock *IfBlock =
989 llvm::BasicBlock::Create(Context, Name: "if", Parent: ModuleCtorFunc);
990 llvm::BasicBlock *ExitBlock =
991 llvm::BasicBlock::Create(Context, Name: "exit", Parent: ModuleCtorFunc);
992 // The name, size, and initialization pattern of this variable is part
993 // of HIP ABI.
994 GpuBinaryHandle = new llvm::GlobalVariable(
995 TheModule, PtrTy, /*isConstant=*/false, Linkage,
996 /*Initializer=*/
997 !RelocatableDeviceCode ? llvm::ConstantPointerNull::get(T: PtrTy)
998 : nullptr,
999 "__hip_gpubin_handle" + (CGM.getLangOpts().CUID.empty()
1000 ? ""
1001 : "_" + CGM.getContext().getCUIDHash()));
1002 GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign());
1003 // Prevent the weak symbol in different shared libraries being merged.
1004 if (Linkage != llvm::GlobalValue::InternalLinkage)
1005 GpuBinaryHandle->setVisibility(llvm::GlobalValue::HiddenVisibility);
1006 Address GpuBinaryAddr(
1007 GpuBinaryHandle, PtrTy,
1008 CharUnits::fromQuantity(Quantity: GpuBinaryHandle->getAlign().valueOrOne()));
1009 {
1010 auto *HandleValue = CtorBuilder.CreateLoad(Addr: GpuBinaryAddr);
1011 llvm::Constant *Zero =
1012 llvm::Constant::getNullValue(Ty: HandleValue->getType());
1013 llvm::Value *EQZero = CtorBuilder.CreateICmpEQ(LHS: HandleValue, RHS: Zero);
1014 CtorBuilder.CreateCondBr(Cond: EQZero, True: IfBlock, False: ExitBlock);
1015 }
1016 {
1017 CtorBuilder.SetInsertPoint(IfBlock);
1018 // GpuBinaryHandle = __hipRegisterFatBinary(&FatbinWrapper);
1019 llvm::CallInst *RegisterFatbinCall =
1020 CtorBuilder.CreateCall(Callee: RegisterFatbinFunc, Args: FatbinWrapper);
1021 CtorBuilder.CreateStore(Val: RegisterFatbinCall, Addr: GpuBinaryAddr);
1022 CtorBuilder.CreateBr(Dest: ExitBlock);
1023 }
1024 {
1025 CtorBuilder.SetInsertPoint(ExitBlock);
1026 // Call __hip_register_globals(GpuBinaryHandle);
1027 if (RegisterGlobalsFunc) {
1028 auto *HandleValue = CtorBuilder.CreateLoad(Addr: GpuBinaryAddr);
1029 CtorBuilder.CreateCall(Callee: RegisterGlobalsFunc, Args: HandleValue);
1030 }
1031 }
1032 } else if (!RelocatableDeviceCode) {
1033 // Register binary with CUDA runtime. This is substantially different in
1034 // default mode vs. separate compilation!
1035 // GpuBinaryHandle = __cudaRegisterFatBinary(&FatbinWrapper);
1036 llvm::CallInst *RegisterFatbinCall =
1037 CtorBuilder.CreateCall(Callee: RegisterFatbinFunc, Args: FatbinWrapper);
1038 GpuBinaryHandle = new llvm::GlobalVariable(
1039 TheModule, PtrTy, false, llvm::GlobalValue::InternalLinkage,
1040 llvm::ConstantPointerNull::get(T: PtrTy), "__cuda_gpubin_handle");
1041 GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign());
1042 CtorBuilder.CreateAlignedStore(Val: RegisterFatbinCall, Addr: GpuBinaryHandle,
1043 Align: CGM.getPointerAlign());
1044
1045 // Call __cuda_register_globals(GpuBinaryHandle);
1046 if (RegisterGlobalsFunc)
1047 CtorBuilder.CreateCall(Callee: RegisterGlobalsFunc, Args: RegisterFatbinCall);
1048
1049 // Call __cudaRegisterFatBinaryEnd(Handle) if this CUDA version needs it.
1050 if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
1051 CudaFeature::CUDA_USES_FATBIN_REGISTER_END)) {
1052 // void __cudaRegisterFatBinaryEnd(void **);
1053 llvm::FunctionCallee RegisterFatbinEndFunc = CGM.CreateRuntimeFunction(
1054 Ty: llvm::FunctionType::get(Result: VoidTy, Params: PtrTy, isVarArg: false),
1055 Name: "__cudaRegisterFatBinaryEnd");
1056 CtorBuilder.CreateCall(Callee: RegisterFatbinEndFunc, Args: RegisterFatbinCall);
1057 }
1058 } else {
1059 // Generate a unique module ID.
1060 // Note that this is unique in a build (with some collision probability
1061 // inherent to MD5 hashing) as long as each compilation sees modules with
1062 // different `SourceFileName`s. Builds using absolute paths or paths
1063 // relative to the same base path should be OK. This is similar to the
1064 // guarantees for ThinLTO and GlobalValue's GUID.
1065 // If desired, a stronger uniqueness guarantee could be computed (with a
1066 // small refactoring) with `llvm::getUniqueModuleId`, which hashes the
1067 // module content (and, therefore, a compile-time tradeoff).
1068 SmallString<64> ModuleID;
1069 llvm::raw_svector_ostream OS(ModuleID);
1070 OS << ModuleIDPrefix
1071 << llvm::format(Fmt: "%" PRIx64,
1072 Vals: llvm::MD5Hash(Str: TheModule.getSourceFileName()));
1073 llvm::Constant *ModuleIDConstant = makeConstantArray(
1074 Str: std::string(ModuleID), Name: "", SectionName: ModuleIDSectionName, Alignment: 32, /*AddNull=*/true);
1075
1076 // Create an alias for the FatbinWrapper that nvcc will look for.
1077 llvm::GlobalAlias::create(Linkage: llvm::GlobalValue::ExternalLinkage,
1078 Name: Twine("__fatbinwrap") + ModuleID, Aliasee: FatbinWrapper);
1079
1080 // void __cudaRegisterLinkedBinary%ModuleID%(void (*)(void *), void *,
1081 // void *, void (*)(void **))
1082 SmallString<128> RegisterLinkedBinaryName("__cudaRegisterLinkedBinary");
1083 RegisterLinkedBinaryName += ModuleID;
1084 llvm::FunctionCallee RegisterLinkedBinaryFunc = CGM.CreateRuntimeFunction(
1085 Ty: getRegisterLinkedBinaryFnTy(), Name: RegisterLinkedBinaryName);
1086
1087 assert(RegisterGlobalsFunc && "Expecting at least dummy function!");
1088 llvm::Value *Args[] = {RegisterGlobalsFunc, FatbinWrapper, ModuleIDConstant,
1089 makeDummyFunction(FnTy: getCallbackFnTy())};
1090 CtorBuilder.CreateCall(Callee: RegisterLinkedBinaryFunc, Args);
1091 }
1092
1093 // Create destructor and register it with atexit() the way NVCC does it. Doing
1094 // it during regular destructor phase worked in CUDA before 9.2 but results in
1095 // double-free in 9.2.
1096 if (llvm::Function *CleanupFn = makeModuleDtorFunction()) {
1097 // extern "C" int atexit(void (*f)(void));
1098 llvm::FunctionType *AtExitTy =
1099 llvm::FunctionType::get(Result: IntTy, Params: CleanupFn->getType(), isVarArg: false);
1100 llvm::FunctionCallee AtExitFunc =
1101 CGM.CreateRuntimeFunction(Ty: AtExitTy, Name: "atexit", ExtraAttrs: llvm::AttributeList(),
1102 /*Local=*/true);
1103 CtorBuilder.CreateCall(Callee: AtExitFunc, Args: CleanupFn);
1104 }
1105
1106 CtorBuilder.CreateRetVoid();
1107 return ModuleCtorFunc;
1108}
1109
1110/// Creates a global destructor function that unregisters the GPU code blob
1111/// registered by constructor.
1112///
1113/// For CUDA:
1114/// \code
1115/// void __cuda_module_dtor() {
1116/// __cudaUnregisterFatBinary(Handle);
1117/// }
1118/// \endcode
1119///
1120/// For HIP:
1121/// \code
1122/// void __hip_module_dtor() {
1123/// if (__hip_gpubin_handle) {
1124/// __hipUnregisterFatBinary(__hip_gpubin_handle);
1125/// __hip_gpubin_handle = 0;
1126/// }
1127/// }
1128/// \endcode
1129llvm::Function *CGNVCUDARuntime::makeModuleDtorFunction() {
1130 // No need for destructor if we don't have a handle to unregister.
1131 if (!GpuBinaryHandle)
1132 return nullptr;
1133
1134 // void __cudaUnregisterFatBinary(void ** handle);
1135 llvm::FunctionCallee UnregisterFatbinFunc = CGM.CreateRuntimeFunction(
1136 Ty: llvm::FunctionType::get(Result: VoidTy, Params: PtrTy, isVarArg: false),
1137 Name: addUnderscoredPrefixToName(FuncName: "UnregisterFatBinary"));
1138
1139 llvm::Function *ModuleDtorFunc = llvm::Function::Create(
1140 Ty: llvm::FunctionType::get(Result: VoidTy, isVarArg: false),
1141 Linkage: llvm::GlobalValue::InternalLinkage,
1142 N: addUnderscoredPrefixToName(FuncName: "_module_dtor"), M: &TheModule);
1143
1144 llvm::BasicBlock *DtorEntryBB =
1145 llvm::BasicBlock::Create(Context, Name: "entry", Parent: ModuleDtorFunc);
1146 CGBuilderTy DtorBuilder(CGM, Context);
1147 DtorBuilder.SetInsertPoint(DtorEntryBB);
1148
1149 Address GpuBinaryAddr(
1150 GpuBinaryHandle, GpuBinaryHandle->getValueType(),
1151 CharUnits::fromQuantity(Quantity: GpuBinaryHandle->getAlign().valueOrOne()));
1152 auto *HandleValue = DtorBuilder.CreateLoad(Addr: GpuBinaryAddr);
1153 // There is only one HIP fat binary per linked module, however there are
1154 // multiple destructor functions. Make sure the fat binary is unregistered
1155 // only once.
1156 if (CGM.getLangOpts().HIP) {
1157 llvm::BasicBlock *IfBlock =
1158 llvm::BasicBlock::Create(Context, Name: "if", Parent: ModuleDtorFunc);
1159 llvm::BasicBlock *ExitBlock =
1160 llvm::BasicBlock::Create(Context, Name: "exit", Parent: ModuleDtorFunc);
1161 llvm::Constant *Zero = llvm::Constant::getNullValue(Ty: HandleValue->getType());
1162 llvm::Value *NEZero = DtorBuilder.CreateICmpNE(LHS: HandleValue, RHS: Zero);
1163 DtorBuilder.CreateCondBr(Cond: NEZero, True: IfBlock, False: ExitBlock);
1164
1165 DtorBuilder.SetInsertPoint(IfBlock);
1166 DtorBuilder.CreateCall(Callee: UnregisterFatbinFunc, Args: HandleValue);
1167 DtorBuilder.CreateStore(Val: Zero, Addr: GpuBinaryAddr);
1168 DtorBuilder.CreateBr(Dest: ExitBlock);
1169
1170 DtorBuilder.SetInsertPoint(ExitBlock);
1171 } else {
1172 DtorBuilder.CreateCall(Callee: UnregisterFatbinFunc, Args: HandleValue);
1173 }
1174 DtorBuilder.CreateRetVoid();
1175 return ModuleDtorFunc;
1176}
1177
1178CGCUDARuntime *CodeGen::CreateNVCUDARuntime(CodeGenModule &CGM) {
1179 return new CGNVCUDARuntime(CGM);
1180}
1181
1182void CGNVCUDARuntime::internalizeDeviceSideVar(
1183 const VarDecl *D, llvm::GlobalValue::LinkageTypes &Linkage) {
1184 // For -fno-gpu-rdc, host-side shadows of external declarations of device-side
1185 // global variables become internal definitions. These have to be internal in
1186 // order to prevent name conflicts with global host variables with the same
1187 // name in a different TUs.
1188 //
1189 // For -fgpu-rdc, the shadow variables should not be internalized because
1190 // they may be accessed by different TU.
1191 if (CGM.getLangOpts().GPURelocatableDeviceCode)
1192 return;
1193
1194 // __shared__ variables are odd. Shadows do get created, but
1195 // they are not registered with the CUDA runtime, so they
1196 // can't really be used to access their device-side
1197 // counterparts. It's not clear yet whether it's nvcc's bug or
1198 // a feature, but we've got to do the same for compatibility.
1199 if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
1200 D->hasAttr<CUDASharedAttr>() ||
1201 D->getType()->isCUDADeviceBuiltinSurfaceType() ||
1202 D->getType()->isCUDADeviceBuiltinTextureType()) {
1203 Linkage = llvm::GlobalValue::InternalLinkage;
1204 }
1205}
1206
1207void CGNVCUDARuntime::handleVarRegistration(const VarDecl *D,
1208 llvm::GlobalVariable &GV) {
1209 if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
1210 // Shadow variables and their properties must be registered with CUDA
1211 // runtime. Skip Extern global variables, which will be registered in
1212 // the TU where they are defined.
1213 //
1214 // Don't register a C++17 inline variable. The local symbol can be
1215 // discarded and referencing a discarded local symbol from outside the
1216 // comdat (__cuda_register_globals) is disallowed by the ELF spec.
1217 //
1218 // HIP managed variables need to be always recorded in device and host
1219 // compilations for transformation.
1220 //
1221 // HIP managed variables and variables in CUDADeviceVarODRUsedByHost are
1222 // added to llvm.compiler-used, therefore they are safe to be registered.
1223 if ((!D->hasExternalStorage() && !D->isInline()) ||
1224 CGM.getContext().CUDADeviceVarODRUsedByHost.contains(key: D) ||
1225 D->hasAttr<HIPManagedAttr>()) {
1226 registerDeviceVar(VD: D, Var&: GV, Extern: !D->hasDefinition(),
1227 Constant: D->hasAttr<CUDAConstantAttr>());
1228 }
1229 } else if (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
1230 D->getType()->isCUDADeviceBuiltinTextureType()) {
1231 // Builtin surfaces and textures and their template arguments are
1232 // also registered with CUDA runtime.
1233 const auto *TD = cast<ClassTemplateSpecializationDecl>(
1234 Val: D->getType()->castAsCXXRecordDecl());
1235 const TemplateArgumentList &Args = TD->getTemplateArgs();
1236 if (TD->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>()) {
1237 assert(Args.size() == 2 &&
1238 "Unexpected number of template arguments of CUDA device "
1239 "builtin surface type.");
1240 auto SurfType = Args[1].getAsIntegral();
1241 if (!D->hasExternalStorage())
1242 registerDeviceSurf(VD: D, Var&: GV, Extern: !D->hasDefinition(), Type: SurfType.getSExtValue());
1243 } else {
1244 assert(Args.size() == 3 &&
1245 "Unexpected number of template arguments of CUDA device "
1246 "builtin texture type.");
1247 auto TexType = Args[1].getAsIntegral();
1248 auto Normalized = Args[2].getAsIntegral();
1249 if (!D->hasExternalStorage())
1250 registerDeviceTex(VD: D, Var&: GV, Extern: !D->hasDefinition(), Type: TexType.getSExtValue(),
1251 Normalized: Normalized.getZExtValue());
1252 }
1253 }
1254}
1255
1256// Transform managed variables to pointers to managed variables in device code.
1257// Each use of the original managed variable is replaced by a load from the
1258// transformed managed variable. The transformed managed variable contains
1259// the address of managed memory which will be allocated by the runtime.
1260void CGNVCUDARuntime::transformManagedVars() {
1261 for (auto &&Info : DeviceVars) {
1262 llvm::GlobalVariable *Var = Info.Var;
1263 if (Info.Flags.getKind() == DeviceVarFlags::Variable &&
1264 Info.Flags.isManaged()) {
1265 auto *ManagedVar = new llvm::GlobalVariable(
1266 CGM.getModule(), Var->getType(),
1267 /*isConstant=*/false, Var->getLinkage(),
1268 /*Init=*/Var->isDeclaration()
1269 ? nullptr
1270 : llvm::ConstantPointerNull::get(T: Var->getType()),
1271 /*Name=*/"", /*InsertBefore=*/nullptr,
1272 llvm::GlobalVariable::NotThreadLocal,
1273 CGM.getContext().getTargetAddressSpace(AS: CGM.getLangOpts().CUDAIsDevice
1274 ? LangAS::cuda_device
1275 : LangAS::Default));
1276 ManagedVar->setDSOLocal(Var->isDSOLocal());
1277 ManagedVar->setVisibility(Var->getVisibility());
1278 ManagedVar->setExternallyInitialized(true);
1279 replaceManagedVar(Var, ManagedVar);
1280 ManagedVar->takeName(V: Var);
1281 Var->setName(Twine(ManagedVar->getName()) + ".managed");
1282 // Keep managed variables even if they are not used in device code since
1283 // they need to be allocated by the runtime.
1284 if (CGM.getLangOpts().CUDAIsDevice && !Var->isDeclaration()) {
1285 assert(!ManagedVar->isDeclaration());
1286 CGM.addCompilerUsedGlobal(GV: Var);
1287 CGM.addCompilerUsedGlobal(GV: ManagedVar);
1288 }
1289 }
1290 }
1291}
1292
1293// Creates offloading entries for all the kernels and globals that must be
1294// registered. The linker will provide a pointer to this section so we can
1295// register the symbols with the linked device image.
1296void CGNVCUDARuntime::createOffloadingEntries() {
1297 llvm::object::OffloadKind Kind = CGM.getLangOpts().HIP
1298 ? llvm::object::OffloadKind::OFK_HIP
1299 : llvm::object::OffloadKind::OFK_Cuda;
1300
1301 llvm::Module &M = CGM.getModule();
1302 for (KernelInfo &I : EmittedKernels)
1303 llvm::offloading::emitOffloadingEntry(
1304 M, Kind, Addr: KernelHandles[I.Kernel->getName()],
1305 Name: getDeviceSideName(ND: cast<NamedDecl>(Val: I.D)), /*Flags=*/Size: 0, /*Data=*/Flags: 0,
1306 Data: llvm::offloading::OffloadGlobalEntry);
1307
1308 for (VarInfo &I : DeviceVars) {
1309 uint64_t VarSize =
1310 CGM.getDataLayout().getTypeAllocSize(Ty: I.Var->getValueType());
1311 int32_t Flags =
1312 (I.Flags.isExtern()
1313 ? static_cast<int32_t>(llvm::offloading::OffloadGlobalExtern)
1314 : 0) |
1315 (I.Flags.isConstant()
1316 ? static_cast<int32_t>(llvm::offloading::OffloadGlobalConstant)
1317 : 0) |
1318 (I.Flags.isNormalized()
1319 ? static_cast<int32_t>(llvm::offloading::OffloadGlobalNormalized)
1320 : 0);
1321 if (I.Flags.getKind() == DeviceVarFlags::Variable) {
1322 if (I.Flags.isManaged()) {
1323 assert(I.Var->getName().ends_with(".managed") &&
1324 "HIP managed variables not transformed");
1325
1326 auto *ManagedVar = M.getNamedGlobal(
1327 Name: I.Var->getName().drop_back(N: StringRef(".managed").size()));
1328 llvm::offloading::emitOffloadingEntry(
1329 M, Kind, Addr: I.Var, Name: getDeviceSideName(ND: I.D), Size: VarSize,
1330 Flags: llvm::offloading::OffloadGlobalManagedEntry | Flags,
1331 /*Data=*/I.Var->getAlign().valueOrOne().value(), AuxAddr: ManagedVar);
1332 } else {
1333 llvm::offloading::emitOffloadingEntry(
1334 M, Kind, Addr: I.Var, Name: getDeviceSideName(ND: I.D), Size: VarSize,
1335 Flags: llvm::offloading::OffloadGlobalEntry | Flags,
1336 /*Data=*/0);
1337 }
1338 } else if (I.Flags.getKind() == DeviceVarFlags::Surface) {
1339 llvm::offloading::emitOffloadingEntry(
1340 M, Kind, Addr: I.Var, Name: getDeviceSideName(ND: I.D), Size: VarSize,
1341 Flags: llvm::offloading::OffloadGlobalSurfaceEntry | Flags,
1342 Data: I.Flags.getSurfTexType());
1343 } else if (I.Flags.getKind() == DeviceVarFlags::Texture) {
1344 llvm::offloading::emitOffloadingEntry(
1345 M, Kind, Addr: I.Var, Name: getDeviceSideName(ND: I.D), Size: VarSize,
1346 Flags: llvm::offloading::OffloadGlobalTextureEntry | Flags,
1347 Data: I.Flags.getSurfTexType());
1348 }
1349 }
1350
1351 // Register the per-TU offload-profiling shadow. The offloading entry
1352 // makes the linker-wrapper emit the host __hipRegisterVar call in the
1353 // combined ctor. Separately emit a per-TU ctor that registers the
1354 // shadow with the profile runtime's drain list.
1355 if (OffloadProfShadow) {
1356 llvm::offloading::emitOffloadingEntry(
1357 M, Kind, Addr: OffloadProfShadow, Name: OffloadProfShadow->getName(),
1358 Size: CGM.getDataLayout().getPointerSize(/*AS=*/0),
1359 Flags: llvm::offloading::OffloadGlobalEntry, /*Data=*/0);
1360
1361 llvm::LLVMContext &Ctx = M.getContext();
1362 auto *PtrTy = llvm::PointerType::getUnqual(C&: Ctx);
1363 llvm::FunctionCallee RegisterShadow = CGM.CreateRuntimeFunction(
1364 Ty: llvm::FunctionType::get(Result: VoidTy, Params: {PtrTy}, isVarArg: false),
1365 Name: "__llvm_profile_offload_register_shadow_variable");
1366 llvm::FunctionCallee RegisterSectionShadow = CGM.CreateRuntimeFunction(
1367 Ty: llvm::FunctionType::get(Result: VoidTy, Params: {PtrTy}, isVarArg: false),
1368 Name: "__llvm_profile_offload_register_section_shadow_variable");
1369 auto *CtorFn = llvm::Function::Create(
1370 Ty: llvm::FunctionType::get(Result: VoidTy, isVarArg: false),
1371 Linkage: llvm::GlobalValue::InternalLinkage,
1372 N: "__llvm_profile_register_shadow." + CGM.getContext().getCUIDHash(), M: &M);
1373 auto *Entry = llvm::BasicBlock::Create(Context&: Ctx, Name: "entry", Parent: CtorFn);
1374 llvm::IRBuilder<> B(Entry);
1375 B.CreateCall(Callee: RegisterShadow, Args: {OffloadProfShadow});
1376 for (const auto &Info : OffloadProfSectionShadows) {
1377 llvm::offloading::emitOffloadingEntry(
1378 M, Kind, Addr: Info.Shadow, Name: Info.DeviceName,
1379 Size: CGM.getDataLayout().getPointerSize(/*AS=*/0),
1380 Flags: llvm::offloading::OffloadGlobalEntry, /*Data=*/0);
1381 B.CreateCall(Callee: RegisterSectionShadow, Args: {Info.Shadow});
1382 }
1383 B.CreateRetVoid();
1384 llvm::appendToGlobalCtors(M, F: CtorFn, /*Priority=*/65535);
1385 }
1386}
1387
1388// For HIP host+device compiles with PGO enabled, emit the host-side shadow for
1389// the per-TU __llvm_profile_sections_<CUID> global. Device-side section table
1390// emission is owned by InstrProfiling so it can be gated on real profile data.
1391void CGNVCUDARuntime::emitOffloadProfilingSections() {
1392 if (!CGM.getLangOpts().HIP)
1393 return;
1394 if (!CGM.getCodeGenOpts().hasProfileInstr())
1395 return;
1396
1397 StringRef CUIDHash = CGM.getContext().getCUIDHash();
1398 if (CUIDHash.empty())
1399 return;
1400
1401 llvm::Module &M = CGM.getModule();
1402 llvm::LLVMContext &Ctx = M.getContext();
1403 std::string Name = ("__llvm_profile_sections_" + CUIDHash).str();
1404
1405 // If the global already exists (e.g. another TU was merged in), don't
1406 // duplicate it.
1407 if (M.getNamedValue(Name))
1408 return;
1409
1410 if (CGM.getLangOpts().CUDAIsDevice) {
1411 // Device side: emit only the per-TU names postfix marker. The sections
1412 // struct is emitted later by the InstrProfiling pass, which emits it only
1413 // when the TU has profile data, avoiding dangling section references.
1414 unsigned GlobalAS = M.getDataLayout().getDefaultGlobalsAddressSpace();
1415 std::string NamesVarPostfixVarName =
1416 std::string(llvm::getInstrProfNamesVarPostfixVarName());
1417 if (!M.getNamedValue(Name: NamesVarPostfixVarName)) {
1418 auto *NamesVarPostfix = llvm::ConstantDataArray::getString(
1419 Context&: Ctx, Initializer: (llvm::Twine("_") + CUIDHash).str(), AddNull: true);
1420 auto *NamesGV = new llvm::GlobalVariable(
1421 M, NamesVarPostfix->getType(), /*isConstant=*/true,
1422 llvm::GlobalValue::PrivateLinkage, NamesVarPostfix,
1423 NamesVarPostfixVarName,
1424 /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
1425 GlobalAS);
1426 CGM.addCompilerUsedGlobal(GV: NamesGV);
1427 }
1428 return;
1429 }
1430
1431 // Host side: emit an opaque void* shadow. Layout doesn't matter — the
1432 // runtime locates it by name via hipGetSymbolAddress and treats it as
1433 // the address of the device-side struct. Registration with the HIP
1434 // runtime is added by makeRegisterGlobalsFn (non-RDC) or
1435 // createOffloadingEntries (RDC).
1436 auto *PtrTy = llvm::PointerType::getUnqual(C&: Ctx);
1437 OffloadProfShadow = new llvm::GlobalVariable(
1438 M, PtrTy, /*isConstant=*/false, llvm::GlobalValue::ExternalLinkage,
1439 llvm::ConstantPointerNull::get(T: PtrTy), Name);
1440 CGM.addCompilerUsedGlobal(GV: OffloadProfShadow);
1441
1442 auto AddSectionShadow = [&](StringRef Kind, const Twine &DeviceName) {
1443 std::string ShadowName =
1444 (Twine("__llvm_profile_shadow_") + Kind + "_" + CUIDHash + "_" +
1445 Twine(OffloadProfSectionShadows.size()))
1446 .str();
1447 auto *Shadow = new llvm::GlobalVariable(
1448 M, PtrTy, /*isConstant=*/false, llvm::GlobalValue::ExternalLinkage,
1449 llvm::ConstantPointerNull::get(T: PtrTy), ShadowName);
1450 CGM.addCompilerUsedGlobal(GV: Shadow);
1451 OffloadProfSectionShadows.push_back(Elt: {.Shadow: Shadow, .DeviceName: DeviceName.str()});
1452 };
1453
1454 // Keep this order in sync with the runtime: data, counters, uniform counters,
1455 // then names.
1456 for (auto &&I : EmittedKernels) {
1457 std::string KernelName = getDeviceSideName(ND: cast<NamedDecl>(Val: I.D));
1458 AddSectionShadow("data", Twine("__profd_") + KernelName);
1459 AddSectionShadow("cnts", Twine("__profc_") + KernelName);
1460 AddSectionShadow("ucnts", Twine("__llvm_prf_unifcnt_") + KernelName);
1461 AddSectionShadow("names",
1462 Twine(llvm::getInstrProfNamesVarName()) + "_" + CUIDHash);
1463 }
1464}
1465
1466// Returns module constructor to be added.
1467llvm::Function *CGNVCUDARuntime::finalizeModule() {
1468 transformManagedVars();
1469 emitOffloadProfilingSections();
1470 if (CGM.getLangOpts().CUDAIsDevice) {
1471 // Mark ODR-used device variables as compiler used to prevent it from being
1472 // eliminated by optimization. This is necessary for device variables
1473 // ODR-used by host functions. Sema correctly marks them as ODR-used no
1474 // matter whether they are ODR-used by device or host functions.
1475 //
1476 // We do not need to do this if the variable has used attribute since it
1477 // has already been added.
1478 //
1479 // Static device variables have been externalized at this point, therefore
1480 // variables with LLVM private or internal linkage need not be added.
1481 for (auto &&Info : DeviceVars) {
1482 auto Kind = Info.Flags.getKind();
1483 if (!Info.Var->isDeclaration() &&
1484 !llvm::GlobalValue::isLocalLinkage(Linkage: Info.Var->getLinkage()) &&
1485 (Kind == DeviceVarFlags::Variable ||
1486 Kind == DeviceVarFlags::Surface ||
1487 Kind == DeviceVarFlags::Texture) &&
1488 Info.D->isUsed() && !Info.D->hasAttr<UsedAttr>()) {
1489 CGM.addCompilerUsedGlobal(GV: Info.Var);
1490 }
1491 }
1492 return nullptr;
1493 }
1494 if (!CGM.getLangOpts().CUDANVCCABI &&
1495 (CGM.getLangOpts().OffloadViaLLVM || RelocatableDeviceCode))
1496 createOffloadingEntries();
1497 else
1498 return makeModuleCtorFunction();
1499
1500 return nullptr;
1501}
1502
1503llvm::GlobalValue *CGNVCUDARuntime::getKernelHandle(llvm::Function *F,
1504 GlobalDecl GD) {
1505 auto Loc = KernelHandles.find(Val: F->getName());
1506 if (Loc != KernelHandles.end()) {
1507 auto OldHandle = Loc->second;
1508 if (KernelStubs[OldHandle] == F)
1509 return OldHandle;
1510
1511 // We've found the function name, but F itself has changed, so we need to
1512 // update the references.
1513 if (CGM.getLangOpts().HIP) {
1514 // For HIP compilation the handle itself does not change, so we only need
1515 // to update the Stub value.
1516 KernelStubs[OldHandle] = F;
1517 return OldHandle;
1518 }
1519 // For non-HIP compilation, erase the old Stub and fall-through to creating
1520 // new entries.
1521 KernelStubs.erase(Val: OldHandle);
1522 }
1523
1524 if (!CGM.getLangOpts().HIP) {
1525 KernelHandles[F->getName()] = F;
1526 KernelStubs[F] = F;
1527 return F;
1528 }
1529
1530 auto *Var = new llvm::GlobalVariable(
1531 TheModule, F->getType(), /*isConstant=*/true, F->getLinkage(),
1532 /*Initializer=*/nullptr,
1533 CGM.getMangledName(
1534 GD: GD.getWithKernelReferenceKind(Kind: KernelReferenceKind::Kernel)));
1535 Var->setAlignment(CGM.getPointerAlign().getAsAlign());
1536 Var->setVisibility(F->getVisibility());
1537 CGM.setDSOLocal(Var);
1538 auto *FD = cast<FunctionDecl>(Val: GD.getDecl());
1539 auto *FT = FD->getPrimaryTemplate();
1540 if (!FT || FT->isThisDeclarationADefinition())
1541 CGM.maybeSetTrivialComdat(D: *FD, GO&: *Var);
1542 KernelHandles[F->getName()] = Var;
1543 KernelStubs[Var] = F;
1544 return Var;
1545}
1546