53 lines · plain
1// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify %s2// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s3 4#include "Inputs/cuda.h"5 6#ifndef __CUDA_ARCH__7// expected-no-diagnostics8#endif9 10// When compiling for device, foo()'s call to host_fn() is an error, because11// foo() is known-emitted.12//13// The trickiness here comes from the fact that the FunctionDecl bar() sees14// foo() does not have the "inline" keyword, so we might incorrectly think that15// 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 would17// prevent us from emitting an error for foo()'s call to host_fn() when we18// eventually see it.19 20void host_fn() {}21#ifdef __CUDA_ARCH__22 // expected-note@-2 {{declared here}}23#endif24 25__host__ __device__ void foo();26__device__ void bar() {27 foo();28#ifdef __CUDA_ARCH__29 // expected-note@-2 {{called by 'bar'}}30#endif31}32inline __host__ __device__ void foo() {33 host_fn();34#ifdef __CUDA_ARCH__35 // expected-error@-2 {{reference to __host__ function}}36#endif37}38 39// This is similar to the above, except there's no error here. This code used40// to trip an assertion due to us noticing, when emitting the definition of41// boom(), that T::operator S() was (incorrectly) considered a priori42// known-emitted.43struct S {};44struct 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