reference, declaration → definition
definition → references, declarations, derived classes, virtual overrides
reference to multiple definitions → definitions
unreferenced

    1
    2
    3
    4
    5
    6
    7
    8
    9
   10
   11
   12
   13
   14
   15
   16
   17
   18
   19
   20
   21
   22
   23
   24
   25
   26
   27
   28
   29
   30
   31
   32
   33
   34
   35
   36
   37
   38
   39
   40
   41
   42
   43
   44
   45
   46
   47
   48
   49
   50
   51
   52
   53
   54
   55
   56
   57

// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -fcuda-is-device -fsyntax-only -verify %s

#include "__clang_cuda_builtin_vars.h"
__attribute__((global))
void kernel(int *out) {
  int i = 0;
  out[i++] = threadIdx.x;
  threadIdx.x = 0; // expected-error {{no setter defined for property 'x'}}
  out[i++] = threadIdx.y;
  threadIdx.y = 0; // expected-error {{no setter defined for property 'y'}}
  out[i++] = threadIdx.z;
  threadIdx.z = 0; // expected-error {{no setter defined for property 'z'}}

  out[i++] = blockIdx.x;
  blockIdx.x = 0; // expected-error {{no setter defined for property 'x'}}
  out[i++] = blockIdx.y;
  blockIdx.y = 0; // expected-error {{no setter defined for property 'y'}}
  out[i++] = blockIdx.z;
  blockIdx.z = 0; // expected-error {{no setter defined for property 'z'}}

  out[i++] = blockDim.x;
  blockDim.x = 0; // expected-error {{no setter defined for property 'x'}}
  out[i++] = blockDim.y;
  blockDim.y = 0; // expected-error {{no setter defined for property 'y'}}
  out[i++] = blockDim.z;
  blockDim.z = 0; // expected-error {{no setter defined for property 'z'}}

  out[i++] = gridDim.x;
  gridDim.x = 0; // expected-error {{no setter defined for property 'x'}}
  out[i++] = gridDim.y;
  gridDim.y = 0; // expected-error {{no setter defined for property 'y'}}
  out[i++] = gridDim.z;
  gridDim.z = 0; // expected-error {{no setter defined for property 'z'}}

  out[i++] = warpSize;
  warpSize = 0; // expected-error {{cannot assign to variable 'warpSize' with const-qualified type 'const int'}}
  // expected-note@__clang_cuda_builtin_vars.h:* {{variable 'warpSize' declared const here}}

  // Make sure we can't construct or assign to the special variables.
  __cuda_builtin_threadIdx_t x; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}}
  // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}

  __cuda_builtin_threadIdx_t y = threadIdx; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}}
  // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}

  threadIdx = threadIdx; // expected-error {{'operator=' is a private member of '__cuda_builtin_threadIdx_t'}}
  // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}

  void *ptr = &threadIdx; // expected-error {{'operator&' is a private member of '__cuda_builtin_threadIdx_t'}}
  // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}

  // Following line should've caused an error as one is not allowed to
  // take address of a built-in variable in CUDA. Alas there's no way
  // to prevent getting address of a 'const int', so the line
  // currently compiles without errors or warnings.
  const void *wsptr = &warpSize;
}