55 lines · plain
1// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -fcuda-is-device -fsyntax-only -verify \2// RUN: -verify-ignore-unexpected=note %s3 4#include "__clang_cuda_builtin_vars.h"5__attribute__((global))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 to50 // take address of a built-in variable in CUDA. Alas there's no way51 // to prevent getting address of a 'const int', so the line52 // currently compiles without errors or warnings.53 const void *wsptr = &warpSize;54}55