75 lines · plain
1// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s2// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s \3// RUN: -fcuda-is-device4// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \5// RUN: -fopenmp %s6// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \7// RUN: -fopenmp %s -fcuda-is-device8 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.20struct 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 28constexpr int HostDevice() { return 0; }29 30// This should be a host-only function, because there's a previous __device__31// overload in <overload.h>.32constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }33 34namespace ns {35// The "using" statement in overload.h should prevent OverloadMe from being36// implicitly host+device.37constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }38} // namespace ns39 40// This is an error, because NonSysHdrOverload was not defined in a system41// header.42__device__ int NonSysHdrOverload() { return 0; }43// expected-note@-1 {{conflicting __device__ function declared here}}44constexpr 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 as48// host-only.49constexpr 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#else72 constexpr HostReturnTy y = OverloadMe();73#endif74}75