[CUDA] Register relocatable GPU binaries

[Hahnfeld]

Summary

nvcc generates a unique registration function for each object file
that contains relocatable device code. Unique names are achieved
with a module id that is also reflected in the function's name.

[Hahnfeld]

I didn't write tests for this yet, but I wanted to get some early feedback on this and show what I have in mind.

lib/CodeGen/CGCUDANV.cpp

333–334

Can we actually have multiple GPU binaries here? If yes, how do I get there?

334–335

@jlebar The same here, probably __NV_CUDA,__nv_module_id?

335–336

@jlebar Could yo help me here as I don't have a Mac? I'd guess it's __NV_CUDA,__nv_relfatbin but I'd feel better if I can get a confirmation...

[Hahnfeld]

Rebase and fix Debug build.

[Hahnfeld]

Still no regression tests.

I did some functional tests though (https://reviews.llvm.org/F5822023): With this patch Clang can generate valid object files with relocatable device code. For linking I still defer to nvcc and I'm not sure if I'm interested in reverse-engineering the needed tools to make this fully work with Clang's Driver: I think the biggest advantage of CUDA in Clang is using LLVM's CodeGen. Note that (in my simple tests) Clang's object files had enough compatibility to mix them with other objects generated by nvcc (see Makefile.mixed)!

[tra]

lib/CodeGen/CGCUDANV.cpp

333–334

Yes. clang --cuda-gpu-arch=sm_35 --cuda-gpu-arch=sm_50... will compile for sm_35 and sm_50 and then will pass the names of GPU-side objects to the host compilation via -fcuda-include-gpubinary.

[Hahnfeld]

lib/CodeGen/CGCUDANV.cpp

333–334

I'm not sure if that's true anymore: I think they are now combined by fatbinary. This seems to be confirmed by test/Driver/cuda-options.cu. If that was the only use case, we may try to get rid of this possibility, let me check this.

[tra]

lib/CodeGen/CGCUDANV.cpp

333–334

You are correct. All GPU binaries are in the single fatbin now.
That said, you could still pass extra -fcuda-include-gpubinary to cc1 manually, but I see no practical reason to do it -- single fatbin serves the purpose better.

We should remove this loop and make CGM.getCodeGenOpts().CudaGpuBinaryFileNames a scalar.

[Hahnfeld]

lib/CodeGen/CGCUDANV.cpp

333–334

Ok, I'll work on this as a preparation patch and rebase this on top. That actually explains why my changes have always been working even though it didn't handle the loop correctly :-)

[Hahnfeld]

Sorry for the long delay. This update rebases the patch against current trunk and adapts the regression test.

[Hahnfeld]

Correct test check prefix.

[Hahnfeld]

Ping

[tra]

lib/CodeGen/CGCUDANV.cpp

298–300

Instead of tracking these through the conditionals of pretty long function, could we make these pointers class fields and init them in the constructor and make accessors return them and, possibly, assert that they are used if RDC is enabled?

test/CodeGenCUDA/device-stub.cu

2–6

Labels could be a bit more descriptive:
CHECK -> ALL
DEFAULT -> NORDC

Long RUN lines could use some re-wrapping.

[Hahnfeld]

Move FunctionTypes to methods and change test prefixes.

lib/CodeGen/CGCUDANV.cpp

298–300

I've removed the caching entirely because that's already done by llvm::FunctionType::get(). These are now called in new methods to avoid duplication.

[tra]

lib/CodeGen/CGCUDANV.cpp

364–407

This can all be folded into the 'else' branch of the 'if' below.

[Hahnfeld]

Move module ID to corresponding else branch.

lib/CodeGen/CGCUDANV.cpp

 // This provides a class for CUDA code generation targeting the NVIDIA CUDA
 // runtime library.
 //
 //===----------------------------------------------------------------------===//
 
 #include "CGCUDARuntime.h"
 #include "CodeGenFunction.h"
 #include "CodeGenModule.h"
-#include "clang/CodeGen/ConstantInitBuilder.h"
 #include "clang/AST/Decl.h"
+#include "clang/CodeGen/ConstantInitBuilder.h"
 #include "llvm/IR/BasicBlock.h"
 #include "llvm/IR/CallSite.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/DerivedTypes.h"
