| 1 | // RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -fcuda-is-device -fsyntax-only -verify %s |
|---|
| 2 | |
|---|
| 3 | #include "__clang_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@__clang_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@__clang_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@__clang_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@__clang_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@__clang_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 | |
|---|