reference, declaration → definition
definition → references, declarations, derived classes, virtual overrides
reference to multiple definitions → definitions
unreferenced

    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

// 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__ variables must be global}}
  __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}}
}
__device__ void device_fn() {
  __constant__ int c; // expected-error {{__constant__ variables must be global}}
}