105 lines · plain
1// Tests handling of CUDA attributes that are bad either because they're2// applied to the wrong sort of thing, or because they're given in illegal3// combinations.4//5// You should be able to run this file through nvcc for compatibility testing.6//7// RUN: %clang_cc1 -fsyntax-only -Wcuda-compat -verify -DEXPECT_INLINE_WARNING %s8// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -Wcuda-compat -verify %s9 10#include "Inputs/cuda.h"11 12// Try applying attributes to functions and variables. Some should generate13// warnings; others not.14__device__ int a1;15__device__ void a2();16__host__ int b1; // expected-warning {{attribute only applies to functions}}17__host__ void b2();18__constant__ int c1;19__constant__ void c2(); // expected-warning {{attribute only applies to variables}}20__shared__ int d1;21__shared__ void d2(); // expected-warning {{attribute only applies to variables}}22__global__ int e1; // expected-warning {{attribute only applies to functions}}23__global__ void e2();24 25// Try all pairs of attributes which can be present on a function or a26// variable. Check both orderings of the attributes, as that can matter in27// clang.28__device__ __host__ void z1();29__device__ __constant__ int z2;30__device__ __shared__ int z3;31__device__ __global__ void z4(); // expected-error {{attributes are not compatible}}32// expected-note@-1 {{conflicting attribute is here}}33 34__host__ __device__ void z5();35__host__ __global__ void z6(); // expected-error {{attributes are not compatible}}36// expected-note@-1 {{conflicting attribute is here}}37 38__constant__ __device__ int z7;39__constant__ __shared__ int z8; // expected-error {{attributes are not compatible}}40// expected-note@-1 {{conflicting attribute is here}}41 42__shared__ __device__ int z9;43__shared__ __constant__ int z10; // expected-error {{attributes are not compatible}}44// expected-note@-1 {{conflicting attribute is here}}45__constant__ __shared__ int z10a; // expected-error {{attributes are not compatible}}46// expected-note@-1 {{conflicting attribute is here}}47 48__global__ __device__ void z11(); // expected-error {{attributes are not compatible}}49// expected-note@-1 {{conflicting attribute is here}}50__global__ __host__ void z12(); // expected-error {{attributes are not compatible}}51// expected-note@-1 {{conflicting attribute is here}}52 53// Make sure GPU-side variables do not allow __attribute((address_space(N)))54// expected-error@+1 {{__constant__, __device__, and __shared__ variables must use default address space}}55__shared__ __attribute__((address_space(999))) int as_s;56// expected-error@+1 {{__constant__, __device__, and __shared__ variables must use default address space}}57__device__ __attribute__((address_space(999))) int as_d;58// expected-error@+1 {{__constant__, __device__, and __shared__ variables must use default address space}}59__constant__ __attribute__((address_space(999))) int as_c;60 61struct S {62 __global__ void foo() {}; // expected-error {{must be a free function or static member function}}63 __global__ static void bar(); // expected-warning {{kernel function 'bar' is a member function}}64 // Although this is implicitly inline, we shouldn't warn.65 __global__ static void baz() {}; // expected-warning {{kernel function 'baz' is a member function}}66};67 68__global__ static inline void foobar() {};69#ifdef EXPECT_INLINE_WARNING70// expected-warning@-2 {{ignored 'inline' attribute on kernel function 'foobar'}}71#endif72 73__constant__ int global_constant;74void host_fn() {75 __constant__ int c; // expected-error {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}76 __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}}77}78__device__ void device_fn() {79 __constant__ int c; // expected-error {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}80}81 82typedef __attribute__((device_builtin_surface_type)) unsigned long long s0_ty; // expected-warning {{'device_builtin_surface_type' attribute only applies to classes}}83typedef __attribute__((device_builtin_texture_type)) unsigned long long t0_ty; // expected-warning {{'device_builtin_texture_type' attribute only applies to classes}}84 85struct __attribute__((device_builtin_surface_type)) s1_ref {}; // expected-error {{illegal device builtin surface reference type 's1_ref' declared here}}86// expected-note@-1 {{'s1_ref' needs to be instantiated from a class template with proper template arguments}}87struct __attribute__((device_builtin_texture_type)) t1_ref {}; // expected-error {{illegal device builtin texture reference type 't1_ref' declared here}}88// expected-note@-1 {{'t1_ref' needs to be instantiated from a class template with proper template arguments}}89 90template <typename T>91struct __attribute__((device_builtin_surface_type)) s2_cls_template {}; // expected-error {{illegal device builtin surface reference class template 's2_cls_template' declared here}}92// expected-note@-1 {{'s2_cls_template' needs to have exactly 2 template parameters}}93template <typename T>94struct __attribute__((device_builtin_texture_type)) t2_cls_template {}; // expected-error {{illegal device builtin texture reference class template 't2_cls_template' declared here}}95// expected-note@-1 {{'t2_cls_template' needs to have exactly 3 template parameters}}96 97template <int val, void *ptr>98struct __attribute__((device_builtin_surface_type)) s3_cls_template {}; // expected-error {{illegal device builtin surface reference class template 's3_cls_template' declared here}}99// expected-note@-1 {{the 1st template parameter of 's3_cls_template' needs to be a type}}100// expected-note@-2 {{the 2nd template parameter of 's3_cls_template' needs to be an integer or enum value}}101template <int val, int type, typename T>102struct __attribute__((device_builtin_texture_type)) t3_cls_template {}; // expected-error {{illegal device builtin texture reference class template 't3_cls_template' declared here}}103// expected-note@-1 {{the 1st template parameter of 't3_cls_template' needs to be a type}}104// expected-note@-2 {{the 3rd template parameter of 't3_cls_template' needs to be an integer or enum value}}105