[CUDA] Let device-side shared variables be initialized with undef

[yaxunl]

Summary

CUDA shared variable should be initialized with undef.

Patch by Greg Rodgers.
Revised and lit test added by Yaxun Liu.

[tra]

test/CodeGenCUDA/device-var-init.cu

112 ↗

(On Diff #140095)

Hmm. shared should not be initialized in NVPTX either. This looks like a bug in NVPTX.

For now you should make shared uninitialized regardless of whether we're compiling for AMDGCN or NVPTX.

[yaxunl]

Revised by Artem's comments.

[tra]

test/CodeGenCUDA/device-var-init.cu

121 ↗

(On Diff #140110)

Perhaps it would make sense to capture there names for NVPTX as well and avoid duplicating all the checks below.

[yaxunl]

test/CodeGenCUDA/device-var-init.cu

121 ↗

(On Diff #140110)

will do when committing.

[rjmccall]

What exactly are you trying to express here? Are you just trying to make these external declarations when compiling for the device because __shared__ variables are actually defined on the host? That should be handled by the frontend by setting up the AST so that these declarations are not definitions.

[tra]

In D44985#1050670, @rjmccall wrote:

What exactly are you trying to express here? Are you just trying to make these external declarations when compiling for the device because __shared__ variables are actually defined on the host? That should be handled by the frontend by setting up the AST so that these declarations are not definitions.

shared vars (at least in CUDA) are weird. Local-scoped ones are implicitly static (which compiler will attempt to zero-init) but in CUDA shared variables can't have static initializers and we don't know the value of such vars when we launch the kernel.

[yaxunl]

In D44985#1050670, @rjmccall wrote:

What exactly are you trying to express here? Are you just trying to make these external declarations when compiling for the device because __shared__ variables are actually defined on the host? That should be handled by the frontend by setting up the AST so that these declarations are not definitions.

No. These variables are not like external symbols defined on the host. They behave like global variables in the kernel code but never initialized. Currently no targets are able to initialize them and it is users' responsibility to initialize them explicitly.

Giving them an initial value will cause error in some backends since they cannot handle them, therefore put undef as initializer.

[rjmccall]

In D44985#1050674, @yaxunl wrote:

In D44985#1050670, @rjmccall wrote:

What exactly are you trying to express here? Are you just trying to make these external declarations when compiling for the device because __shared__ variables are actually defined on the host? That should be handled by the frontend by setting up the AST so that these declarations are not definitions.

No. These variables are not like external symbols defined on the host. They behave like global variables in the kernel code but never initialized. Currently no targets are able to initialize them and it is users' responsibility to initialize them explicitly.

Giving them an initial value will cause error in some backends since they cannot handle them, therefore put undef as initializer.

So undef is being used as a special marker to the backends that it's okay not to try to initialize these variables?

[yaxunl]

In D44985#1050876, @rjmccall wrote:

In D44985#1050674, @yaxunl wrote:

In D44985#1050670, @rjmccall wrote:

What exactly are you trying to express here? Are you just trying to make these external declarations when compiling for the device because __shared__ variables are actually defined on the host? That should be handled by the frontend by setting up the AST so that these declarations are not definitions.

No. These variables are not like external symbols defined on the host. They behave like global variables in the kernel code but never initialized. Currently no targets are able to initialize them and it is users' responsibility to initialize them explicitly.

Giving them an initial value will cause error in some backends since they cannot handle them, therefore put undef as initializer.

So undef is being used as a special marker to the backends that it's okay not to try to initialize these variables?

I think undef as the initializer tells the llvm passes and backend that this global variable contains undefined value. I am not sure if this is better than without an initializer. I saw code in CodeGenModule::getOrCreateStaticVarDecl

// Local address space cannot have an initializer.
llvm::Constant *Init = nullptr;
if (Ty.getAddressSpace() != LangAS::opencl_local)
  Init = EmitNullConstant(Ty);
else
  Init = llvm::UndefValue::get(LTy);

which means OpenCL static variable in local address space (equivalent to CUDA shared address space) gets an undef initializer.

For CUDA shared variable, in CodeGenFunction::EmitStaticVarDecl, it first goes through call of CodeGenModule::getOrCreateStaticVarDecl and gets a zeroinitializer, then it reaches line 400

// Whatever initializer such variable may have when it gets here is
  // a no-op and should not be emitted.
  bool isCudaSharedVar = getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
                         D.hasAttr<CUDASharedAttr>();
  // If this value has an initializer, emit it.
  if (D.getInit() && !isCudaSharedVar)
    var = AddInitializerToStaticVarDecl(D, var);

Although this disables adding initializer from D, var already has a zeroinitializer from CodeGenModule::getOrCreateStaticVarDecl, therefore its initializer needs to be overwritten by undef.

Probably a better solution would be do it in CodeGenModule::getOrCreateStaticVarDecl, side by side by the OpenCL code.

[rjmccall]

In D44985#1051840, @yaxunl wrote:

In D44985#1050876, @rjmccall wrote:

In D44985#1050674, @yaxunl wrote:

In D44985#1050670, @rjmccall wrote:

What exactly are you trying to express here? Are you just trying to make these external declarations when compiling for the device because __shared__ variables are actually defined on the host? That should be handled by the frontend by setting up the AST so that these declarations are not definitions.

No. These variables are not like external symbols defined on the host. They behave like global variables in the kernel code but never initialized. Currently no targets are able to initialize them and it is users' responsibility to initialize them explicitly.

Giving them an initial value will cause error in some backends since they cannot handle them, therefore put undef as initializer.

So undef is being used as a special marker to the backends that it's okay not to try to initialize these variables?

I think undef as the initializer tells the llvm passes and backend that this global variable contains undefined value. I am not sure if this is better than without an initializer. I saw code in CodeGenModule::getOrCreateStaticVarDecl

// Local address space cannot have an initializer.
llvm::Constant *Init = nullptr;
if (Ty.getAddressSpace() != LangAS::opencl_local)
  Init = EmitNullConstant(Ty);
else
  Init = llvm::UndefValue::get(LTy);

which means OpenCL static variable in local address space (equivalent to CUDA shared address space) gets an undef initializer.

For CUDA shared variable, in CodeGenFunction::EmitStaticVarDecl, it first goes through call of CodeGenModule::getOrCreateStaticVarDecl and gets a zeroinitializer, then it reaches line 400

// Whatever initializer such variable may have when it gets here is
  // a no-op and should not be emitted.
  bool isCudaSharedVar = getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
                         D.hasAttr<CUDASharedAttr>();
  // If this value has an initializer, emit it.
  if (D.getInit() && !isCudaSharedVar)
    var = AddInitializerToStaticVarDecl(D, var);

Although this disables adding initializer from D, var already has a zeroinitializer from CodeGenModule::getOrCreateStaticVarDecl, therefore its initializer needs to be overwritten by undef.

Probably a better solution would be do it in CodeGenModule::getOrCreateStaticVarDecl, side by side by the OpenCL code.

Yes, I agree, just updating the condition to trigger if either language mode is set is the right fix.

[yaxunl]

Revised by John's comments. Also simplified the test by Artem's comments.

[tra]

lib/CodeGen/CGDecl.cpp

235–240 ↗

(On Diff #140310)

This is too hard to read. Inverting it makes it somewhat easier to understand -- either opencl_local or device-side CUDA shared are undef.

if (Ty.getAddressSpace() == LangAS::opencl_local 
    || (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
        D.hasAttr<CUDASharedAttr>()))
  Init = llvm::UndefValue::get(LTy);
else
  Init = EmitNullConstant(Ty);

[yaxunl]

Revised by Artem's comment. Improve readability of condition.

[tra]

Still LGTM.

[rjmccall]

lib/CodeGen/CGDecl.cpp

235–240 ↗

(On Diff #140310)

I assume getLangOpts().CUDAIsDevice implies getLangOpts().CUDA, so you really only need to check CUDAIsDevice. But it might be faster still to just check for the attribute.

[yaxunl]

Only check attribute.

lib/CodeGen/CGDecl.cpp

235–240 ↗

(On Diff #140310)

Right. I think only attribute needs to be checked.

[rjmccall]

LGTM, thanks.

cfe/trunk/lib/CodeGen/CGDecl.cpp

if (D.hasAttr<AsmLabelAttr>())
     Name = getMangledName(&D);
   else
     Name = getStaticDeclName(*this, D);
 
   llvm::Type *LTy = getTypes().ConvertTypeForMem(Ty);
   LangAS AS = GetGlobalVarAddressSpace(&D);
   unsigned TargetAS = getContext().getTargetAddressSpace(AS);
 
-  // Local address space cannot have an initializer.
+  // OpenCL variables in local address space and CUDA shared
+  // variables cannot have an initializer.
   llvm::Constant *Init = nullptr;
-  if (Ty.getAddressSpace() != LangAS::opencl_local)
-    Init = EmitNullConstant(Ty);
-  else
+  if (Ty.getAddressSpace() == LangAS::opencl_local ||
+      D.hasAttr<CUDASharedAttr>())
     Init = llvm::UndefValue::get(LTy);
+  else
+    Init = EmitNullConstant(Ty);
 
   llvm::GlobalVariable *GV = new llvm::GlobalVariable(
       getModule(), LTy, Ty.isConstant(getContext()), Linkage, Init, Name,
       nullptr, llvm::GlobalVariable::NotThreadLocal, TargetAS);
   GV->setAlignment(getContext().getDeclAlign(&D).getQuantity());
 
   if (supportsCOMDAT() && GV->isWeakForLinker())
     GV->setComdat(TheModule.getOrInsertComdat(GV->getName()));

cfe/trunk/test/CodeGenCUDA/address-spaces.cu

 // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple amdgcn | FileCheck %s
 
 // Verifies Clang emits correct address spaces and addrspacecast instructions
 // for CUDA code.
 
 #include "Inputs/cuda.h"
 
 // CHECK: @i = addrspace(1) externally_initialized global
 __device__ int i;
 
 // 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: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct undef
+// CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float undef
+// CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] undef
+// CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float undef
+// CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float undef
 // 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++;
 
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*)
   k++;
 
   __shared__ int lk;
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
   lk++;
 }
 
 __device__ void func0() {
   __shared__ MyStruct a;
   MyStruct *ap = &a; // composite type
   ap->data1 = 1;
   ap->data2 = 2;
 }
 // CHECK: define void @_Z5func0v()
-// CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %ap
+// CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %{{.*}}
 
 __device__ void callee(float *ap) {
   *ap = 1.0f;
 }
 
 __device__ void func1() {
   __shared__ float a;
   callee(&a); // implicit cast from parameters
 }
 // CHECK: define void @_Z5func1v()
 // CHECK: call void @_Z6calleePf(float* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float*))
 
 __device__ void func2() {
   __shared__ float a[256];
   float *ap = &a[128]; // implicit cast from a decayed array
   *ap = 1.0f;
 }
 // CHECK: define void @_Z5func2v()
-// CHECK: store float* getelementptr inbounds ([256 x float], [256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i32 0, i32 128), float** %ap
+// CHECK: store float* getelementptr inbounds ([256 x float], [256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i{{32|64}} 0, i{{32|64}} 128), float** %{{.*}}
 
 __device__ void func3() {
   __shared__ float a;
   float *ap = reinterpret_cast<float *>(&a); // explicit cast
   *ap = 1.0f;
 }
 // CHECK: define void @_Z5func3v()
-// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %ap
+// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %{{.*}}
 
 __device__ void func4() {
   __shared__ float a;
   float *ap = (float *)&a; // explicit c-style cast
   *ap = 1.0f;
 }
 // CHECK: define void @_Z5func4v()
-// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %ap
+// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %{{.*}}
 
 __shared__ float b;
 
 __device__ float *func5() {
   return &b; // implicit cast from a return value
 }
 // CHECK: define float* @_Z5func5v()
 // CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)

cfe/trunk/test/CodeGenCUDA/device-var-init.cu

 // REQUIRES: nvptx-registered-target
+// REQUIRES: amdgpu-registered-target
 
 // Make sure we don't allow dynamic initialization for device
 // variables, but accept empty constructors allowed by CUDA.
 
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \
-// RUN:     -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck %s
+// RUN:     -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,NVPTX %s
+
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 \
+// RUN:     -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,AMDGCN %s
 
 #ifdef __clang__
 #include "Inputs/cuda.h"
 #endif
 
 // Use the types we share with Sema tests.
 #include "Inputs/cuda-initializers.h"
 

 // Calling empty base class initializer is OK
 __device__ EC_I_EC d_ec_i_ec;
 // CHECK: @d_ec_i_ec = addrspace(1) externally_initialized global %struct.EC_I_EC zeroinitializer,
 __shared__ EC_I_EC s_ec_i_ec;
 // CHECK: @s_ec_i_ec = addrspace(3) global %struct.EC_I_EC undef,
 __constant__ EC_I_EC c_ec_i_ec;
 // CHECK: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer,
 
+// CHECK: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef
+// CHECK: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef
+
 // We should not emit global initializers for device-side variables.
 // CHECK-NOT: @__cxx_global_var_init
 
 // Make sure that initialization restrictions do not apply to local
 // variables.
 __device__ void df() {
+  // NVPTX:  %[[ec:.*]] = alloca %struct.EC
+  // NVPTX:  %[[ed:.*]] = alloca %struct.ED
+  // NVPTX:  %[[ecd:.*]] = alloca %struct.ECD
+  // NVPTX:  %[[etc:.*]] = alloca %struct.ETC
+  // NVPTX:  %[[uc:.*]] = alloca %struct.UC
+  // NVPTX:  %[[ud:.*]] = alloca %struct.UD
+  // NVPTX:  %[[eci:.*]] = alloca %struct.ECI
+  // NVPTX:  %[[nec:.*]] = alloca %struct.NEC
+  // NVPTX:  %[[ned:.*]] = alloca %struct.NED
+  // NVPTX:  %[[ncv:.*]] = alloca %struct.NCV
+  // NVPTX:  %[[vd:.*]] = alloca %struct.VD
+  // NVPTX:  %[[ncf:.*]] = alloca %struct.NCF
+  // NVPTX:  %[[ncfs:.*]] = alloca %struct.NCFS
+  // NVPTX:  %[[utc:.*]] = alloca %struct.UTC
+  // NVPTX:  %[[netc:.*]] = alloca %struct.NETC
+  // NVPTX:  %[[ec_i_ec:.*]] = alloca %struct.EC_I_EC
+  // NVPTX:  %[[ec_i_ec1:.*]] = alloca %struct.EC_I_EC1
+  // NVPTX:  %[[t_v_t:.*]] = alloca %struct.T_V_T
+  // NVPTX:  %[[t_b_nec:.*]] = alloca %struct.T_B_NEC
+  // NVPTX:  %[[t_f_nec:.*]] = alloca %struct.T_F_NEC
+  // NVPTX:  %[[t_fa_nec:.*]] = alloca %struct.T_FA_NEC
+  // NVPTX:  %[[t_b_ned:.*]] = alloca %struct.T_B_NED
+  // NVPTX:  %[[t_f_ned:.*]] = alloca %struct.T_F_NED
+  // NVPTX:  %[[t_fa_ned:.*]] = alloca %struct.T_FA_NED
+  // AMDGCN:  %[[ec:.*]] = addrspacecast %struct.EC addrspace(5)* %ec to %struct.EC*
+  // AMDGCN:  %[[ed:.*]] = addrspacecast %struct.ED addrspace(5)* %ed to %struct.ED*
+  // AMDGCN:  %[[ecd:.*]] = addrspacecast %struct.ECD addrspace(5)* %ecd to %struct.ECD*
+  // AMDGCN:  %[[etc:.*]] = addrspacecast %struct.ETC addrspace(5)* %etc to %struct.ETC*
+  // AMDGCN:  %[[uc:.*]] = addrspacecast %struct.UC addrspace(5)* %uc to %struct.UC*
+  // AMDGCN:  %[[ud:.*]] = addrspacecast %struct.UD addrspace(5)* %ud to %struct.UD*
+  // AMDGCN:  %[[eci:.*]] = addrspacecast %struct.ECI addrspace(5)* %eci to %struct.ECI*
+  // AMDGCN:  %[[nec:.*]] = addrspacecast %struct.NEC addrspace(5)* %nec to %struct.NEC*
+  // AMDGCN:  %[[ned:.*]] = addrspacecast %struct.NED addrspace(5)* %ned to %struct.NED*
+  // AMDGCN:  %[[ncv:.*]] = addrspacecast %struct.NCV addrspace(5)* %ncv to %struct.NCV*
+  // AMDGCN:  %[[vd:.*]] = addrspacecast %struct.VD addrspace(5)* %vd to %struct.VD*
+  // AMDGCN:  %[[ncf:.*]] = addrspacecast %struct.NCF addrspace(5)* %ncf to %struct.NCF*
+  // AMDGCN:  %[[ncfs:.*]] = addrspacecast %struct.NCFS addrspace(5)* %ncfs to %struct.NCFS*
+  // AMDGCN:  %[[utc:.*]] = addrspacecast %struct.UTC addrspace(5)* %utc to %struct.UTC*
+  // AMDGCN:  %[[netc:.*]] = addrspacecast %struct.NETC addrspace(5)* %netc to %struct.NETC*
+  // AMDGCN:  %[[ec_i_ec:.*]] = addrspacecast %struct.EC_I_EC addrspace(5)* %ec_i_ec to %struct.EC_I_EC*
+  // AMDGCN:  %[[ec_i_ec1:.*]] = addrspacecast %struct.EC_I_EC1 addrspace(5)* %ec_i_ec1 to %struct.EC_I_EC1*
+  // AMDGCN:  %[[t_v_t:.*]] = addrspacecast %struct.T_V_T addrspace(5)* %t_v_t to %struct.T_V_T*
+  // AMDGCN:  %[[t_b_nec:.*]] = addrspacecast %struct.T_B_NEC addrspace(5)* %t_b_nec to %struct.T_B_NEC*
+  // AMDGCN:  %[[t_f_nec:.*]] = addrspacecast %struct.T_F_NEC addrspace(5)* %t_f_nec to %struct.T_F_NEC*
+  // AMDGCN:  %[[t_fa_nec:.*]] = addrspacecast %struct.T_FA_NEC addrspace(5)* %t_fa_nec to %struct.T_FA_NEC*
+  // AMDGCN:  %[[t_b_ned:.*]] = addrspacecast %struct.T_B_NED addrspace(5)* %t_b_ned to %struct.T_B_NED*
+  // AMDGCN:  %[[t_f_ned:.*]] = addrspacecast %struct.T_F_NED addrspace(5)* %t_f_ned to %struct.T_F_NED*
+  // AMDGCN:  %[[t_fa_ned:.*]] = addrspacecast %struct.T_FA_NED addrspace(5)* %t_fa_ned to %struct.T_FA_NED*
+
   T t;
   // CHECK-NOT: call
   EC ec;
-  // CHECK:   call void @_ZN2ECC1Ev(%struct.EC* %ec)
+  // CHECK:  call void @_ZN2ECC1Ev(%struct.EC* %[[ec]])
   ED ed;
   // CHECK-NOT: call
   ECD ecd;
-  // CHECK:   call void @_ZN3ECDC1Ev(%struct.ECD* %ecd)
+  // CHECK:  call void @_ZN3ECDC1Ev(%struct.ECD* %[[ecd]])
   ETC etc;
-  // CHECK:   call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
+  // CHECK:  call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %[[etc]])
   UC uc;
   // undefined constructor -- not allowed
-  // CHECK:   call void @_ZN2UCC1Ev(%struct.UC* %uc)
+  // CHECK:  call void @_ZN2UCC1Ev(%struct.UC* %[[uc]])
   UD ud;
   // undefined destructor -- not allowed
   // CHECK-NOT: call
   ECI eci;
   // empty constructor w/ initializer list -- not allowed
-  // CHECK:   call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
+  // CHECK:  call void @_ZN3ECIC1Ev(%struct.ECI* %[[eci]])
   NEC nec;
   // non-empty constructor -- not allowed
-  // CHECK:   call void @_ZN3NECC1Ev(%struct.NEC* %nec)
+  // CHECK:  call void @_ZN3NECC1Ev(%struct.NEC* %[[nec]])
   // non-empty destructor -- not allowed
   NED ned;
   // no-constructor,  virtual method -- not allowed
-  // CHECK:   call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
+  // CHECK:  call void @_ZN3NCVC1Ev(%struct.NCV* %[[ncv]])
   NCV ncv;
   // CHECK-NOT: call
   VD vd;
-  // CHECK:   call void @_ZN2VDC1Ev(%struct.VD* %vd)
+  // CHECK:  call void @_ZN2VDC1Ev(%struct.VD* %[[vd]])
   NCF ncf;
-  // CHECK:   call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
+  // CHECK:   call void @_ZN3NCFC1Ev(%struct.NCF* %[[ncf]])
   NCFS ncfs;
-  // CHECK:   call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs)
+  // CHECK:  call void @_ZN4NCFSC1Ev(%struct.NCFS* %[[ncfs]])
   UTC utc;
-  // CHECK:   call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc)
+  // CHECK:  call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %[[utc]])
   NETC netc;
-  // CHECK:   call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc)
+  // CHECK:  call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %[[netc]])
   T_B_T t_b_t;
   // CHECK-NOT: call
   T_F_T t_f_t;
   // CHECK-NOT: call
   T_FA_T t_fa_t;
   // CHECK-NOT: call
   EC_I_EC ec_i_ec;
-  // CHECK:   call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec)
+  // CHECK:  call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %[[ec_i_ec]])
   EC_I_EC1 ec_i_ec1;
-  // CHECK:   call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1)
+  // CHECK:  call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %[[ec_i_ec1]])
   T_V_T t_v_t;
-  // CHECK:   call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t)
+  // CHECK:  call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %[[t_v_t]])
   T_B_NEC t_b_nec;
-  // CHECK:   call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
+  // CHECK:  call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %[[t_b_nec]])
   T_F_NEC t_f_nec;
-  // CHECK:   call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
+  // CHECK:  call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %[[t_f_nec]])
   T_FA_NEC t_fa_nec;
-  // CHECK:   call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
+  // CHECK:  call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %[[t_fa_nec]])
   T_B_NED t_b_ned;
   // CHECK-NOT: call
   T_F_NED t_f_ned;
   // CHECK-NOT: call
   T_FA_NED t_fa_ned;
   // CHECK-NOT: call
   static __shared__ EC s_ec;
   // CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
   static __shared__ ETC s_etc;
   // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
 
   // anchor point separating constructors and destructors
   df(); // CHECK: call void @_Z2dfv()
 
   // Verify that we only call non-empty destructors
-  // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned)
-  // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned)
-  // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned)
-  // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %vd)
-  // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %ned)
-  // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %ud)
-  // CHECK-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %ecd)
-  // CHECK-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %ed)
+  // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %[[t_fa_ned]])
+  // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %[[t_f_ned]])
+  // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %[[t_b_ned]])
+  // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %[[vd]])
+  // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %[[ned]])
+  // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %[[ud]])
+  // CHECK-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %[[ecd]])
+  // CHECK-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %[[ed]])
 
   // CHECK-NEXT: ret void
 }
 
 // We should not emit global init function.
 // CHECK-NOT: @_GLOBAL__sub_I