+#include "llvm/Support/Format.h"
 
 using namespace clang;
 using namespace CodeGen;
 
 namespace {
 
 class CGNVCUDARuntime : public CGCUDARuntime {
 
 private:
   llvm::IntegerType *IntTy, *SizeTy;
   llvm::Type *VoidTy;
   llvm::PointerType *CharPtrTy, *VoidPtrTy, *VoidPtrPtrTy;
 
   /// Convenience reference to LLVM Context
   llvm::LLVMContext &Context;
   /// Convenience reference to the current module
   llvm::Module &TheModule;
   /// Keeps track of kernel launch stubs emitted in this module
   llvm::SmallVector<llvm::Function *, 16> EmittedKernels;
   llvm::SmallVector<std::pair<llvm::GlobalVariable *, unsigned>, 16> DeviceVars;
   /// Keeps track of variable containing handle of GPU binary. Populated by
   /// ModuleCtorFunction() and used to create corresponding cleanup calls in
   /// ModuleDtorFunction()
   llvm::GlobalVariable *GpuBinaryHandle = nullptr;
+  /// Whether we generate relocatable device code.
+  bool RelocatableDeviceCode;
 
   llvm::Constant *getSetupArgumentFn() const;
   llvm::Constant *getLaunchFn() const;
 
+  llvm::FunctionType *getRegisterGlobalsFnTy() const;
+  llvm::FunctionType *getCallbackFnTy() const;
+  llvm::FunctionType *getRegisterLinkedBinaryFnTy() const;
+
   /// Creates a function to register all kernel stubs generated in this module.
   llvm::Function *makeRegisterGlobalsFn();
 
   /// Helper function that generates a constant string and returns a pointer to
   /// the start of the string.  The result of this function can be used anywhere
   /// where the C code specifies const char*.
   llvm::Constant *makeConstantString(const std::string &Str,
                                      const std::string &Name = "",
                                      const std::string &SectionName = "",
                                      unsigned Alignment = 0) {
     llvm::Constant *Zeros[] = {llvm::ConstantInt::get(SizeTy, 0),
                                llvm::ConstantInt::get(SizeTy, 0)};
     auto ConstStr = CGM.GetAddrOfConstantCString(Str, Name.c_str());
     llvm::GlobalVariable *GV =
         cast<llvm::GlobalVariable>(ConstStr.getPointer());
     if (!SectionName.empty())
       GV->setSection(SectionName);
     if (Alignment)
       GV->setAlignment(Alignment);
 
     return llvm::ConstantExpr::getGetElementPtr(ConstStr.getElementType(),
                                                 ConstStr.getPointer(), Zeros);
   }
 
+  /// Helper function that generates an empty dummy function returning void.
+  llvm::Function *makeDummyFunction(llvm::FunctionType *FnTy) {
+    assert(FnTy->getReturnType()->isVoidTy() &&
+           "Can only generate dummy functions returning void!");
+    llvm::Function *DummyFunc = llvm::Function::Create(
+        FnTy, llvm::GlobalValue::InternalLinkage, "dummy", &TheModule);
+
+    llvm::BasicBlock *DummyBlock =
+        llvm::BasicBlock::Create(Context, "", DummyFunc);
+    CGBuilderTy FuncBuilder(CGM, Context);
+    FuncBuilder.SetInsertPoint(DummyBlock);
+    FuncBuilder.CreateRetVoid();
+
+    return DummyFunc;
+  }
+
   void emitDeviceStubBody(CodeGenFunction &CGF, FunctionArgList &Args);
 
 public:
   CGNVCUDARuntime(CodeGenModule &CGM);
 
   void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) override;
   void registerDeviceVar(llvm::GlobalVariable &Var, unsigned Flags) override {
     DeviceVars.push_back(std::make_pair(&Var, Flags));
   }
 
   /// Creates module constructor function
   llvm::Function *makeModuleCtorFunction() override;
   /// Creates module destructor function
   llvm::Function *makeModuleDtorFunction() override;
 };
 
 }
 
 CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
     : CGCUDARuntime(CGM), Context(CGM.getLLVMContext()),
-      TheModule(CGM.getModule()) {
+      TheModule(CGM.getModule()),
+      RelocatableDeviceCode(CGM.getLangOpts().CUDARelocatableDeviceCode) {
   CodeGen::CodeGenTypes &Types = CGM.getTypes();
   ASTContext &Ctx = CGM.getContext();
 
   IntTy = CGM.IntTy;
   SizeTy = CGM.SizeTy;
   VoidTy = CGM.VoidTy;
 
   CharPtrTy = llvm::PointerType::getUnqual(Types.ConvertType(Ctx.CharTy));

 }
 
 llvm::Constant *CGNVCUDARuntime::getLaunchFn() const {
   // cudaError_t cudaLaunch(char *)
   return CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(IntTy, CharPtrTy, false), "cudaLaunch");
 }
 
