clang/test/SemaCUDA/cuda-builtin-vars.cu

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