tra created this revision.
tra added reviewers: jlebar, yaxunl.
Herald added subscribers: bixia, sanjoy.
The host-side code can't (and should not) access the values that may
only exist on the device side. E.g. address of a __device__ function
does not exist on the host side as we don't generate the code for it there
This is similar to what nvcc does and allows us to compile code which has
device-side variables with initializers that only exist on device side:
__device__ void func() {}
__device__ funcptr_t Funcs[] = {func};
https://reviews.llvm.org/D55663
Files:
clang/lib/CodeGen/CodeGenModule.cpp
clang/test/CodeGenCUDA/device-var-init.cu
Index: clang/test/CodeGenCUDA/device-var-init.cu
===================================================================
--- clang/test/CodeGenCUDA/device-var-init.cu
+++ clang/test/CodeGenCUDA/device-var-init.cu
@@ -5,10 +5,12 @@
// 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 -check-prefixes=CHECK,NVPTX %s
+// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,NVPTX %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -std=c++11 \
+// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=HOST %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
+// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,AMDGCN %s
#ifdef __clang__
#include "Inputs/cuda.h"
@@ -18,105 +20,140 @@
#include "Inputs/cuda-initializers.h"
__device__ int d_v;
-// CHECK: @d_v = addrspace(1) externally_initialized global i32 0,
+// DEVICE: @d_v = addrspace(1) externally_initialized global i32 0,
+// HOST: @d_v = internal global i32 undef,
__shared__ int s_v;
-// CHECK: @s_v = addrspace(3) global i32 undef,
+// DEVICE: @s_v = addrspace(3) global i32 undef,
+// HOST: @s_v = internal global i32 undef,
__constant__ int c_v;
-// CHECK: addrspace(4) externally_initialized global i32 0,
+// DEVICE: addrspace(4) externally_initialized global i32 0,
+// HOST: @c_v = internal global i32 undef,
__device__ int d_v_i = 1;
-// CHECK: @d_v_i = addrspace(1) externally_initialized global i32 1,
+// DEVICE: @d_v_i = addrspace(1) externally_initialized global i32 1,
+// HOST: @d_v_i = internal global i32 undef,
// trivial constructor -- allowed
__device__ T d_t;
-// CHECK: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer
+// DEVICE: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer
+// HOST: @d_t = internal global %struct.T undef,
__shared__ T s_t;
-// CHECK: @s_t = addrspace(3) global %struct.T undef,
+// DEVICE: @s_t = addrspace(3) global %struct.T undef,
+// HOST: @s_t = internal global %struct.T undef,
__constant__ T c_t;
-// CHECK: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer,
+// DEVICE: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer,
+// HOST: @c_t = internal global %struct.T undef,
__device__ T d_t_i = {2};
-// CHECK: @d_t_i = addrspace(1) externally_initialized global %struct.T { i32 2 },
+// DEVICE: @d_t_i = addrspace(1) externally_initialized global %struct.T { i32 2 },
+// HOST: @d_t_i = internal global %struct.T undef,
__constant__ T c_t_i = {2};
-// CHECK: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 },
+// DEVICE: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 },
+// HOST: @c_t_i = internal global %struct.T undef,
// empty constructor
__device__ EC d_ec;
-// CHECK: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer,
+// DEVICE: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer,
+// HOST: @d_ec = internal global %struct.EC undef,
__shared__ EC s_ec;
-// CHECK: @s_ec = addrspace(3) global %struct.EC undef,
+// DEVICE: @s_ec = addrspace(3) global %struct.EC undef,
+// HOST: @s_ec = internal global %struct.EC undef,
__constant__ EC c_ec;
-// CHECK: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer,
+// DEVICE: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer,
+// HOST: @c_ec = internal global %struct.EC undef
// empty destructor
__device__ ED d_ed;
-// CHECK: @d_ed = addrspace(1) externally_initialized global %struct.ED zeroinitializer,
+// DEVICE: @d_ed = addrspace(1) externally_initialized global %struct.ED zeroinitializer,
+// HOST: @d_ed = internal global %struct.ED undef,
__shared__ ED s_ed;
-// CHECK: @s_ed = addrspace(3) global %struct.ED undef,
+// DEVICE: @s_ed = addrspace(3) global %struct.ED undef,
+// HOST: @s_ed = internal global %struct.ED undef,
__constant__ ED c_ed;
-// CHECK: @c_ed = addrspace(4) externally_initialized global %struct.ED zeroinitializer,
+// DEVICE: @c_ed = addrspace(4) externally_initialized global %struct.ED zeroinitializer,
+// HOST: @c_ed = internal global %struct.ED undef,
__device__ ECD d_ecd;
-// CHECK: @d_ecd = addrspace(1) externally_initialized global %struct.ECD zeroinitializer,
+// DEVICE: @d_ecd = addrspace(1) externally_initialized global %struct.ECD zeroinitializer,
+// HOST: @d_ecd = internal global %struct.ECD undef,
__shared__ ECD s_ecd;
-// CHECK: @s_ecd = addrspace(3) global %struct.ECD undef,
+// DEVICE: @s_ecd = addrspace(3) global %struct.ECD undef,
+// HOST: @s_ecd = internal global %struct.ECD undef,
__constant__ ECD c_ecd;
-// CHECK: @c_ecd = addrspace(4) externally_initialized global %struct.ECD zeroinitializer,
+// DEVICE: @c_ecd = addrspace(4) externally_initialized global %struct.ECD zeroinitializer,
+// HOST: @c_ecd = internal global %struct.ECD undef,
// empty templated constructor -- allowed with no arguments
__device__ ETC d_etc;
-// CHECK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer,
+// DEVICE: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer,
+// HOST: @d_etc = internal global %struct.ETC undef,
__shared__ ETC s_etc;
-// CHECK: @s_etc = addrspace(3) global %struct.ETC undef,
+// DEVICE: @s_etc = addrspace(3) global %struct.ETC undef,
+// HOST: @s_etc = internal global %struct.ETC undef,
__constant__ ETC c_etc;
-// CHECK: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer,
+// DEVICE: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer,
+// HOST: @c_etc = internal global %struct.ETC undef,
__device__ NCFS d_ncfs;
-// CHECK: @d_ncfs = addrspace(1) externally_initialized global %struct.NCFS { i32 3 }
+// DEVICE: @d_ncfs = addrspace(1) externally_initialized global %struct.NCFS { i32 3 }
+// HOST: @d_ncfs = internal global %struct.NCFS undef,
__constant__ NCFS c_ncfs;
-// CHECK: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 }
+// DEVICE: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 }
+// HOST: @c_ncfs = internal global %struct.NCFS undef,
// Regular base class -- allowed
__device__ T_B_T d_t_b_t;
-// CHECK: @d_t_b_t = addrspace(1) externally_initialized global %struct.T_B_T zeroinitializer,
+// DEVICE: @d_t_b_t = addrspace(1) externally_initialized global %struct.T_B_T zeroinitializer,
+// HOST: @d_t_b_t = internal global %struct.T_B_T undef,
__shared__ T_B_T s_t_b_t;
-// CHECK: @s_t_b_t = addrspace(3) global %struct.T_B_T undef,
+// DEVICE: @s_t_b_t = addrspace(3) global %struct.T_B_T undef,
+// HOST: @s_t_b_t = internal global %struct.T_B_T undef,
__constant__ T_B_T c_t_b_t;
-// CHECK: @c_t_b_t = addrspace(4) externally_initialized global %struct.T_B_T zeroinitializer,
+// DEVICE: @c_t_b_t = addrspace(4) externally_initialized global %struct.T_B_T zeroinitializer,
+// HOST: @c_t_b_t = internal global %struct.T_B_T undef,
// Incapsulated object of allowed class -- allowed
__device__ T_F_T d_t_f_t;
-// CHECK: @d_t_f_t = addrspace(1) externally_initialized global %struct.T_F_T zeroinitializer,
+// DEVICE: @d_t_f_t = addrspace(1) externally_initialized global %struct.T_F_T zeroinitializer,
+// HOST: @d_t_f_t = internal global %struct.T_F_T undef,
__shared__ T_F_T s_t_f_t;
-// CHECK: @s_t_f_t = addrspace(3) global %struct.T_F_T undef,
+// DEVICE: @s_t_f_t = addrspace(3) global %struct.T_F_T undef,
+// HOST: @s_t_f_t = internal global %struct.T_F_T undef,
__constant__ T_F_T c_t_f_t;
-// CHECK: @c_t_f_t = addrspace(4) externally_initialized global %struct.T_F_T zeroinitializer,
+// DEVICE: @c_t_f_t = addrspace(4) externally_initialized global %struct.T_F_T zeroinitializer,
+// HOST: @c_t_f_t = internal global %struct.T_F_T undef,
// array of allowed objects -- allowed
__device__ T_FA_T d_t_fa_t;
-// CHECK: @d_t_fa_t = addrspace(1) externally_initialized global %struct.T_FA_T zeroinitializer,
+// DEVICE: @d_t_fa_t = addrspace(1) externally_initialized global %struct.T_FA_T zeroinitializer,
+// HOST: @d_t_fa_t = internal global %struct.T_FA_T undef,
__shared__ T_FA_T s_t_fa_t;
-// CHECK: @s_t_fa_t = addrspace(3) global %struct.T_FA_T undef,
+// DEVICE: @s_t_fa_t = addrspace(3) global %struct.T_FA_T undef,
+// HOST: @s_t_fa_t = internal global %struct.T_FA_T undef,
__constant__ T_FA_T c_t_fa_t;
-// CHECK: @c_t_fa_t = addrspace(4) externally_initialized global %struct.T_FA_T zeroinitializer,
+// DEVICE: @c_t_fa_t = addrspace(4) externally_initialized global %struct.T_FA_T zeroinitializer,
+// HOST: @c_t_fa_t = internal global %struct.T_FA_T undef,
// 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,
+// DEVICE: @d_ec_i_ec = addrspace(1) externally_initialized global %struct.EC_I_EC zeroinitializer,
+// HOST: @d_ec_i_ec = internal global %struct.EC_I_EC undef,
__shared__ EC_I_EC s_ec_i_ec;
-// CHECK: @s_ec_i_ec = addrspace(3) global %struct.EC_I_EC undef,
+// DEVICE: @s_ec_i_ec = addrspace(3) global %struct.EC_I_EC undef,
+// HOST: @s_ec_i_ec = internal 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,
+// DEVICE: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer,
+// HOST: @c_ec_i_ec = internal global %struct.EC_I_EC undef,
-// CHECK: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef
-// CHECK: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef
+// DEVICE: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef
+// DEVICE: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef
-// CHECK: @_ZZ2dfvE11const_array = internal addrspace(4) constant [5 x i32] [i32 1, i32 2, i32 3, i32 4, i32 5]
-// CHECK: @_ZZ2dfvE9const_int = internal addrspace(4) constant i32 123
+// DEVICE: @_ZZ2dfvE11const_array = internal addrspace(4) constant [5 x i32] [i32 1, i32 2, i32 3, i32 4, i32 5]
+// DEVICE: @_ZZ2dfvE9const_int = internal addrspace(4) constant i32 123
// We should not emit global initializers for device-side variables.
-// CHECK-NOT: @__cxx_global_var_init
+// DEVICE-NOT: @__cxx_global_var_init
// Make sure that initialization restrictions do not apply to local
// variables.
@@ -171,90 +208,90 @@
// AMDGCN: %[[t_fa_ned:.*]] = addrspacecast %struct.T_FA_NED addrspace(5)* %t_fa_ned to %struct.T_FA_NED*
T t;
- // CHECK-NOT: call
+ // DEVICE-NOT: call
EC ec;
- // CHECK: call void @_ZN2ECC1Ev(%struct.EC* %[[ec]])
+ // DEVICE: call void @_ZN2ECC1Ev(%struct.EC* %[[ec]])
ED ed;
- // CHECK-NOT: call
+ // DEVICE-NOT: call
ECD ecd;
- // CHECK: call void @_ZN3ECDC1Ev(%struct.ECD* %[[ecd]])
+ // DEVICE: call void @_ZN3ECDC1Ev(%struct.ECD* %[[ecd]])
ETC etc;
- // CHECK: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %[[etc]])
+ // DEVICE: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %[[etc]])
UC uc;
// undefined constructor -- not allowed
- // CHECK: call void @_ZN2UCC1Ev(%struct.UC* %[[uc]])
+ // DEVICE: call void @_ZN2UCC1Ev(%struct.UC* %[[uc]])
UD ud;
// undefined destructor -- not allowed
- // CHECK-NOT: call
+ // DEVICE-NOT: call
ECI eci;
// empty constructor w/ initializer list -- not allowed
- // CHECK: call void @_ZN3ECIC1Ev(%struct.ECI* %[[eci]])
+ // DEVICE: call void @_ZN3ECIC1Ev(%struct.ECI* %[[eci]])
NEC nec;
// non-empty constructor -- not allowed
- // CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %[[nec]])
+ // DEVICE: 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]])
+ // DEVICE: call void @_ZN3NCVC1Ev(%struct.NCV* %[[ncv]])
NCV ncv;
- // CHECK-NOT: call
+ // DEVICE-NOT: call
VD vd;
- // CHECK: call void @_ZN2VDC1Ev(%struct.VD* %[[vd]])
+ // DEVICE: call void @_ZN2VDC1Ev(%struct.VD* %[[vd]])
NCF ncf;
- // CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %[[ncf]])
+ // DEVICE: call void @_ZN3NCFC1Ev(%struct.NCF* %[[ncf]])
NCFS ncfs;
- // CHECK: call void @_ZN4NCFSC1Ev(%struct.NCFS* %[[ncfs]])
+ // DEVICE: call void @_ZN4NCFSC1Ev(%struct.NCFS* %[[ncfs]])
UTC utc;
- // CHECK: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %[[utc]])
+ // DEVICE: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %[[utc]])
NETC netc;
- // CHECK: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %[[netc]])
+ // DEVICE: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %[[netc]])
T_B_T t_b_t;
- // CHECK-NOT: call
+ // DEVICE-NOT: call
T_F_T t_f_t;
- // CHECK-NOT: call
+ // DEVICE-NOT: call
T_FA_T t_fa_t;
- // CHECK-NOT: call
+ // DEVICE-NOT: call
EC_I_EC ec_i_ec;
- // CHECK: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %[[ec_i_ec]])
+ // DEVICE: 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]])
+ // DEVICE: 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]])
+ // DEVICE: 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]])
+ // DEVICE: 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]])
+ // DEVICE: 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]])
+ // DEVICE: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %[[t_fa_nec]])
T_B_NED t_b_ned;
- // CHECK-NOT: call
+ // DEVICE-NOT: call
T_F_NED t_f_ned;
- // CHECK-NOT: call
+ // DEVICE-NOT: call
T_FA_NED t_fa_ned;
- // CHECK-NOT: call
+ // DEVICE-NOT: call
static __shared__ EC s_ec;
- // CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
+ // DEVICE-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*))
+ // DEVICE-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
static const int const_array[] = {1, 2, 3, 4, 5};
static const int const_int = 123;
// anchor point separating constructors and destructors
- df(); // CHECK: call void @_Z2dfv()
+ df(); // DEVICE: 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]])
+ // DEVICE-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %[[t_fa_ned]])
+ // DEVICE-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %[[t_f_ned]])
+ // DEVICE-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %[[t_b_ned]])
+ // DEVICE-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %[[vd]])
+ // DEVICE-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %[[ned]])
+ // DEVICE-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %[[ud]])
+ // DEVICE-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %[[ecd]])
+ // DEVICE-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %[[ed]])
- // CHECK-NEXT: ret void
+ // DEVICE-NEXT: ret void
}
// We should not emit global init function.
-// CHECK-NOT: @_GLOBAL__sub_I
+// DEVICE-NOT: @_GLOBAL__sub_I
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -3485,8 +3485,15 @@
// CUDA E.2.4.1 "__shared__ variables cannot have an initialization
// as part of their declaration." Sema has already checked for
// error cases, so we just need to set Init to UndefValue.
- if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
- D->hasAttr<CUDASharedAttr>())
+ bool IsCUDASharedVar =
+ getLangOpts().CUDAIsDevice && D->hasAttr<CUDASharedAttr>();
+ // Shadows of initialized device-side global variables are also left
+ // undefined.
+ bool IsCUDAShadowVar =
+ !getLangOpts().CUDAIsDevice &&
+ (D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() ||
+ D->hasAttr<CUDASharedAttr>());
+ if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar))
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
else if (!InitExpr) {
// This is a tentative definition; tentative definitions are
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