+llvm::FunctionType *CGNVCUDARuntime::getRegisterGlobalsFnTy() const {
+  return llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false);
+}
+
+llvm::FunctionType *CGNVCUDARuntime::getCallbackFnTy() const {
+  return llvm::FunctionType::get(VoidTy, VoidPtrTy, false);
+}
+
+llvm::FunctionType *CGNVCUDARuntime::getRegisterLinkedBinaryFnTy() const {
+  auto CallbackFnTy = getCallbackFnTy();
+  auto RegisterGlobalsFnTy = getRegisterGlobalsFnTy();
+  llvm::Type *Params[] = {RegisterGlobalsFnTy->getPointerTo(), VoidPtrTy,
+                          VoidPtrTy, CallbackFnTy->getPointerTo()};
+  return llvm::FunctionType::get(VoidTy, Params, false);
+}
+
 void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
                                      FunctionArgList &Args) {
   EmittedKernels.push_back(CGF.CurFn);
   emitDeviceStubBody(CGF, Args);
 }
 
 void CGNVCUDARuntime::emitDeviceStubBody(CodeGenFunction &CGF,
                                          FunctionArgList &Args) {

 /// }
 /// \endcode
 llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
   // No need to register anything
   if (EmittedKernels.empty() && DeviceVars.empty())
     return nullptr;
 
   llvm::Function *RegisterKernelsFunc = llvm::Function::Create(
-      llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false),
-      llvm::GlobalValue::InternalLinkage, "__cuda_register_globals", &TheModule);
+      getRegisterGlobalsFnTy(), llvm::GlobalValue::InternalLinkage,
+      "__cuda_register_globals", &TheModule);
   llvm::BasicBlock *EntryBB =
       llvm::BasicBlock::Create(Context, "entry", RegisterKernelsFunc);
   CGBuilderTy Builder(CGM, Context);
   Builder.SetInsertPoint(EntryBB);
 
   // void __cudaRegisterFunction(void **, const char *, char *, const char *,
   //                             int, uint3*, uint3*, dim3*, dim3*, int*)
   llvm::Type *RegisterFuncParams[] = {

 /// }
 /// \endcode
 llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
   // No need to generate ctors/dtors if there is no GPU binary.
   std::string GpuBinaryFileName = CGM.getCodeGenOpts().CudaGpuBinaryFileName;
   if (GpuBinaryFileName.empty())
     return nullptr;
 
   // void __cuda_register_globals(void* handle);
   llvm::Function *RegisterGlobalsFunc = makeRegisterGlobalsFn();
+  // We always need a function to pass in as callback. Create a dummy
+  // implementation if we don't need to register anything.
+  if (RelocatableDeviceCode && !RegisterGlobalsFunc)
+    RegisterGlobalsFunc = makeDummyFunction(getRegisterGlobalsFnTy());
+
   // void ** __cudaRegisterFatBinary(void *);
   llvm::Constant *RegisterFatbinFunc = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(VoidPtrPtrTy, VoidPtrTy, false),
       "__cudaRegisterFatBinary");
   // struct { int magic, int version, void * gpu_binary, void * dont_care };
   llvm::StructType *FatbinWrapperTy =
       llvm::StructType::get(IntTy, IntTy, VoidPtrTy, VoidPtrTy);
 
llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
   llvm::Function *ModuleCtorFunc = llvm::Function::Create(
       llvm::FunctionType::get(VoidTy, VoidPtrTy, false),
       llvm::GlobalValue::InternalLinkage, "__cuda_module_ctor", &TheModule);
   llvm::BasicBlock *CtorEntryBB =
       llvm::BasicBlock::Create(Context, "entry", ModuleCtorFunc);
   CGBuilderTy CtorBuilder(CGM, Context);
 
   CtorBuilder.SetInsertPoint(CtorEntryBB);
 
-  const char *FatbinConstantName =
+  const char *FatbinConstantName;
+  if (RelocatableDeviceCode)
+    // TODO: Figure out how this is called on mac OS!
+    FatbinConstantName = "__nv_relfatbin";
+  else
+    FatbinConstantName =
         CGM.getTriple().isMacOSX() ? "__NV_CUDA,__nv_fatbin" : ".nv_fatbin";
   // NVIDIA's cuobjdump looks for fatbins in this section.
   const char *FatbinSectionName =
       CGM.getTriple().isMacOSX() ? "__NV_CUDA,__fatbin" : ".nvFatBinSegment";
