[CUDA] Change initializer for CUDA device code based on CUDA documentation.

[wengxt]

Summary

According to CUDA documentation, global variables declared with device,
constant can be initialized from host code, so mark them as
externally initialized. Because shared variables cannot have an
initialization as part of their declaration and since the value maybe kept
across different kernel invocation, the value of shared is effectively
undefined instead of zero initialized.

Wrongly using zero initializer may cause illegitimate optimization, e.g.
removing unused constant variable because it's not updated in the device
code and the value is initialized with zero.

Test Plan: test/CodeGenCUDA/address-spaces.cu

Patch by Xuetian Weng

[jingyue]

Looks good in general.

According to CUDA documentation, global variable with device,

global variables

externally initialized. For shared variable, its value might be kept
by GPU Memory among different kernel invocation and it is not possible
to initialize it. So change shared variable it to undefined value.

Just say: shared variables cannot have an initialization as part of their declaration. Whether being kept across invocations is undefined.

lib/CodeGen/CodeGenModule.cpp

1998

Then, this should be a hard error instead of an unsupported feature, right? ErrorUnsupported sounds like we want to support it later.

[jingyue]

+Eli

[jingyue]

+Artem

[tra]

+1 to jingyue's comments.

lib/CodeGen/CodeGenModule.cpp

1995

please use getLangOpts()

2086–2096

Cosmetic nit: no closing quote marks in both citations.

You may want to leave only one copy if "is accessible..." as it applies to both B.2.1 and B.2.2.

[wengxt]

Just say: shared variables cannot have an initialization as part of their declaration. Whether being kept across invocations is undefined.

My intention is to explain that why shared varible should be undefined instead of zero-initialized. Being unable to have initialization as declaration does not assume that it may contain random value.

[jingyue]

OK. Then some rewording needs to be done.

The reason shared variables should be initialized as undef is that CUDA's language manual explicitly says that it can't be initialized.

The reason they shouldn't be initialized as zero is sort of what you explained in the second paragraph: compilers might perform optimizations that leverage the zero initializer. However, such optimizations are illegimate because, at runtime, data previously stored to shared memory might be left over.

Is this more accurate?

[wengxt]

update based on comments.

[wengxt]

make shared with initialization a hard error.

[jingyue]

LGTM

[tra]

Following code compiles with nvcc, but errors out with clang and this patch due to InitExpr check in CodeGenModule.cpp. It looks like the check needs to be more selective.

struct c {
  c() {}
};

__shared__ c var;

[tra]

In D12241#236830, @tra wrote:

Following code compiles with nvcc, but errors out with clang and this patch due to InitExpr check in CodeGenModule.cpp. It looks like the check needs to be more selective.

struct c {
  c() {}
};

__shared__ c var;

Looks like we don't even need explicit constructor to fail:

struct c { };
__shared__ c var;

I believe it's something that we *do* want to support.

lib/CodeGen/CodeGenModule.cpp

void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D) {
   QualType ASTTy = D->getType();
   CXXRecordDecl *RD = ASTTy->getBaseElementTypeUnsafe()->getAsCXXRecordDecl();
   bool NeedsGlobalCtor = false;
   bool NeedsGlobalDtor = RD && !RD->hasTrivialDestructor();
 
   const VarDecl *InitDecl;
   const Expr *InitExpr = D->getAnyInitializer(InitDecl);
 
-  if (!InitExpr) {
+  // CUDA E.2.4.1 "__shared__ variables cannot have an initialization as part
+  // of their declaration."
+  if (getLangOpts().CPlusPlus && getLangOpts().CUDAIsDevice
+      && D->hasAttr<CUDASharedAttr>()) {
+    if (InitExpr)
+      ErrorUnsupported(D, "static initializer");
+    Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
+  } else if (!InitExpr) {
     // This is a tentative definition; tentative definitions are
     // implicitly initialized with { 0 }.
     //
     // Note that tentative definitions are only emitted at the end of
     // a translation unit, so they should never have incomplete
     // type. In addition, EmitTentativeDefinition makes sure that we
     // never attempt to emit a tentative definition if a real one
     // exists. A use may still exists, however, so we still may need
if (!GV ||
     cast<llvm::GlobalValue>(Entry)->eraseFromParent();
   }
 
   MaybeHandleStaticInExternC(D, GV);
 
   if (D->hasAttr<AnnotateAttr>())
     AddGlobalAnnotations(D, GV);
 
+  // CUDA B.2.1 "The __device__ qualifier declares a variable that resides on
+  // the device. [...]"
+  // CUDA B.2.2 "The __constant__ qualifier, optionally used together with
+  // __device__, declares a variable that: [...]
+  // Is accessible from all the threads within the grid and from the host
+  // through the runtime library (cudaGetSymbolAddress() / cudaGetSymbolSize()
+  // / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol())."
+  if (GV && LangOpts.CUDA && LangOpts.CUDAIsDevice &&
+      (D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>())) {
+    GV->setExternallyInitialized(true);
+  }
   GV->setInitializer(Init);
 
   // If it is safe to mark the global 'constant', do so now.
   GV->setConstant(!NeedsGlobalCtor && !NeedsGlobalDtor &&
                   isTypeConstant(D->getType(), true));
 
   // If it is in a read-only section, mark it 'constant'.
   if (const SectionAttr *SA = D->getAttr<SectionAttr>()) {

test/CodeGenCUDA/address-spaces.cu

 // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
 
 // Verifies Clang emits correct address spaces and addrspacecast instructions
 // for CUDA code.
 
 #include "Inputs/cuda.h"
 
-// CHECK: @i = addrspace(1) global
+// CHECK: @i = addrspace(1) externally_initialized global
 __device__ int i;
 
-// CHECK: @j = addrspace(4) global
+// CHECK: @j = addrspace(4) externally_initialized global
 __constant__ int j;
 
 // CHECK: @k = addrspace(3) global
 __shared__ int k;
 
 struct MyStruct {
   int data1;
   int data2;
 };
 
 // CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct zeroinitializer
 // CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] zeroinitializer
 // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00
-// CHECK: @b = addrspace(3) global float 0.000000e+00
+// CHECK: @b = addrspace(3) global float undef
 
 __device__ void foo() {
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
   i++;
 
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @j to i32*)
   j++;
 

test/CodeGenCUDA/filter-decl.cu

 // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - | FileCheck -check-prefix=CHECK-HOST %s
 // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefix=CHECK-DEVICE %s
 
 #include "Inputs/cuda.h"
 
 // This has to be at the top of the file as that's where file-scope
 // asm ends up.
 // CHECK-HOST: module asm "file scope asm is host only"
 // CHECK-DEVICE-NOT: module asm "file scope asm is host only"
 __asm__("file scope asm is host only");
 
-// CHECK-HOST-NOT: constantdata = global
-// CHECK-DEVICE: constantdata = global
+// CHECK-HOST-NOT: constantdata = externally_initialized global
+// CHECK-DEVICE: constantdata = externally_initialized global
 __constant__ char constantdata[256];
 
-// CHECK-HOST-NOT: devicedata = global
-// CHECK-DEVICE: devicedata = global
+// CHECK-HOST-NOT: devicedata = externally_initialized global
+// CHECK-DEVICE: devicedata = externally_initialized global
 __device__ char devicedata[256];
 
 // CHECK-HOST-NOT: shareddata = global
 // CHECK-DEVICE: shareddata = global
 __shared__ char shareddata[256];
 
 // CHECK-HOST: hostdata = global
 // CHECK-DEVICE-NOT: hostdata = global