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
   73
   74
   75
   76
   77
   78
   79
   80
   81
   82
   83
   84
   85
   86
   87
   88
   89
   90
   91
   92
   93
   94
   95
   96
   97
   98
   99
  100
  101
  102
  103
  104
  105
  106
  107
  108
  109
  110
  111
  112
  113
  114
  115
  116
  117
  118
  119
  120
  121
  122
  123
  124
  125
  126
  127
  128
  129
  130
  131
  132
  133
  134
  135
  136
  137
  138
  139
  140

// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \
// RUN:   -emit-llvm -o /dev/null -verify -verify-ignore-unexpected=note

// Note: This test won't work with -fsyntax-only, because some of these errors
// are emitted during codegen.

#include "Inputs/cuda.h"

extern "C" void host_fn() {}
// expected-note@-1 7 {{'host_fn' declared here}}

struct Dummy {};

struct S {
  S() {}
  // expected-note@-1 2 {{'S' declared here}}
  ~S() { host_fn(); }
  // expected-note@-1 {{'~S' declared here}}
  int x;
};

struct T {
  __host__ __device__ void hd() { host_fn(); }
  // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}

  // No error; this is (implicitly) inline and is never called, so isn't
  // codegen'ed.
  __host__ __device__ void hd2() { host_fn(); }

  __host__ __device__ void hd3();

  void h() {}
  // expected-note@-1 {{'h' declared here}}

  void operator+();
  // expected-note@-1 {{'operator+' declared here}}

  void operator-(const T&) {}
  // expected-note@-1 {{'operator-' declared here}}

  operator Dummy() { return Dummy(); }
  // expected-note@-1 {{'operator Dummy' declared here}}

  __host__ void operator delete(void *) { host_fn(); };
  __device__ void operator delete(void*, __SIZE_TYPE__);
};

struct U {
  __device__ void operator delete(void*, __SIZE_TYPE__) = delete;
  __host__ __device__ void operator delete(void*);
};

__host__ __device__ void T::hd3() {
  host_fn();
  // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
}

template <typename T> __host__ __device__ void hd2() { host_fn(); }
// expected-error@-1 2 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
__global__ void kernel() { hd2<int>(); }

__host__ __device__ void hd() { host_fn(); }
// expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}

template <typename T> __host__ __device__ void hd3() { host_fn(); }
// expected-error@-1 2 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
__device__ void device_fn() { hd3<int>(); }

// No error because this is never instantiated.
template <typename T> __host__ __device__ void hd4() { host_fn(); }

__host__ __device__ void local_var() {
  S s;
  // expected-error@-1 {{reference to __host__ function 'S' in __host__ __device__ function}}
}

__host__ __device__ void placement_new(char *ptr) {
  ::new(ptr) S();
  // expected-error@-1 {{reference to __host__ function 'S' in __host__ __device__ function}}
}

__host__ __device__ void explicit_destructor(S *s) {
  s->~S();
  // expected-error@-1 {{reference to __host__ function '~S' in __host__ __device__ function}}
}

__host__ __device__ void class_specific_delete(T *t, U *u) {
  delete t; // ok, call sized device delete even though host has preferable non-sized version
  delete u; // ok, call non-sized HD delete rather than sized D delete
}

__host__ __device__ void hd_member_fn() {
  T t;
  // Necessary to trigger an error on T::hd.  It's (implicitly) inline, so
  // isn't codegen'ed until we call it.
  t.hd();
}

__host__ __device__ void h_member_fn() {
  T t;
  t.h();
  // expected-error@-1 {{reference to __host__ function 'h' in __host__ __device__ function}}
}

__host__ __device__ void fn_ptr() {
  auto* ptr = &host_fn;
  // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
}

template <typename T>
__host__ __device__ void fn_ptr_template() {
  auto* ptr = &host_fn;  // Not an error because the template isn't instantiated.
}

__host__ __device__ void unaryOp() {
  T t;
  (void) +t; // expected-error {{reference to __host__ function 'operator+' in __host__ __device__ function}}
}

__host__ __device__ void binaryOp() {
  T t;
  (void) (t - t); // expected-error {{reference to __host__ function 'operator-' in __host__ __device__ function}}
}

__host__ __device__ void implicitConversion() {
  T t;
  Dummy d = t; // expected-error {{reference to __host__ function 'operator Dummy' in __host__ __device__ function}}
}

template <typename T>
struct TmplStruct {
  template <typename U> __host__ __device__ void fn() {}
};

template <>
template <>
__host__ __device__ void TmplStruct<int>::fn<int>() { host_fn(); }
// expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}

__device__ void double_specialization() { TmplStruct<int>().fn<int>(); }