r349087 - [CUDA] Make all host-side shadows of device-side variables undef.

Artem Belevich via cfe-commits cfe-commits at lists.llvm.org
Thu Dec 13 13:43:05 PST 2018


Author: tra
Date: Thu Dec 13 13:43:04 2018
New Revision: 349087

URL: http://llvm.org/viewvc/llvm-project?rev=349087&view=rev
Log:
[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

Modified:
    cfe/trunk/lib/CodeGen/CodeGenModule.cpp
    cfe/trunk/test/CodeGenCUDA/device-var-init.cu

Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=349087&r1=349086&r2=349087&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Thu Dec 13 13:43:04 2018
@@ -3485,8 +3485,15 @@ void CodeGenModule::EmitGlobalVarDefinit
   // 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

Modified: cfe/trunk/test/CodeGenCUDA/device-var-init.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/device-var-init.cu?rev=349087&r1=349086&r2=349087&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/device-var-init.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/device-var-init.cu Thu Dec 13 13:43:04 2018
@@ -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




More information about the cfe-commits mailing list