303 lines · plain
1// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify=dev %s2// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify=host %s3 4// host-no-diagnostics5 6#include "Inputs/cuda.h"7 8int func();9 10struct A {11 int x;12 static int host_var;13};14 15int A::host_var; // dev-note {{host variable declared here}}16 17namespace X {18 int host_var; // dev-note {{host variable declared here}}19}20 21// struct with non-empty ctor.22struct B1 {23 int x;24 B1() { x = 1; }25};26 27// struct with non-empty dtor.28struct B2 {29 int x;30 B2() {}31 ~B2() { x = 0; }32};33 34static int static_host_var; // dev-note {{host variable declared here}}35 36__device__ int global_dev_var;37__constant__ int global_constant_var;38__shared__ int global_shared_var;39 40int global_host_var; // dev-note 8{{host variable declared here}}41const int global_const_var = 1;42constexpr int global_constexpr_var = 1;43 44int global_host_array[2] = {1, 2}; // dev-note {{host variable declared here}}45const int global_const_array[2] = {1, 2};46constexpr int global_constexpr_array[2] = {1, 2};47 48A global_host_struct_var{1}; // dev-note 2{{host variable declared here}}49const A global_const_struct_var{1};50constexpr A global_constexpr_struct_var{1};51 52// Check const host var initialized with non-empty ctor is not allowed in53// device function.54const B1 b1; // dev-note {{const variable cannot be emitted on device side due to dynamic initialization}}55 56// Check const host var having non-empty dtor is not allowed in device function.57const B2 b2; // dev-note {{const variable cannot be emitted on device side due to dynamic initialization}}58 59// Check const host var initialized by non-constant initializer is not allowed60// in device function.61const int b3 = func(); // dev-note {{const variable cannot be emitted on device side due to dynamic initialization}}62 63template<typename F>64__global__ void kernel(F f) { f(); } // dev-note2 {{called by 'kernel<(lambda}}65 66__device__ void dev_fun(int *out) {67 // Check access device variables are allowed.68 int &ref_dev_var = global_dev_var;69 int &ref_constant_var = global_constant_var;70 int &ref_shared_var = global_shared_var;71 *out = ref_dev_var;72 *out = ref_constant_var;73 *out = ref_shared_var;74 *out = global_dev_var;75 *out = global_constant_var;76 *out = global_shared_var;77 78 // Check access of non-const host variables are not allowed.79 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}80 *out = global_const_var;81 *out = global_constexpr_var;82 *out = b1.x; // dev-error {{reference to __host__ variable 'b1' in __device__ function}}83 *out = b2.x; // dev-error {{reference to __host__ variable 'b2' in __device__ function}}84 *out = b3; // dev-error {{reference to __host__ variable 'b3' in __device__ function}}85 global_host_var = 1; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}86 87 // Check reference of non-constexpr host variables are not allowed.88 int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}89 const int &ref_const_var = global_const_var;90 const int &ref_constexpr_var = global_constexpr_var;91 *out = ref_host_var;92 *out = ref_constexpr_var;93 *out = ref_const_var;94 95 // Check access member of non-constexpr struct type host variable is not allowed.96 *out = global_host_struct_var.x; // dev-error {{reference to __host__ variable 'global_host_struct_var' in __device__ function}}97 *out = global_const_struct_var.x;98 *out = global_constexpr_struct_var.x;99 global_host_struct_var.x = 1; // dev-error {{reference to __host__ variable 'global_host_struct_var' in __device__ function}}100 101 // Check address taking of non-constexpr host variables is not allowed.102 int *p = &global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}103 const int *cp = &global_const_var;104 const int *cp2 = &global_constexpr_var;105 106 // Check access elements of non-constexpr host array is not allowed.107 *out = global_host_array[1]; // dev-error {{reference to __host__ variable 'global_host_array' in __device__ function}}108 *out = global_const_array[1];109 *out = global_constexpr_array[1];110 111 // Check ODR-use of host variables in namespace is not allowed.112 *out = X::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}}113 114 // Check ODR-use of static host variables in class or file scope is not115 // allowed.116 *out = A::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}}117 *out = static_host_var; // dev-error {{reference to __host__ variable 'static_host_var' in __device__ function}}118 119 // Check function-scope static variable is allowed.120 static int static_var;121 *out = static_var;122 123 // Check non-ODR use of host varirables are allowed.124 *out = sizeof(global_host_var);125 *out = sizeof(global_host_struct_var.x);126 decltype(global_host_var) var1;127 decltype(global_host_struct_var.x) var2;128}129 130__global__ void global_fun(int *out) {131 int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}}132 int &ref_dev_var = global_dev_var;133 int &ref_constant_var = global_constant_var;134 int &ref_shared_var = global_shared_var;135 const int &ref_constexpr_var = global_constexpr_var;136 const int &ref_const_var = global_const_var;137 138 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}}139 *out = global_dev_var;140 *out = global_constant_var;141 *out = global_shared_var;142 *out = global_constexpr_var;143 *out = global_const_var;144 145 *out = ref_host_var;146 *out = ref_dev_var;147 *out = ref_constant_var;148 *out = ref_shared_var;149 *out = ref_constexpr_var;150 *out = ref_const_var;151}152 153__host__ __device__ void host_dev_fun(int *out) {154 int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}155 int &ref_dev_var = global_dev_var;156 int &ref_constant_var = global_constant_var;157 int &ref_shared_var = global_shared_var;158 const int &ref_constexpr_var = global_constexpr_var;159 const int &ref_const_var = global_const_var;160 161 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}162 *out = global_dev_var;163 *out = global_constant_var;164 *out = global_shared_var;165 *out = global_constexpr_var;166 *out = global_const_var;167 168 *out = ref_host_var;169 *out = ref_dev_var;170 *out = ref_constant_var;171 *out = ref_shared_var;172 *out = ref_constexpr_var;173 *out = ref_const_var;174}175 176inline __host__ __device__ void inline_host_dev_fun(int *out) {177 int &ref_host_var = global_host_var;178 int &ref_dev_var = global_dev_var;179 int &ref_constant_var = global_constant_var;180 int &ref_shared_var = global_shared_var;181 const int &ref_constexpr_var = global_constexpr_var;182 const int &ref_const_var = global_const_var;183 184 *out = global_host_var;185 *out = global_dev_var;186 *out = global_constant_var;187 *out = global_shared_var;188 *out = global_constexpr_var;189 *out = global_const_var;190 191 *out = ref_host_var;192 *out = ref_dev_var;193 *out = ref_constant_var;194 *out = ref_shared_var;195 *out = ref_constexpr_var;196 *out = ref_const_var;197}198 199void dev_lambda_capture_by_ref(int *out) {200 int &ref_host_var = global_host_var;201 kernel<<<1,1>>>([&]() {202 int &ref_dev_var = global_dev_var;203 int &ref_constant_var = global_constant_var;204 int &ref_shared_var = global_shared_var;205 const int &ref_constexpr_var = global_constexpr_var;206 const int &ref_const_var = global_const_var;207 208 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}209 // dev-error@-1 {{capture host variable 'out' by reference in device or host device lambda function}}210 *out = global_dev_var;211 *out = global_constant_var;212 *out = global_shared_var;213 *out = global_constexpr_var;214 *out = global_const_var;215 216 *out = ref_host_var; // dev-error {{capture host variable 'ref_host_var' by reference in device or host device lambda function}}217 *out = ref_dev_var;218 *out = ref_constant_var;219 *out = ref_shared_var;220 *out = ref_constexpr_var;221 *out = ref_const_var;222 });223}224 225void dev_lambda_capture_by_copy(int *out) {226 int &ref_host_var = global_host_var;227 kernel<<<1,1>>>([=]() {228 int &ref_dev_var = global_dev_var;229 int &ref_constant_var = global_constant_var;230 int &ref_shared_var = global_shared_var;231 const int &ref_constexpr_var = global_constexpr_var;232 const int &ref_const_var = global_const_var;233 234 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}235 *out = global_dev_var;236 *out = global_constant_var;237 *out = global_shared_var;238 *out = global_constexpr_var;239 *out = global_const_var;240 241 *out = ref_host_var;242 *out = ref_dev_var;243 *out = ref_constant_var;244 *out = ref_shared_var;245 *out = ref_constexpr_var;246 *out = ref_const_var;247 });248}249 250// Texture references are special. As far as C++ is concerned they are host251// variables that are referenced from device code. However, they are handled252// very differently by the compiler under the hood and such references are253// allowed. Compiler should produce no warning here, but it should diagnose the254// same case without the device_builtin_texture_type attribute.255template <class, int = 1, int = 1>256struct __attribute__((device_builtin_texture_type)) texture {257 static texture<int> ref;258 __device__ void c() {259 auto &x = ref;260 }261};262 263template <class, int = 1, int = 1>264struct not_a_texture {265 static not_a_texture<int> ref;266 __device__ void c() {267 auto &x = ref; // dev-error {{reference to __host__ variable 'ref' in __device__ function}}268 }269};270 271template<>272not_a_texture<int> not_a_texture<int>::ref; // dev-note {{host variable declared here}}273 274__device__ void test_not_a_texture() {275 not_a_texture<int> inst;276 inst.c(); // dev-note {{in instantiation of member function 'not_a_texture<int>::c' requested here}}277}278 279// Test static variable in host function used by device function.280void test_static_var_host() {281 for (int i = 0; i < 10; i++) {282 static int x; // dev-note {{host variable declared here}}283 struct A {284 __device__ int f() {285 return x; // dev-error{{reference to __host__ variable 'x' in __device__ function}}286 }287 };288 }289}290 291// Test static variable in device function used by device function.292__device__ void test_static_var_device() {293 for (int i = 0; i < 10; i++) {294 static int x;295 int y = x;296 struct A {297 __device__ int f() {298 return x;299 }300 };301 }302}303