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