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
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
|
// Tests handling of CUDA attributes that are bad either because they're
// applied to the wrong sort of thing, or because they're given in illegal
// combinations.
//
// You should be able to run this file through nvcc for compatibility testing.
//
// RUN: %clang_cc1 -fsyntax-only -Wcuda-compat -verify -DEXPECT_INLINE_WARNING %s
// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -Wcuda-compat -verify %s
#include "Inputs/cuda.h"
// Try applying attributes to functions and variables. Some should generate
// warnings; others not.
__device__ int a1;
__device__ void a2();
__host__ int b1; // expected-warning {{attribute only applies to functions}}
__host__ void b2();
__constant__ int c1;
__constant__ void c2(); // expected-warning {{attribute only applies to variables}}
__shared__ int d1;
__shared__ void d2(); // expected-warning {{attribute only applies to variables}}
__global__ int e1; // expected-warning {{attribute only applies to functions}}
__global__ void e2();
// Try all pairs of attributes which can be present on a function or a
// variable. Check both orderings of the attributes, as that can matter in
// clang.
__device__ __host__ void z1();
__device__ __constant__ int z2;
__device__ __shared__ int z3;
__device__ __global__ void z4(); // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__host__ __device__ void z5();
__host__ __global__ void z6(); // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__constant__ __device__ int z7;
__constant__ __shared__ int z8; // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__shared__ __device__ int z9;
__shared__ __constant__ int z10; // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__constant__ __shared__ int z10a; // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__global__ __device__ void z11(); // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__global__ __host__ void z12(); // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
struct S {
__global__ void foo() {}; // expected-error {{must be a free function or static member function}}
__global__ static void bar(); // expected-warning {{kernel function 'bar' is a member function}}
// Although this is implicitly inline, we shouldn't warn.
__global__ static void baz() {}; // expected-warning {{kernel function 'baz' is a member function}}
};
__global__ static inline void foobar() {};
#ifdef EXPECT_INLINE_WARNING
// expected-warning@-2 {{ignored 'inline' attribute on kernel function 'foobar'}}
#endif
__constant__ int global_constant;
void host_fn() {
__constant__ int c; // expected-error {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
__shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}}
}
__device__ void device_fn() {
__constant__ int c; // expected-error {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
}
typedef __attribute__((device_builtin_surface_type)) unsigned long long s0_ty; // expected-warning {{'device_builtin_surface_type' attribute only applies to classes}}
typedef __attribute__((device_builtin_texture_type)) unsigned long long t0_ty; // expected-warning {{'device_builtin_texture_type' attribute only applies to classes}}
struct __attribute__((device_builtin_surface_type)) s1_ref {}; // expected-error {{illegal device builtin surface reference type 's1_ref' declared here}}
// expected-note@-1 {{'s1_ref' needs to be instantiated from a class template with proper template arguments}}
struct __attribute__((device_builtin_texture_type)) t1_ref {}; // expected-error {{illegal device builtin texture reference type 't1_ref' declared here}}
// expected-note@-1 {{'t1_ref' needs to be instantiated from a class template with proper template arguments}}
template <typename T>
struct __attribute__((device_builtin_surface_type)) s2_cls_template {}; // expected-error {{illegal device builtin surface reference class template 's2_cls_template' declared here}}
// expected-note@-1 {{'s2_cls_template' needs to have exactly 2 template parameters}}
template <typename T>
struct __attribute__((device_builtin_texture_type)) t2_cls_template {}; // expected-error {{illegal device builtin texture reference class template 't2_cls_template' declared here}}
// expected-note@-1 {{'t2_cls_template' needs to have exactly 3 template parameters}}
template <int val, void *ptr>
struct __attribute__((device_builtin_surface_type)) s3_cls_template {}; // expected-error {{illegal device builtin surface reference class template 's3_cls_template' declared here}}
// expected-note@-1 {{the 1st template parameter of 's3_cls_template' needs to be a type}}
// expected-note@-2 {{the 2nd template parameter of 's3_cls_template' needs to be an integer or enum value}}
template <int val, int type, typename T>
struct __attribute__((device_builtin_texture_type)) t3_cls_template {}; // expected-error {{illegal device builtin texture reference class template 't3_cls_template' declared here}}
// expected-note@-1 {{the 1st template parameter of 't3_cls_template' needs to be a type}}
// expected-note@-2 {{the 3rd template parameter of 't3_cls_template' needs to be an integer or enum value}}
|
