34 lines · plain
1// RUN: %clang_cc1 -fsyntax-only -verify %s2// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s3#include "Inputs/cuda.h"4 5struct 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 12void 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 right15// away, without instantiation.16template <typename T>17__global__ void tkernel_nd_const(__grid_constant__ const S arg, T dummy) {}18template <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.22template <typename T>23__global__ void tkernel_const(__grid_constant__ const T arg) {}24 25template <typename T>26__global__ void tkernel(__grid_constant__ T arg) {} // expected-error {{__grid_constant__ is only allowed on const-qualified kernel parameters}}27 28void 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