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