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

// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s
// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s \
// RUN:            -fcuda-is-device
// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \
// RUN:            -fopenmp %s
// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \
// RUN:            -fopenmp %s -fcuda-is-device

#include "Inputs/cuda.h"

// Declares one function and pulls it into namespace ns:
//
//   __device__ int OverloadMe();
//   namespace ns { using ::OverloadMe; }
//
// Clang cares that this is done in a system header.
#include <overload.h>

// Opaque type used to determine which overload we're invoking.
struct HostReturnTy {};

// These shouldn't become host+device because they already have attributes.
__host__ constexpr int HostOnly() { return 0; }
// expected-note@-1 0+ {{not viable}}
__device__ constexpr int DeviceOnly() { return 0; }
// expected-note@-1 0+ {{not viable}}

constexpr int HostDevice() { return 0; }

// This should be a host-only function, because there's a previous __device__
// overload in <overload.h>.
constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }

namespace ns {
// The "using" statement in overload.h should prevent OverloadMe from being
// implicitly host+device.
constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }
}  // namespace ns

// This is an error, because NonSysHdrOverload was not defined in a system
// header.
__device__ int NonSysHdrOverload() { return 0; }
// expected-note@-1 {{conflicting __device__ function declared here}}
constexpr int NonSysHdrOverload() { return 0; }
// expected-error@-1 {{constexpr function 'NonSysHdrOverload' without __host__ or __device__ attributes}}

// Variadic device functions are not allowed, so this is just treated as
// host-only.
constexpr void Variadic(const char*, ...);
// expected-note@-1 {{call to __host__ function from __device__ function}}

__host__ void HostFn() {
  HostOnly();
  DeviceOnly(); // expected-error {{no matching function}}
  HostReturnTy x = OverloadMe();
  HostReturnTy y = ns::OverloadMe();
  Variadic("abc", 42);
}

__device__ void DeviceFn() {
  HostOnly(); // expected-error {{no matching function}}
  DeviceOnly();
  int x = OverloadMe();
  int y = ns::OverloadMe();
  Variadic("abc", 42); // expected-error {{no matching function}}
}

__host__ __device__ void HostDeviceFn() {
#ifdef __CUDA_ARCH__
  int y = OverloadMe();
#else
  constexpr HostReturnTy y = OverloadMe();
#endif
}