+  // TODO: Figure out how this is called on mac OS!
+  const char *NVModuleIDSectionName = "__nv_module_id";
 
   // Create initialized wrapper structure that points to the loaded GPU binary
   ConstantInitBuilder Builder(CGM);
   auto Values = Builder.beginStruct(FatbinWrapperTy);
   // Fatbin wrapper magic.
   Values.addInt(IntTy, 0x466243b1);
   // Fatbin version.
   Values.addInt(IntTy, 1);
   // Data.
   Values.add(makeConstantString(GpuBinaryOrErr.get()->getBuffer(), "",
                                 FatbinConstantName, 8));
   // Unused in fatbin v1.
   Values.add(llvm::ConstantPointerNull::get(VoidPtrTy));
   llvm::GlobalVariable *FatbinWrapper = Values.finishAndCreateGlobal(
       "__cuda_fatbin_wrapper", CGM.getPointerAlign(),
       /*constant*/ true);
   FatbinWrapper->setSection(FatbinSectionName);
 
+  // Register binary with CUDA runtime. This is substantially different in
+  // default mode vs. separate compilation!
+  if (!RelocatableDeviceCode) {
     // GpuBinaryHandle = __cudaRegisterFatBinary(&FatbinWrapper);
     llvm::CallInst *RegisterFatbinCall = CtorBuilder.CreateCall(
-      RegisterFatbinFunc, CtorBuilder.CreateBitCast(FatbinWrapper, VoidPtrTy));
+        RegisterFatbinFunc,
+        CtorBuilder.CreateBitCast(FatbinWrapper, VoidPtrTy));
     GpuBinaryHandle = new llvm::GlobalVariable(
         TheModule, VoidPtrPtrTy, false, llvm::GlobalValue::InternalLinkage,
         llvm::ConstantPointerNull::get(VoidPtrPtrTy), "__cuda_gpubin_handle");
     CtorBuilder.CreateAlignedStore(RegisterFatbinCall, GpuBinaryHandle,
                                    CGM.getPointerAlign());
 
     // Call __cuda_register_globals(GpuBinaryHandle);
     if (RegisterGlobalsFunc)
       CtorBuilder.CreateCall(RegisterGlobalsFunc, RegisterFatbinCall);
+  } else {
+    // Generate a unique module ID.
+    SmallString<64> NVModuleID;
+    llvm::raw_svector_ostream OS(NVModuleID);
+    OS << "__nv_" << llvm::format("%x", FatbinWrapper->getGUID());
+    llvm::Constant *NVModuleIDConstant =
+        makeConstantString(NVModuleID.str(), "", NVModuleIDSectionName, 32);
+
+    // Create an alias for the FatbinWrapper that nvcc will look for.
+    llvm::GlobalAlias::create(llvm::GlobalValue::ExternalLinkage,
+                              Twine("__fatbinwrap") + NVModuleID,
+                              FatbinWrapper);
+
+    // void __cudaRegisterLinkedBinary%NVModuleID%(void (*)(void *), void *,
+    // void *, void (*)(void **))
+    SmallString<128> RegisterLinkedBinaryName("__cudaRegisterLinkedBinary");
+    RegisterLinkedBinaryName += NVModuleID;
+    llvm::Constant *RegisterLinkedBinaryFunc = CGM.CreateRuntimeFunction(
+        getRegisterLinkedBinaryFnTy(), RegisterLinkedBinaryName);
+
+    assert(RegisterGlobalsFunc && "Expecting at least dummy function!");
+    llvm::Value *Args[] = {RegisterGlobalsFunc,
+                           CtorBuilder.CreateBitCast(FatbinWrapper, VoidPtrTy),
+                           NVModuleIDConstant,
+                           makeDummyFunction(getCallbackFnTy())};
+    CtorBuilder.CreateCall(RegisterLinkedBinaryFunc, Args);
+  }
 
   CtorBuilder.CreateRetVoid();
   return ModuleCtorFunc;
 }
 
 /// Creates a global destructor function that unregisters the GPU code blob
 /// registered by constructor.
 /// \code
 /// void __cuda_module_dtor(void*) {

