1 // RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify=dev %s 2 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify=host %s 3 4 // host-no-diagnostics 5 6 #include "Inputs/cuda.h" 7 8 int global_host_var; 9 __device__ int global_dev_var; 10 __constant__ int global_constant_var; 11 __shared__ int global_shared_var; 12 constexpr int global_constexpr_var = 1; 13 const int global_const_var = 1; 14 15 template<typename F> kernel(F f)16__global__ void kernel(F f) { f(); } // dev-note2 {{called by 'kernel<(lambda}} 17 dev_fun(int * out)18__device__ void dev_fun(int *out) { 19 int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}} 20 int &ref_dev_var = global_dev_var; 21 int &ref_constant_var = global_constant_var; 22 int &ref_shared_var = global_shared_var; 23 const int &ref_constexpr_var = global_constexpr_var; 24 const int &ref_const_var = global_const_var; 25 26 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}} 27 *out = global_dev_var; 28 *out = global_constant_var; 29 *out = global_shared_var; 30 *out = global_constexpr_var; 31 *out = global_const_var; 32 33 *out = ref_host_var; 34 *out = ref_dev_var; 35 *out = ref_constant_var; 36 *out = ref_shared_var; 37 *out = ref_constexpr_var; 38 *out = ref_const_var; 39 } 40 global_fun(int * out)41__global__ void global_fun(int *out) { 42 int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}} 43 int &ref_dev_var = global_dev_var; 44 int &ref_constant_var = global_constant_var; 45 int &ref_shared_var = global_shared_var; 46 const int &ref_constexpr_var = global_constexpr_var; 47 const int &ref_const_var = global_const_var; 48 49 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}} 50 *out = global_dev_var; 51 *out = global_constant_var; 52 *out = global_shared_var; 53 *out = global_constexpr_var; 54 *out = global_const_var; 55 56 *out = ref_host_var; 57 *out = ref_dev_var; 58 *out = ref_constant_var; 59 *out = ref_shared_var; 60 *out = ref_constexpr_var; 61 *out = ref_const_var; 62 } 63 host_dev_fun(int * out)64__host__ __device__ void host_dev_fun(int *out) { 65 int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}} 66 int &ref_dev_var = global_dev_var; 67 int &ref_constant_var = global_constant_var; 68 int &ref_shared_var = global_shared_var; 69 const int &ref_constexpr_var = global_constexpr_var; 70 const int &ref_const_var = global_const_var; 71 72 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}} 73 *out = global_dev_var; 74 *out = global_constant_var; 75 *out = global_shared_var; 76 *out = global_constexpr_var; 77 *out = global_const_var; 78 79 *out = ref_host_var; 80 *out = ref_dev_var; 81 *out = ref_constant_var; 82 *out = ref_shared_var; 83 *out = ref_constexpr_var; 84 *out = ref_const_var; 85 } 86 inline_host_dev_fun(int * out)87inline __host__ __device__ void inline_host_dev_fun(int *out) { 88 int &ref_host_var = global_host_var; 89 int &ref_dev_var = global_dev_var; 90 int &ref_constant_var = global_constant_var; 91 int &ref_shared_var = global_shared_var; 92 const int &ref_constexpr_var = global_constexpr_var; 93 const int &ref_const_var = global_const_var; 94 95 *out = global_host_var; 96 *out = global_dev_var; 97 *out = global_constant_var; 98 *out = global_shared_var; 99 *out = global_constexpr_var; 100 *out = global_const_var; 101 102 *out = ref_host_var; 103 *out = ref_dev_var; 104 *out = ref_constant_var; 105 *out = ref_shared_var; 106 *out = ref_constexpr_var; 107 *out = ref_const_var; 108 } 109 dev_lambda_capture_by_ref(int * out)110void dev_lambda_capture_by_ref(int *out) { 111 int &ref_host_var = global_host_var; 112 kernel<<<1,1>>>([&]() { 113 int &ref_dev_var = global_dev_var; 114 int &ref_constant_var = global_constant_var; 115 int &ref_shared_var = global_shared_var; 116 const int &ref_constexpr_var = global_constexpr_var; 117 const int &ref_const_var = global_const_var; 118 119 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}} 120 // dev-error@-1 {{capture host variable 'out' by reference in device or host device lambda function}} 121 *out = global_dev_var; 122 *out = global_constant_var; 123 *out = global_shared_var; 124 *out = global_constexpr_var; 125 *out = global_const_var; 126 127 *out = ref_host_var; // dev-error {{capture host variable 'ref_host_var' by reference in device or host device lambda function}} 128 *out = ref_dev_var; 129 *out = ref_constant_var; 130 *out = ref_shared_var; 131 *out = ref_constexpr_var; 132 *out = ref_const_var; 133 }); 134 } 135 dev_lambda_capture_by_copy(int * out)136void dev_lambda_capture_by_copy(int *out) { 137 int &ref_host_var = global_host_var; 138 kernel<<<1,1>>>([=]() { 139 int &ref_dev_var = global_dev_var; 140 int &ref_constant_var = global_constant_var; 141 int &ref_shared_var = global_shared_var; 142 const int &ref_constexpr_var = global_constexpr_var; 143 const int &ref_const_var = global_const_var; 144 145 *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}} 146 *out = global_dev_var; 147 *out = global_constant_var; 148 *out = global_shared_var; 149 *out = global_constexpr_var; 150 *out = global_const_var; 151 152 *out = ref_host_var; 153 *out = ref_dev_var; 154 *out = ref_constant_var; 155 *out = ref_shared_var; 156 *out = ref_constexpr_var; 157 *out = ref_const_var; 158 }); 159 } 160 161