| 1 | // RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify %s |
|---|
| 2 | // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s |
|---|
| 3 | |
|---|
| 4 | #include "Inputs/cuda.h" |
|---|
| 5 | |
|---|
| 6 | #ifndef __CUDA_ARCH__ |
|---|
| 7 | // expected-no-diagnostics |
|---|
| 8 | #endif |
|---|
| 9 | |
|---|
| 10 | // When compiling for device, foo()'s call to host_fn() is an error, because |
|---|
| 11 | // foo() is known-emitted. |
|---|
| 12 | // |
|---|
| 13 | // The trickiness here comes from the fact that the FunctionDecl bar() sees |
|---|
| 14 | // foo() does not have the "inline" keyword, so we might incorrectly think that |
|---|
| 15 | // foo() is a priori known-emitted. This would prevent us from marking foo() |
|---|
| 16 | // as known-emitted when we see the call from bar() to foo(), which would |
|---|
| 17 | // prevent us from emitting an error for foo()'s call to host_fn() when we |
|---|
| 18 | // eventually see it. |
|---|
| 19 | |
|---|
| 20 | void host_fn() {} |
|---|
| 21 | #ifdef __CUDA_ARCH__ |
|---|
| 22 | // expected-note@-2 {{declared here}} |
|---|
| 23 | #endif |
|---|
| 24 | |
|---|
| 25 | __host__ __device__ void foo(); |
|---|
| 26 | __device__ void bar() { |
|---|
| 27 | foo(); |
|---|
| 28 | #ifdef __CUDA_ARCH__ |
|---|
| 29 | // expected-note@-2 {{called by 'bar'}} |
|---|
| 30 | #endif |
|---|
| 31 | } |
|---|
| 32 | inline __host__ __device__ void foo() { |
|---|
| 33 | host_fn(); |
|---|
| 34 | #ifdef __CUDA_ARCH__ |
|---|
| 35 | // expected-error@-2 {{reference to __host__ function}} |
|---|
| 36 | #endif |
|---|
| 37 | } |
|---|
| 38 | |
|---|
| 39 | // This is similar to the above, except there's no error here. This code used |
|---|
| 40 | // to trip an assertion due to us noticing, when emitting the definition of |
|---|
| 41 | // boom(), that T::operator S() was (incorrectly) considered a priori |
|---|
| 42 | // known-emitted. |
|---|
| 43 | struct S {}; |
|---|
| 44 | struct T { |
|---|
| 45 | __device__ operator S() const; |
|---|
| 46 | }; |
|---|
| 47 | __device__ inline T::operator S() const { return S(); } |
|---|
| 48 | |
|---|
| 49 | __device__ T t; |
|---|
| 50 | __device__ void boom() { |
|---|
| 51 | S s = t; |
|---|
| 52 | } |
|---|
| 53 | |
|---|