test/CodeGenCUDA/device-stub.cu

 // RUN: echo "GPU binary would be here" > %t
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o - | FileCheck %s
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o -  -DNOGLOBALS \
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN:     -fcuda-include-gpubinary %t -o - \
+// RUN:   | FileCheck %s --check-prefixes=ALL,NORDC
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN:     -fcuda-include-gpubinary %t -o -  -DNOGLOBALS \
 // RUN:   | FileCheck %s -check-prefix=NOGLOBALS
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefix=NOGPUBIN
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN:     -fcuda-rdc -fcuda-include-gpubinary %t -o - \
+// RUN:   | FileCheck %s --check-prefixes=ALL,RDC
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - \
+// RUN:   | FileCheck %s -check-prefix=NOGPUBIN
 
 #include "Inputs/cuda.h"
 
 #ifndef NOGLOBALS
-// CHECK-DAG: @device_var = internal global i32
+// ALL-DAG: @device_var = internal global i32
 __device__ int device_var;
 
-// CHECK-DAG: @constant_var = internal global i32
+// ALL-DAG: @constant_var = internal global i32
 __constant__ int constant_var;
 
-// CHECK-DAG: @shared_var = internal global i32
+// ALL-DAG: @shared_var = internal global i32
 __shared__ int shared_var;
 
 // Make sure host globals don't get internalized...
-// CHECK-DAG: @host_var = global i32
+// ALL-DAG: @host_var = global i32
 int host_var;
 // ... and that extern vars remain external.
-// CHECK-DAG: @ext_host_var = external global i32
+// ALL-DAG: @ext_host_var = external global i32
 extern int ext_host_var;
 
 // Shadows for external device-side variables are *definitions* of
 // those variables.
-// CHECK-DAG: @ext_device_var = internal global i32
+// ALL-DAG: @ext_device_var = internal global i32
 extern __device__ int ext_device_var;
-// CHECK-DAG: @ext_device_var = internal global i32
+// ALL-DAG: @ext_device_var = internal global i32
 extern __constant__ int ext_constant_var;
 
 void use_pointers() {
   int *p;
   p = &device_var;
   p = &constant_var;
   p = &shared_var;
   p = &host_var;
   p = &ext_device_var;
   p = &ext_constant_var;
   p = &ext_host_var;
 }
 
 // Make sure that all parts of GPU code init/cleanup are there:
 // * constant unnamed string with the kernel name
-// CHECK: private unnamed_addr constant{{.*}}kernelfunc{{.*}}\00"
+// ALL: private unnamed_addr constant{{.*}}kernelfunc{{.*}}\00"
 // * constant unnamed string with GPU binary
-// CHECK: private unnamed_addr constant{{.*GPU binary would be here.*}}\00"
-// CHECK-SAME: section ".nv_fatbin", align 8
+// ALL: private unnamed_addr constant{{.*GPU binary would be here.*}}\00"
+// NORDC-SAME: section ".nv_fatbin", align 8
+// RDC-SAME: section "__nv_relfatbin", align 8
 // * constant struct that wraps GPU binary
-// CHECK: @__cuda_fatbin_wrapper = internal constant { i32, i32, i8*, i8* } 
-// CHECK-SAME: { i32 1180844977, i32 1, {{.*}}, i8* null }
-// CHECK-SAME: section ".nvFatBinSegment"
+// ALL: @__cuda_fatbin_wrapper = internal constant { i32, i32, i8*, i8* } 
+// ALL-SAME: { i32 1180844977, i32 1, {{.*}}, i8* null }
+// ALL-SAME: section ".nvFatBinSegment"
 // * variable to save GPU binary handle after initialization
-// CHECK: @__cuda_gpubin_handle = internal global i8** null
-// * Make sure our constructor/destructor was added to global ctor/dtor list.
-// CHECK: @llvm.global_ctors = appending global {{.*}}@__cuda_module_ctor
-// CHECK: @llvm.global_dtors = appending global {{.*}}@__cuda_module_dtor
+// NORDC: @__cuda_gpubin_handle = internal global i8** null
+// * constant unnamed string with NVModuleID
+// RDC: [[MODULE_ID_GLOBAL:@.*]] = private unnamed_addr constant
+// RDC-SAME: c"[[MODULE_ID:.+]]\00", section "__nv_module_id", align 32
+// * Make sure our constructor was added to global ctor list.
+// ALL: @llvm.global_ctors = appending global {{.*}}@__cuda_module_ctor
+// * In separate mode we also register a destructor.
+// NORDC: @llvm.global_dtors = appending global {{.*}}@__cuda_module_dtor
+// * Alias to global symbol containing the NVModuleID.
+// RDC: @__fatbinwrap[[MODULE_ID]] = alias { i32, i32, i8*, i8* }
+// RDC-SAME: { i32, i32, i8*, i8* }* @__cuda_fatbin_wrapper
 
 // Test that we build the correct number of calls to cudaSetupArgument followed
 // by a call to cudaLaunch.
 
