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))
kernel(int * out)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