[CUDA] Use only the GVALinkage on function definitions.

[jlebar]

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.

[jlebar]

Friendly ping

[tra]

LGTM.

clang/test/SemaCUDA/add-inline-in-definition.cu

13–14 ↗

(On Diff #76808)

"bar() sees for foo()" does not sound right. Did you mean "bar() sees *that* foo()"?

15 ↗

(On Diff #76808)

Extra whitespace.

[jlebar]

Thank you for the review! Submitting...

clang/test/SemaCUDA/add-inline-in-definition.cu

13–14 ↗

(On Diff #76808)

Yes, this is a mistake, thanks!

15 ↗

(On Diff #76808)

We seem to be pretty evenly-split wrt one or two spaces here in llvm, agreed irl it's nbd.

cfe/trunk/lib/Sema/SemaCUDA.cpp

static bool IsKnownEmitted(Sema &S, FunctionDecl *FD) {
   // doesn't count for our purposes here.)
   Sema::CUDAFunctionTarget T = S.IdentifyCUDATarget(FD);
   if (S.getLangOpts().CUDAIsDevice && T == Sema::CFT_Host)
     return false;
   if (!S.getLangOpts().CUDAIsDevice &&
       (T == Sema::CFT_Device || T == Sema::CFT_Global))
     return false;
 
-  // Externally-visible and similar functions are always emitted.
-  if (!isDiscardableGVALinkage(S.getASTContext().GetGVALinkageForFunction(FD)))
+  // Check whether this function is externally visible -- if so, it's
+  // known-emitted.
+  //
+  // We have to check the GVA linkage of the function's *definition* -- if we
+  // only have a declaration, we don't know whether or not the function will be
+  // emitted, because (say) the definition could include "inline".
+  FunctionDecl *Def = FD->getDefinition();
+
+  // We may currently be parsing the body of FD, in which case
+  // FD->getDefinition() will be null, but we still want to treat FD as though
+  // it's a definition.
+  if (!Def && FD->willHaveBody())
+    Def = FD;
+
+  if (Def &&
+      !isDiscardableGVALinkage(S.getASTContext().GetGVALinkageForFunction(Def)))
     return true;
 
   // Otherwise, the function is known-emitted if it's in our set of
   // known-emitted functions.
   return S.CUDAKnownEmittedFns.count(FD) > 0;
 }
 
 Sema::CUDADiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,

cfe/trunk/test/SemaCUDA/add-inline-in-definition.cu

+// 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;
+}