clang/test/SemaCUDA/grid-constant.cu

   1 // RUN: %clang_cc1 -fsyntax-only -verify %s
   2 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
   3 #include "Inputs/cuda.h"
   4
   5 struct S {};
   6
   7 __global__ void kernel_struct(__grid_constant__ const S arg) {}
   8 __global__ void kernel_scalar(__grid_constant__ const int arg) {}
   9
  10 __global__ void gc_kernel_non_const(__grid_constant__ S arg) {} // expected-error {{__grid_constant__ is only allowed on const-qualified kernel parameters}}
  11
  12 void non_kernel(__grid_constant__ S arg) {} // expected-error {{__grid_constant__ is only allowed on const-qualified kernel parameters}}
  13
  14 // templates w/ non-dependent argument types get diagnosed right
  15 // away, without instantiation.
  16 template <typename T>
  17 __global__ void tkernel_nd_const(__grid_constant__ const S arg, T dummy) {}
  18 template <typename T>
  19 __global__ void tkernel_nd_non_const(__grid_constant__ S arg, T dummy) {} // expected-error {{__grid_constant__ is only allowed on const-qualified kernel parameters}}
  20
  21 // dependent arguments get diagnosed after instantiation.
  22 template <typename T>
  23 __global__ void tkernel_const(__grid_constant__ const T arg) {}
  24
  25 template <typename T>
  26 __global__ void tkernel(__grid_constant__ T arg) {} // expected-error {{__grid_constant__ is only allowed on const-qualified kernel parameters}}
  27
  28 void foo() {
  29   tkernel_const<const S><<<1,1>>>({});
  30   tkernel_const<S><<<1,1>>>({});
  31   tkernel<const S><<<1,1>>>({});
  32   tkernel<S><<<1,1>>>({}); // expected-note {{in instantiation of function template specialization 'tkernel<S>' requested here}}
  33 }