From 7b05666a19c1310d2d3b7e05d08313695b4eaa92 Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Thu, 13 Dec 2018 21:43:04 +0000 Subject: [PATCH] [CUDA] Make all host-side shadows of device-side variables undef. 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. Differential Revision: https://reviews.llvm.org/D55663 llvm-svn: 349087 --- clang/lib/CodeGen/CodeGenModule.cpp | 11 +- clang/test/CodeGenCUDA/device-var-init.cu | 203 ++++++++++++++++++------------ 2 files changed, 129 insertions(+), 85 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index ce4558f..1b7bb2c 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3485,8 +3485,15 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, // 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()) + bool IsCUDASharedVar = + getLangOpts().CUDAIsDevice && D->hasAttr(); + // Shadows of initialized device-side global variables are also left + // undefined. + bool IsCUDAShadowVar = + !getLangOpts().CUDAIsDevice && + (D->hasAttr() || D->hasAttr() || + D->hasAttr()); + if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar)) Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy)); else if (!InitExpr) { // This is a tentative definition; tentative definitions are diff --git a/clang/test/CodeGenCUDA/device-var-init.cu b/clang/test/CodeGenCUDA/device-var-init.cu index f96e42d..af42e69 100644 --- a/clang/test/CodeGenCUDA/device-var-init.cu +++ b/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 @@ __device__ void df() { // 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 -- 2.7.4