[CUDA] Use only the GVALinkage on function definitions.

Summary:
Previously we'd look at the GVALinkage of whatever FunctionDecl you
happened to be calling.

This is not right.  In the absence of the gnu_inline attribute, to be
handled separately, the function definition determines the function's
linkage.  So we need to wait until we get a def before we can know
whether something is known-emitted.

Reviewers: tra

Subscribers: cfe-commits, rsmith

Differential Revision: https://reviews.llvm.org/D26268

llvm-svn: 286313

1 files changed, 52 insertions, 0 deletions

diff --git a/clang/test/SemaCUDA/add-inline-in-definition.cu b/clang/test/SemaCUDA/add-inline-in-definition.cu
new file mode 100644
index 00000000000..33176c580d4
--- /dev/null
+++ b/clang/test/SemaCUDA/add-inline-in-definition.cu
@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+#ifndef __CUDA_ARCH__
+// expected-no-diagnostics
+#endif
+
+// When compiling for device, foo()'s call to host_fn() is an error, because
+// foo() is known-emitted.
+//
+// The trickiness here comes from the fact that the FunctionDecl bar() sees
+// foo() does not have the "inline" keyword, so we might incorrectly think that
+// foo() is a priori known-emitted.  This would prevent us from marking foo()
+// as known-emitted when we see the call from bar() to foo(), which would
+// prevent us from emitting an error for foo()'s call to host_fn() when we
+// eventually see it.
+
+void host_fn() {}
+#ifdef __CUDA_ARCH__
+  // expected-note@-2 {{declared here}}
+#endif
+
+__host__ __device__ void foo();
+__device__ void bar() {
+  foo();
+#ifdef __CUDA_ARCH__
+  // expected-note@-2 {{called by 'bar'}}
+#endif
+}
+inline __host__ __device__ void foo() {
+  host_fn();
+#ifdef __CUDA_ARCH__
+  // expected-error@-2 {{reference to __host__ function}}
+#endif
+}
+
+// This is similar to the above, except there's no error here.  This code used
+// to trip an assertion due to us noticing, when emitting the definition of
+// boom(), that T::operator S() was (incorrectly) considered a priori
+// known-emitted.
+struct S {};
+struct T {
+  __device__ operator S() const;
+};
+__device__ inline T::operator S() const { return S(); }
+
+__device__ T t;
+__device__ void boom() {
+  S s = t;
+}