1 // RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -fcuda-is-device -fsyntax-only -verify \ 2 // RUN: -verify-ignore-unexpected=note %s 3 4 #include "__clang_cuda_builtin_vars.h" 5 __attribute__((global)) kernel(int * out)6void kernel(int *out) { 7 int i = 0; 8 out[i++] = threadIdx.x; 9 threadIdx.x = 0; // expected-error {{no setter defined for property 'x'}} 10 out[i++] = threadIdx.y; 11 threadIdx.y = 0; // expected-error {{no setter defined for property 'y'}} 12 out[i++] = threadIdx.z; 13 threadIdx.z = 0; // expected-error {{no setter defined for property 'z'}} 14 15 out[i++] = blockIdx.x; 16 blockIdx.x = 0; // expected-error {{no setter defined for property 'x'}} 17 out[i++] = blockIdx.y; 18 blockIdx.y = 0; // expected-error {{no setter defined for property 'y'}} 19 out[i++] = blockIdx.z; 20 blockIdx.z = 0; // expected-error {{no setter defined for property 'z'}} 21 22 out[i++] = blockDim.x; 23 blockDim.x = 0; // expected-error {{no setter defined for property 'x'}} 24 out[i++] = blockDim.y; 25 blockDim.y = 0; // expected-error {{no setter defined for property 'y'}} 26 out[i++] = blockDim.z; 27 blockDim.z = 0; // expected-error {{no setter defined for property 'z'}} 28 29 out[i++] = gridDim.x; 30 gridDim.x = 0; // expected-error {{no setter defined for property 'x'}} 31 out[i++] = gridDim.y; 32 gridDim.y = 0; // expected-error {{no setter defined for property 'y'}} 33 out[i++] = gridDim.z; 34 gridDim.z = 0; // expected-error {{no setter defined for property 'z'}} 35 36 out[i++] = warpSize; 37 warpSize = 0; // expected-error {{cannot assign to variable 'warpSize' with const-qualified type 'const int'}} 38 // expected-note@__clang_cuda_builtin_vars.h:* {{variable 'warpSize' declared const here}} 39 40 // Make sure we can't construct or assign to the special variables. 41 __cuda_builtin_threadIdx_t x; // expected-error {{call to deleted constructor of '__cuda_builtin_threadIdx_t'}} 42 43 __cuda_builtin_threadIdx_t y = threadIdx; // expected-error {{call to deleted constructor of '__cuda_builtin_threadIdx_t'}} 44 45 threadIdx = threadIdx; // expected-error {{overload resolution selected deleted operator '='}} 46 47 void *ptr = &threadIdx; // expected-error {{overload resolution selected deleted operator '&'}} 48 49 // Following line should've caused an error as one is not allowed to 50 // take address of a built-in variable in CUDA. Alas there's no way 51 // to prevent getting address of a 'const int', so the line 52 // currently compiles without errors or warnings. 53 const void *wsptr = &warpSize; 54 } 55