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 } 34