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