-// CHECK: define{{.*}}kernelfunc
-// CHECK: call{{.*}}cudaSetupArgument
-// CHECK: call{{.*}}cudaSetupArgument
-// CHECK: call{{.*}}cudaSetupArgument
-// CHECK: call{{.*}}cudaLaunch
+// ALL: define{{.*}}kernelfunc
+// ALL: call{{.*}}cudaSetupArgument
+// ALL: call{{.*}}cudaSetupArgument
+// ALL: call{{.*}}cudaSetupArgument
+// ALL: call{{.*}}cudaLaunch
 __global__ void kernelfunc(int i, int j, int k) {}
 
 // Test that we've built correct kernel launch sequence.
-// CHECK: define{{.*}}hostfunc
-// CHECK: call{{.*}}cudaConfigureCall
-// CHECK: call{{.*}}kernelfunc
+// ALL: define{{.*}}hostfunc
+// ALL: call{{.*}}cudaConfigureCall
+// ALL: call{{.*}}kernelfunc
 void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
 #endif
 
 // Test that we've built a function to register kernels and global vars.
-// CHECK: define internal void @__cuda_register_globals
-// CHECK: call{{.*}}cudaRegisterFunction(i8** %0, {{.*}}kernelfunc
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0
-// CHECK: ret void
-
-// Test that we've built constructor..
-// CHECK: define internal void @__cuda_module_ctor
-//   .. that calls __cudaRegisterFatBinary(&__cuda_fatbin_wrapper)
-// CHECK: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
+// ALL: define internal void @__cuda_register_globals
+// ALL: call{{.*}}cudaRegisterFunction(i8** %0, {{.*}}kernelfunc
+// ALL-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0
+// ALL-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0
+// ALL-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0
+// ALL-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0
+// ALL: ret void
+
+// Test that we've built a constructor.
+// ALL: define internal void @__cuda_module_ctor
+
+// In separate mode it calls __cudaRegisterFatBinary(&__cuda_fatbin_wrapper)
+// NORDC: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
 //   .. stores return value in __cuda_gpubin_handle
-// CHECK-NEXT: store{{.*}}__cuda_gpubin_handle
+// NORDC-NEXT: store{{.*}}__cuda_gpubin_handle
 //   .. and then calls __cuda_register_globals
-// CHECK-NEXT: call void @__cuda_register_globals
+// NORDC-NEXT: call void @__cuda_register_globals
+
+// With relocatable device code we call __cudaRegisterLinkedBinary%NVModuleID%
+// RDC: call{{.*}}__cudaRegisterLinkedBinary[[MODULE_ID]](
+// RDC-SAME: __cuda_register_globals, {{.*}}__cuda_fatbin_wrapper
+// RDC-SAME: [[MODULE_ID_GLOBAL]]
 
 // Test that we've created destructor.
-// CHECK: define internal void @__cuda_module_dtor
-// CHECK: load{{.*}}__cuda_gpubin_handle
-// CHECK-NEXT: call void @__cudaUnregisterFatBinary
+// NORDC: define internal void @__cuda_module_dtor
+// NORDC: load{{.*}}__cuda_gpubin_handle
+// NORDC-NEXT: call void @__cudaUnregisterFatBinary
 
 // There should be no __cuda_register_globals if we have no
 // device-side globals, but we still need to register GPU binary.
 // Skip GPU binary string first.
 // NOGLOBALS: @0 = private unnamed_addr constant{{.*}}
 // NOGLOBALS-NOT: define internal void @__cuda_register_globals
 // NOGLOBALS: define internal void @__cuda_module_ctor
 // NOGLOBALS: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
 // NOGLOBALS-NOT: call void @__cuda_register_globals
 // NOGLOBALS: define internal void @__cuda_module_dtor
 // NOGLOBALS: call void @__cudaUnregisterFatBinary
 
 // There should be no constructors/destructors if we have no GPU binary.
 // NOGPUBIN-NOT: define internal void @__cuda_register_globals
 // NOGPUBIN-NOT: define internal void @__cuda_module_ctor
 // NOGPUBIN-NOT: define internal void @__cuda_module_dtor