Mercurial > hg > CbC > CbC_llvm

comparison clang/test/SemaCUDA/host-device-constexpr.cu @ 150:1d019706d866

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
LLVM10
author anatofuz
date Thu, 13 Feb 2020 15:10:13 +0900
parents
children
comparison
equal deleted inserted replaced
147:c2174574ed3a 150:1d019706d866
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 }