clang/test/SemaCUDA/host-device-constexpr.cu

   1 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s
   2 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s \
   3 // RUN:            -fcuda-is-device
   4 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \
   5 // RUN:            -fopenmp %s
   6 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \
   7 // RUN:            -fopenmp %s -fcuda-is-device
   8
   9 #include "Inputs/cuda.h"
  10
  11 // Declares one function and pulls it into namespace ns:
  12 //
  13 //   __device__ int OverloadMe();
  14 //   namespace ns { using ::OverloadMe; }
  15 //
  16 // Clang cares that this is done in a system header.
  17 #include <overload.h>
  18
  19 // Opaque type used to determine which overload we're invoking.
  20 struct HostReturnTy {};
  21
  22 // These shouldn't become host+device because they already have attributes.
  23 __host__ constexpr int HostOnly() { return 0; }
  24 // expected-note@-1 0+ {{not viable}}
  25 __device__ constexpr int DeviceOnly() { return 0; }
  26 // expected-note@-1 0+ {{not viable}}
  27
  28 constexpr int HostDevice() { return 0; }
  29
  30 // This should be a host-only function, because there's a previous __device__
  31 // overload in <overload.h>.
  32 constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }
  33
  34 namespace ns {
  35 // The "using" statement in overload.h should prevent OverloadMe from being
  36 // implicitly host+device.
  37 constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }
  38 }  // namespace ns
  39
  40 // This is an error, because NonSysHdrOverload was not defined in a system
  41 // header.
  42 __device__ int NonSysHdrOverload() { return 0; }
  43 // expected-note@-1 {{conflicting __device__ function declared here}}
  44 constexpr int NonSysHdrOverload() { return 0; }
  45 // expected-error@-1 {{constexpr function 'NonSysHdrOverload' without __host__ or __device__ attributes}}
  46
  47 // Variadic device functions are not allowed, so this is just treated as
  48 // host-only.
  49 constexpr void Variadic(const char*, ...);
  50 // expected-note@-1 {{call to __host__ function from __device__ function}}
  51
  52 __host__ void HostFn() {
  53   HostOnly();
  54   DeviceOnly(); // expected-error {{no matching function}}
  55   HostReturnTy x = OverloadMe();
  56   HostReturnTy y = ns::OverloadMe();
  57   Variadic("abc", 42);
  58 }
  59
  60 __device__ void DeviceFn() {
  61   HostOnly(); // expected-error {{no matching function}}
  62   DeviceOnly();
  63   int x = OverloadMe();
  64   int y = ns::OverloadMe();
  65   Variadic("abc", 42); // expected-error {{no matching function}}
  66 }
  67
  68 __host__ __device__ void HostDeviceFn() {
  69 #ifdef __CUDA_ARCH__
  70   int y = OverloadMe();
  71 #else
  72   constexpr HostReturnTy y = OverloadMe();
  73 #endif
  74 }