[CUDA] Tweak attribute-based overload resolution to match nvcc behavior.

[tra]

Summary

This is an artefact of split-mode CUDA compilation that we need to
mimic. HD functions are sometimes allowed to call H or D functions. Due
to split compilation mode device-side compilation will not see host-only
function and thus they will not be considered at all. For clang both H
and D variants will become function overloads visible to
compiler. Normally target attribute is considered only if C++ rules can
not determine which function is better. However in this case we need to
ignore functions that would not be present during current compilation
phase before we apply normal overload resolution rules.

Changes:

• introduced another level of call preference to better describe possible call combinations.
• removed WrongSide functions from consideration if the set contains SameSide function.
• disabled H->D and D->H and G->H calls. These combinations are not allowed by CUDA and we were reluctantly allowing them to work around device-side calls to math functions in std namespace. We no longer need it after r258880.

[jingyue]

include/clang/Sema/Sema.h

8805–8808

Why not merging CFP_SameSide and CFP_Native? You can conceptually treat HD as H in host mode or D in device mode. Then, it seems that we can just reuse CFP_Native to express the same meaning.

[tra]

When overload set contains h and HD functions that are otherwise equal for
overload resolution, you want to be able to tell which one is better.

[jingyue]

lib/Sema/SemaCUDA.cpp

126–140

Maybe a clearer way is to replace QuestionableResult with CFP_WrongSide in this if block, and after this if block say

if (DisabletargetCallChecks && CFP_WrongSide)
  return CFP_Never;

lib/Sema/SemaOverload.cpp

8544

Why can't we just return CFP1>CFP2? What's a counter example for that?

8548

You assume some numeric orders between these CFPs. May be worth adding a comment in CFP's definition.

[tra]

Addressed Jingyue's comments.
Fixed function-overload.cu tests to reflect stricter call target checks.

[tra]

lib/Sema/SemaCUDA.cpp

126–140

I've rearranged the code to make it easier to follow. I hope.

lib/Sema/SemaOverload.cpp

8544

You can look at it this way -- if both callees are viable during compilation (i.e. we can generate code for both) we want to use C++ overload resolution rules to pick best one (and can't return here). If one is viable and another is not, then the viable one always wins and that's what the check is for.

See template_vs_hd_function() in the test for example when we should not return here.

[jlebar]

Looks sane to me. Just some suggestions on the comments.

lib/Sema/SemaCUDA.cpp

71

If we're going to use symbols rather than letters, could we use 4, 3, 2, 1, 0? I think that would be easier to follow.

127–139

Nit: "It's OK to call a mode-matching function from an HD function."

lib/Sema/SemaOverload.cpp

8536

Since we have language lawyers on the team, suggest adding articles to comment:

If an HD function calls a function which has host-only and device-only overloads, nvcc sees only the host-side function during host compilation and only the device function during device-side compilation. (This appears to be a side-effect of its splitting of host and device code into separate TUs.) Alas we need to be compatible with existing code that relies on this, so if we see such a case, return the better variant right away.

I actually might suggest rephrasing this a bit more, to something like:

When performing host-side compilation, nvcc doesn't see device functions, and similarly when performing device-side compilation, nvcc doesn't see host functions. (This is a consequence of the fact that it splits host and device code into separate TUs.) We see all functions in both compilation modes, so to match nvcc's behavior, we need to exclude some overload candidates from consideration based only on their host/device attributes. Specifically, if one candidate call is WrongSide and the other is Native or SameSide, we ignore the WrongSide candidate. If we don't return early here, we'll consider the CUDA target attributes again later in this function, as a tiebreaker between calls with otherwise identical priority according to the regular C++ overloading rules.

test/CodeGenCUDA/function-overload.cu

97

call

[jingyue]

test/SemaCUDA/function-overload.cu

138–140

This line seems to have a trailing space.

[tra]

Updated the way WrongSide functions are removed from consideration during overload resolution.
Previous version could provide inconsistent results depending on the order of candidates in the overload set.
This version removes WrongSide functions from the set before we perform pair-wise comparison of candidates.

Added another test case to make sure HD->{H,H} or HD->{D,D} overloads are resolved consistently during host and device compilation.

[tra]

Previously accepted version was dependent on order of functions in overload set.
In order to make ordering consistent, WrongSide functions are now removed from the
set before pair-wise comparison of candidates. Please take a look.

lib/Sema/SemaCUDA.cpp

71

I considered that but found that numbers require more mental effort to map them back to what's supposed to happen for particular call combination. Letters were not sufficiently distinct visually (hard to tell l/f h/n at a glance). Current scheme works best for me. '+*o' == good. '-.' == bad.

[tra]

@jingyue, @jlebar: can you take a look at the updated version?

[jlebar]

Mostly comments on comments.

lib/Sema/SemaCUDA.cpp

71

What if we used the following mapping?

N = native
HD = host+device
SS = same-side
WS = wrong-side
- = never

This mimics how we were writing on the whiteboard.

115–119

I'm not sure "fallback" is the right word to use here anymore, as HD --> HD gets priority HD.

127–139

Suggest "compilation-mode matching"

lib/Sema/SemaOverload.cpp

8731

We should be consistent wrt whether we call it a compilation pass, or (as above) compilation mode, or whatever. (I think "mode" may be right.)

Also please add an article: "During a particular compilation mode".

8738

Maybe name this something more closely tied to what it is -- e.g. ContainsSameSideCall -- and then add a comment in the if statement saying "Remove wrong-side calls from consideration."

8752

This is an indentation mouthful -- can we pull the lambda out, maybe?

8758

Suggest auto*, so it's clear we're not copying things.

8772

auto*

test/CodeGenCUDA/function-overload.cu

81

Nit: Suggest American spelling, "artifact", which is much more common in llvm codebase.

Maybe also remove this sentence, or move it down somewhere later -- this feels like a bad "topic sentence" for the paragraph. e.g.

HD functions are sometimes allowed to call H or D functions -- this is an artifact of the source-to-source splitting performed by nvcc that we need to mimic.

86

This is setting up a contrast between nvcc and clang, so suggest connecting the phrases with "but" or "in contrast". Also suggest being specific that we're talking about nvcc -- since split-mode compilation isn't a thing in clang, if we just talk about it generally, it's not clear what we're referring to. e.g.

During device mode compilation in nvcc, host functions aren't present at all, so don't participate in overloading. But in clang, H and D functions are present in both compilation modes. Clang normally uses the target attribute as a tiebreaker between overloads with otherwise identical priority, but in order to match nvcc's behavior, we sometimes need to wholly discard overloads that would not be present during compilation under nvcc.

98

This was a bit hard to follow, suggest rewording slightly:

Here we expect to call the templated function during host compilation, even if -fcuda-disable-target-call-checks is passed, and even though C++ overload rules prefer the non-templated function.

110

call the overloaded function

128

If we have a mix

in the overload set, normal

162

when the overload set

[jingyue]

I'll defer to Justin's approval.

[tra]

Addressed @jlebar's comments.

[jlebar]

lgtm

include/clang/Sema/Sema.h

enum CUDAFunctionTarget {
     CFT_Global,
     CFT_Host,
     CFT_HostDevice,
     CFT_InvalidTarget
   };
 
   CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D);
 
+  // CUDA function call preference. Must be ordered numerically from
+  // worst to best.
   enum CUDAFunctionPreference {
     CFP_Never,      // Invalid caller/callee combination.
-    CFP_LastResort, // Lowest priority. Only in effect if
+    CFP_WrongSide,  // Calls from host-device to host or device
+                    // function that do not match current compilation
+                    // mode. Only in effect if
                     // LangOpts.CUDADisableTargetCallChecks is true.
-    CFP_Fallback,   // Low priority caller/callee combination
-    CFP_Best,       // Preferred caller/callee combination
+    CFP_HostDevice, // Any calls to host/device functions.
+    CFP_SameSide,   // Calls from host-device to host or device
+                    // function matching current compilation mode.
+    CFP_Native,     // host-to-host or device-to-device calls.
   };
 
   /// Identifies relative preference of a given Caller/Callee
   /// combination, based on their host/device attributes.
   /// \param Caller function which needs address of \p Callee.
   ///               nullptr in case of global context.
   /// \param Callee target function
   ///

lib/Sema/SemaCUDA.cpp

 
 // * CUDA Call preference table
 //
 // F - from,
 // T - to
 // Ph - preference in host mode
 // Pd - preference in device mode
 // H  - handled in (x)
-// Preferences: b-best, f-fallback, l-last resort, n-never.
+// Preferences: N:native, HD:host-device, SS:same side, WS:wrong side, --:never.
 //
-// | F  | T  | Ph | Pd |  H  |
-// |----+----+----+----+-----+
-// | d  | d  | b  | b  | (b) |
-// | d  | g  | n  | n  | (a) |
-// | d  | h  | l  | l  | (e) |
-// | d  | hd | f  | f  | (c) |
-// | g  | d  | b  | b  | (b) |
-// | g  | g  | n  | n  | (a) |
-// | g  | h  | l  | l  | (e) |
-// | g  | hd | f  | f  | (c) |
-// | h  | d  | l  | l  | (e) |
-// | h  | g  | b  | b  | (b) |
-// | h  | h  | b  | b  | (b) |
-// | h  | hd | f  | f  | (c) |
-// | hd | d  | l  | f  | (d) |
-// | hd | g  | f  | n  |(d/a)|
-// | hd | h  | f  | l  | (d) |
-// | hd | hd | b  | b  | (b) |
+// | F  | T  | Ph  | Pd  |  H  |
+// |----+----+-----+-----+-----+
+// | d  | d  | N   | N   | (c) |
+// | d  | g  | --  | --  | (a) |
+// | d  | h  | --  | --  | (e) |
+// | d  | hd | HD  | HD  | (b) |
+// | g  | d  | N   | N   | (c) |
+// | g  | g  | --  | --  | (a) |
+// | g  | h  | --  | --  | (e) |
+// | g  | hd | HD  | HD  | (b) |
+// | h  | d  | --  | --  | (e) |
+// | h  | g  | N   | N   | (c) |
+// | h  | h  | N   | N   | (c) |
+// | h  | hd | HD  | HD  | (b) |
+// | hd | d  | WS  | SS  | (d) |
+// | hd | g  | SS  | --  |(d/a)|
+// | hd | h  | SS  | WS  | (d) |
+// | hd | hd | HD  | HD  | (b) |
 
 Sema::CUDAFunctionPreference
 Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
                              const FunctionDecl *Callee) {
   assert(getLangOpts().CUDATargetOverloads &&
          "Should not be called w/o enabled target overloads.");
 
   assert(Callee && "Callee must be valid.");
   CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee);
   CUDAFunctionTarget CallerTarget =
       (Caller != nullptr) ? IdentifyCUDATarget(Caller) : Sema::CFT_Host;
 
   // If one of the targets is invalid, the check always fails, no matter what
   // the other target is.
   if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
     return CFP_Never;
 
   // (a) Can't call global from some contexts until we support CUDA's
   // dynamic parallelism.
   if (CalleeTarget == CFT_Global &&
       (CallerTarget == CFT_Global || CallerTarget == CFT_Device ||
        (CallerTarget == CFT_HostDevice && getLangOpts().CUDAIsDevice)))
     return CFP_Never;
 
-  // (b) Best case scenarios
+  // (b) Calling HostDevice is OK for everyone.
+  if (CalleeTarget == CFT_HostDevice)
+    return CFP_HostDevice;
+
+  // (c) Best case scenarios
   if (CalleeTarget == CallerTarget ||
       (CallerTarget == CFT_Host && CalleeTarget == CFT_Global) ||
       (CallerTarget == CFT_Global && CalleeTarget == CFT_Device))
-    return CFP_Best;
-
-  // (c) Calling HostDevice is OK as a fallback that works for everyone.
-  if (CalleeTarget == CFT_HostDevice)
-    return CFP_Fallback;
-
-  // Figure out what should be returned 'last resort' cases. Normally
-  // those would not be allowed, but we'll consider them if
-  // CUDADisableTargetCallChecks is true.
-  CUDAFunctionPreference QuestionableResult =
-      getLangOpts().CUDADisableTargetCallChecks ? CFP_LastResort : CFP_Never;
+    return CFP_Native;
 
   // (d) HostDevice behavior depends on compilation mode.
   if (CallerTarget == CFT_HostDevice) {
-    // Calling a function that matches compilation mode is OK.
-    // Calling a function from the other side is frowned upon.
-    if (getLangOpts().CUDAIsDevice)
-      return CalleeTarget == CFT_Device ? CFP_Fallback : QuestionableResult;
-    else
-      return (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)
-                 ? CFP_Fallback
-                 : QuestionableResult;
+    // It's OK to call a compilation-mode matching function from an HD one.
+    if ((getLangOpts().CUDAIsDevice && CalleeTarget == CFT_Device) ||
+        (!getLangOpts().CUDAIsDevice &&
+         (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)))
+      return CFP_SameSide;
+
+    // We'll allow calls to non-mode-matching functions if target call
+    // checks are disabled. This is needed to avoid complaining about
+    // HD->H calls when we compile for device side and vice versa.
+    if (getLangOpts().CUDADisableTargetCallChecks)
+      return CFP_WrongSide;
+
+    return CFP_Never;
   }
 
   // (e) Calling across device/host boundary is not something you should do.
   if ((CallerTarget == CFT_Host && CalleeTarget == CFT_Device) ||
       (CallerTarget == CFT_Device && CalleeTarget == CFT_Host) ||
       (CallerTarget == CFT_Global && CalleeTarget == CFT_Host))
-    return QuestionableResult;
+    return CFP_Never;
 
   llvm_unreachable("All cases should've been handled by now.");
 }
 
 bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
                            const FunctionDecl *Callee) {
   // With target overloads enabled, we only disallow calling
   // combinations with CFP_Never.

lib/Sema/SemaOverload.cpp

else if (!Cand1.Viable)
     return false;
 
   // C++ [over.match.best]p1:
   //
   //   -- if F is a static member function, ICS1(F) is defined such
   //      that ICS1(F) is neither better nor worse than ICS1(G) for
   //      any function G, and, symmetrically, ICS1(G) is neither
   //      better nor worse than ICS1(F).
   unsigned StartArg = 0;
   if (Cand1.IgnoreObjectArgument || Cand2.IgnoreObjectArgument)
     StartArg = 1;
 
   // C++ [over.match.best]p1:
   //   A viable function F1 is defined to be a better function than another
   //   viable function F2 if for all arguments i, ICSi(F1) is not a worse
   //   conversion sequence than ICSi(F2), and then...
   unsigned NumArgs = Cand1.NumConversions;
   assert(Cand2.NumConversions == NumArgs && "Overload candidate mismatch");
   bool HasBetterConversion = false;
   for (unsigned ArgIdx = StartArg; ArgIdx < NumArgs; ++ArgIdx) {
     switch (CompareImplicitConversionSequences(S, Loc,
                                                Cand1.Conversions[ArgIdx],
                                                Cand2.Conversions[ArgIdx])) {
     case ImplicitConversionSequence::Better:
       // Cand1 has a better conversion sequence.
       HasBetterConversion = true;
       break;
 
     case ImplicitConversionSequence::Worse:

 /// \param Best If overload resolution was successful or found a deleted
 /// function, \p Best points to the candidate function found.
 ///
 /// \returns The result of overload resolution.
 OverloadingResult
 OverloadCandidateSet::BestViableFunction(Sema &S, SourceLocation Loc,
                                          iterator &Best,
                                          bool UserDefinedConversion) {
+  llvm::SmallVector<OverloadCandidate *, 16> Candidates;
+  std::transform(begin(), end(), std::back_inserter(Candidates),
+                 [](OverloadCandidate &Cand) { return &Cand; });
+
+  // [CUDA] HD->H or HD->D calls are technically not allowed by CUDA
+  // but accepted by both clang and NVCC. However during a particular
+  // compilation mode only one call variant is viable. We need to
+  // exclude non-viable overload candidates from consideration based
+  // only on their host/device attributes. Specifically, if one
+  // candidate call is WrongSide and the other is SameSide, we ignore
+  // the WrongSide candidate.
+  if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads) {
+    const FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
+    bool ContainsSameSideCandidate =
+        llvm::any_of(Candidates, [&](OverloadCandidate *Cand) {
+          return Cand->Function &&
+                 S.IdentifyCUDAPreference(Caller, Cand->Function) ==
+                     Sema::CFP_SameSide;
+        });
+    if (ContainsSameSideCandidate) {
+      auto IsWrongSideCandidate = [&](OverloadCandidate *Cand) {
+        return Cand->Function &&
+               S.IdentifyCUDAPreference(Caller, Cand->Function) ==
+                   Sema::CFP_WrongSide;
+      };
+      Candidates.erase(std::remove_if(Candidates.begin(), Candidates.end(),
+                                      IsWrongSideCandidate),
+                       Candidates.end());
+    }
+  }
+
   // Find the best viable function.
   Best = end();
-  for (iterator Cand = begin(); Cand != end(); ++Cand) {
+  for (auto *Cand : Candidates)
     if (Cand->Viable)
       if (Best == end() || isBetterOverloadCandidate(S, *Cand, *Best, Loc,
                                                      UserDefinedConversion))
         Best = Cand;
-  }
 
   // If we didn't find any viable functions, abort.
   if (Best == end())
     return OR_No_Viable_Function;
 
   llvm::SmallVector<const NamedDecl *, 4> EquivalentCands;
 
   // Make sure that this function is better than every other viable
   // function. If not, we have an ambiguity.
-  for (iterator Cand = begin(); Cand != end(); ++Cand) {
+  for (auto *Cand : Candidates) {
     if (Cand->Viable &&
         Cand != Best &&
         !isBetterOverloadCandidate(S, *Best, *Cand, Loc,
                                    UserDefinedConversion)) {
       if (S.isEquivalentInternalLinkageDeclaration(Best->Function,
                                                    Cand->Function)) {
         EquivalentCands.push_back(Cand->Function);
         continue;

test/CodeGenCUDA/function-overload.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
 // Make sure we handle target overloads correctly.
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
 // RUN:     -fcuda-target-overloads -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
 // RUN:     -fcuda-target-overloads -emit-llvm -o - %s \
-// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
+// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \
+// RUN:       -check-prefix=CHECK-DEVICE-STRICT %s
 
 // Check target overloads handling with disabled call target checks.
 // RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -emit-llvm \
 // RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST \
 // RUN:    -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-HOST-NC %s
 // RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -emit-llvm \
 // RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads \

 // CHECK-HOST-LABEL: define i32 @ch()
 extern "C" __host__ int ch(void) {return 13;}
 // CHECK-HOST:     ret i32 13
 
 // CHECK-BOTH-LABEL: define i32 @chd()
 extern "C" __host__ __device__ int chd(void) {return 14;}
 // CHECK-BOTH:     ret i32 14
 
+// HD functions are sometimes allowed to call H or D functions -- this
+// is an artifact of the source-to-source splitting performed by nvcc
+// that we need to mimic. During device mode compilation in nvcc, host
+// functions aren't present at all, so don't participate in
+// overloading. But in clang, H and D functions are present in both
+// compilation modes. Clang normally uses the target attribute as a
+// tiebreaker between overloads with otherwise identical priority, but
+// in order to match nvcc's behavior, we sometimes need to wholly
+// discard overloads that would not be present during compilation
+// under nvcc.
+
+template <typename T> T template_vs_function(T arg) { return 15; }
+__device__ float template_vs_function(float arg) { return 16; }
+
+// Here we expect to call the templated function during host
+// compilation, even if -fcuda-disable-target-call-checks is passed,
+// and even though C++ overload rules prefer the non-templated
+// function.
+// CHECK-BOTH-LABEL: define void @_Z5hd_tfv()
+__host__ __device__ void hd_tf(void) {
+  template_vs_function(1.0f);
+  // CHECK-HOST: call float @_Z20template_vs_functionIfET_S0_(float
+  // CHECK-DEVICE: call float @_Z20template_vs_functionf(float
+  template_vs_function(2.0);
+  // CHECK-HOST: call double @_Z20template_vs_functionIdET_S0_(double
+  // CHECK-DEVICE: call float @_Z20template_vs_functionf(float
+}
+
+// Calls from __host__ and __device__ functions should always call the
+// overloaded function that matches their mode.
+// CHECK-HOST-LABEL: define void @_Z4h_tfv()
+__host__ void h_tf() {
+  template_vs_function(1.0f);
+  // CHECK-HOST: call float @_Z20template_vs_functionIfET_S0_(float
+  template_vs_function(2.0);
+  // CHECK-HOST: call double @_Z20template_vs_functionIdET_S0_(double
+}
+
+// CHECK-DEVICE-LABEL: define void @_Z4d_tfv()
+__device__ void d_tf() {
+  template_vs_function(1.0f);
+  // CHECK-DEVICE: call float @_Z20template_vs_functionf(float
+  template_vs_function(2.0);
+  // CHECK-DEVICE: call float @_Z20template_vs_functionf(float
+}
+
+// In case we have a mix of HD and H-only or D-only candidates in the
+// overload set, normal C++ overload resolution rules apply first.
+template <typename T> T template_vs_hd_function(T arg) { return 15; }
+__host__ __device__ float template_vs_hd_function(float arg) { return 16; }
+
+// CHECK-BOTH-LABEL: define void @_Z7hd_thdfv()
+__host__ __device__ void hd_thdf() {
+  template_vs_hd_function(1.0f);
+  // CHECK-HOST: call float @_Z23template_vs_hd_functionf(float
+  // CHECK-DEVICE: call float @_Z23template_vs_hd_functionf(float
+  template_vs_hd_function(1);
+  // CHECK-HOST: call i32 @_Z23template_vs_hd_functionIiET_S0_(i32
+  // CHECK-DEVICE-STRICT: call float @_Z23template_vs_hd_functionf(float
+  // CHECK-DEVICE-NC: call i32 @_Z23template_vs_hd_functionIiET_S0_(i32
+}
+
+// CHECK-HOST-LABEL: define void @_Z6h_thdfv()
+__host__ void h_thdf() {
+  template_vs_hd_function(1.0f);
+  // CHECK-HOST: call float @_Z23template_vs_hd_functionf(float
+  template_vs_hd_function(1);
+  // CHECK-HOST: call i32 @_Z23template_vs_hd_functionIiET_S0_(i32
+}
+
+// CHECK-DEVICE-LABEL: define void @_Z6d_thdfv()
+__device__ void d_thdf() {
+  template_vs_hd_function(1.0f);
+  // CHECK-DEVICE: call float @_Z23template_vs_hd_functionf(float
+  template_vs_hd_function(1);
+  // Host-only function template is not callable with strict call checks,
+  // so for device side HD function will be the only choice.
+  // CHECK-DEVICE: call float @_Z23template_vs_hd_functionf(float
+}
+
+// Check that overloads still work the same way on both host and
+// device side when the overload set contains only functions from one
+// side of compilation.
+__device__ float device_only_function(int arg) { return 17; }
+__device__ float device_only_function(float arg) { return 18; }
+
+__host__ float host_only_function(int arg) { return 19; }
+__host__ float host_only_function(float arg) { return 20; }
+
+// CHECK-BOTH-LABEL: define void @_Z6hd_dofv()
+__host__ __device__ void hd_dof() {
+#ifdef NOCHECKS
+  device_only_function(1.0f);
+  // CHECK-BOTH-NC: call float @_Z20device_only_functionf(float
+  device_only_function(1);
+  // CHECK-BOTH-NC: call float @_Z20device_only_functioni(i32
+  host_only_function(1.0f);
+  // CHECK-BOTH-NC: call float @_Z18host_only_functionf(float
+  host_only_function(1);
+  // CHECK-BOTH-NC: call float @_Z18host_only_functioni(i32
+#endif
+}
+
+
 // CHECK-HOST-LABEL: define void @_Z5hostfv()
 __host__ void hostf(void) {
-#if defined (NOCHECKS)
-  fp_t dp = d;   // CHECK-HOST-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
-  fp_t cdp = cd; // CHECK-HOST-NC: store {{.*}} @cd, {{.*}} %cdp,
-#endif
   fp_t hp = h; // CHECK-HOST: store {{.*}} @_Z1hv, {{.*}} %hp,
   fp_t chp = ch; // CHECK-HOST: store {{.*}} @ch, {{.*}} %chp,
   fp_t dhp = dh; // CHECK-HOST: store {{.*}} @_Z2dhv, {{.*}} %dhp,
   fp_t cdhp = cdh; // CHECK-HOST: store {{.*}} @cdh, {{.*}} %cdhp,
   fp_t hdp = hd; // CHECK-HOST: store {{.*}} @_Z2hdv, {{.*}} %hdp,
   fp_t chdp = chd; // CHECK-HOST: store {{.*}} @chd, {{.*}} %chdp,
   gp_t gp = g; // CHECK-HOST: store {{.*}} @_Z1gv, {{.*}} %gp,
 
-#if defined (NOCHECKS)
-  d();     // CHECK-HOST-NC: call i32 @_Z1dv()
-  cd();    // CHECK-HOST-NC: call i32 @cd()
-#endif
   h();     // CHECK-HOST: call i32 @_Z1hv()
   ch();    // CHECK-HOST: call i32 @ch()
   dh();    // CHECK-HOST: call i32 @_Z2dhv()
   cdh();   // CHECK-HOST: call i32 @cdh()
   g<<<0,0>>>();  // CHECK-HOST: call void @_Z1gv()
 }
 
 // CHECK-DEVICE-LABEL: define void @_Z7devicefv()
 __device__ void devicef(void) {
   fp_t dp = d;   // CHECK-DEVICE: store {{.*}} @_Z1dv, {{.*}} %dp,
   fp_t cdp = cd; // CHECK-DEVICE: store {{.*}} @cd, {{.*}} %cdp,
-#if defined (NOCHECKS)
-  fp_t hp = h; // CHECK-DEVICE-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
-  fp_t chp = ch; // CHECK-DEVICE-NC: store {{.*}} @ch, {{.*}} %chp,
-#endif
   fp_t dhp = dh; // CHECK-DEVICE: store {{.*}} @_Z2dhv, {{.*}} %dhp,
   fp_t cdhp = cdh; // CHECK-DEVICE: store {{.*}} @cdh, {{.*}} %cdhp,
   fp_t hdp = hd; // CHECK-DEVICE: store {{.*}} @_Z2hdv, {{.*}} %hdp,
   fp_t chdp = chd; // CHECK-DEVICE: store {{.*}} @chd, {{.*}} %chdp,
 
   d();     // CHECK-DEVICE: call i32 @_Z1dv()
   cd();    // CHECK-DEVICE: call i32 @cd()
-#if defined (NOCHECKS)
-  h();     // CHECK-DEVICE-NC: call i32 @_Z1hv()
-  ch();    // CHECK-DEVICE-NC: call i32 @ch()
-#endif
   dh();    // CHECK-DEVICE: call i32 @_Z2dhv()
   cdh();   // CHECK-DEVICE: call i32 @cdh()
 }
 
 // CHECK-BOTH-LABEL: define void @_Z11hostdevicefv()
 __host__ __device__ void hostdevicef(void) {
 #if defined (NOCHECKS)
   fp_t dp = d;   // CHECK-BOTH-NC: store {{.*}} @_Z1dv, {{.*}} %dp,

test/SemaCUDA/function-overload.cu

 __device__ int d(void) { return 8; }
 __host__ int h(void) { return 9; }
 __global__ void g(void) {}
 extern "C" __device__ int cd(void) {return 10;}
 extern "C" __host__ int ch(void) {return 11;}
 
 __host__ void hostf(void) {
   fp_t dp = d;
-  fp_t cdp = cd;
-#if !defined(NOCHECKS)
-  // expected-error@-3 {{reference to __device__ function 'd' in __host__ function}}
+  // expected-error@-1 {{reference to __device__ function 'd' in __host__ function}}
   // expected-note@65 {{'d' declared here}}
-  // expected-error@-4 {{reference to __device__ function 'cd' in __host__ function}}
+  fp_t cdp = cd;
+  // expected-error@-1 {{reference to __device__ function 'cd' in __host__ function}}
   // expected-note@68 {{'cd' declared here}}
-#endif
   fp_t hp = h;
   fp_t chp = ch;
   fp_t dhp = dh;
   fp_t cdhp = cdh;
   gp_t gp = g;
 
   d();
-  cd();
-#if !defined(NOCHECKS)
-  // expected-error@-3 {{no matching function for call to 'd'}}
+  // expected-error@-1 {{no matching function for call to 'd'}}
   // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ function}}
-  // expected-error@-4 {{no matching function for call to 'cd'}}
+  cd();
+  // expected-error@-1 {{no matching function for call to 'cd'}}
   // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ function}}
-#endif
   h();
   ch();
   dh();
   cdh();
   g(); // expected-error {{call to global function g not configured}}
   g<<<0,0>>>();
 }
 
 
 __device__ void devicef(void) {
   fp_t dp = d;
   fp_t cdp = cd;
   fp_t hp = h;
-  fp_t chp = ch;
-#if !defined(NOCHECKS)
-  // expected-error@-3 {{reference to __host__ function 'h' in __device__ function}}
+  // expected-error@-1 {{reference to __host__ function 'h' in __device__ function}}
   // expected-note@66 {{'h' declared here}}
-  // expected-error@-4 {{reference to __host__ function 'ch' in __device__ function}}
+  fp_t chp = ch;
+  // expected-error@-1 {{reference to __host__ function 'ch' in __device__ function}}
   // expected-note@69 {{'ch' declared here}}
-#endif
   fp_t dhp = dh;
   fp_t cdhp = cdh;
   gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
                // expected-note@67 {{'g' declared here}}
 
   d();
   cd();
-  h();
-  ch();
-#if !defined(NOCHECKS)
-  // expected-error@-3 {{no matching function for call to 'h'}}
+  h(); // expected-error {{no matching function for call to 'h'}}
   // expected-note@66 {{candidate function not viable: call to __host__ function from __device__ function}}
-  // expected-error@-4 {{no matching function for call to 'ch'}}
+  ch(); // expected-error {{no matching function for call to 'ch'}}
   // expected-note@69 {{candidate function not viable: call to __host__ function from __device__ function}}
-#endif
   dh();
   cdh();
   g(); // expected-error {{no matching function for call to 'g'}}
   // expected-note@67 {{candidate function not viable: call to __global__ function from __device__ function}}
   g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}}
   // expected-note@67 {{'g' declared here}}
 }
 
 __global__ void globalf(void) {
   fp_t dp = d;
   fp_t cdp = cd;
   fp_t hp = h;
-  fp_t chp = ch;
-#if !defined(NOCHECKS)
-  // expected-error@-3 {{reference to __host__ function 'h' in __global__ function}}
+  // expected-error@-1 {{reference to __host__ function 'h' in __global__ function}}
   // expected-note@66 {{'h' declared here}}
-  // expected-error@-4 {{reference to __host__ function 'ch' in __global__ function}}
+  fp_t chp = ch;
+  // expected-error@-1 {{reference to __host__ function 'ch' in __global__ function}}
   // expected-note@69 {{'ch' declared here}}
-#endif
   fp_t dhp = dh;
   fp_t cdhp = cdh;
-  gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __global__ function}}
-               // expected-note@67 {{'g' declared here}}
+  gp_t gp = g;
+  // expected-error@-1 {{reference to __global__ function 'g' in __global__ function}}
 
   d();
   cd();
   h();
-  ch();
-#if !defined(NOCHECKS)
-  // expected-error@-3 {{no matching function for call to 'h'}}
+  // expected-error@-1 {{no matching function for call to 'h'}}
   // expected-note@66 {{candidate function not viable: call to __host__ function from __global__ function}}
-  // expected-error@-4 {{no matching function for call to 'ch'}}
+  ch();
+  // expected-error@-1 {{no matching function for call to 'ch'}}
   // expected-note@69 {{candidate function not viable: call to __host__ function from __global__ function}}
-#endif
   dh();
   cdh();
   g(); // expected-error {{no matching function for call to 'g'}}
   // expected-note@67 {{candidate function not viable: call to __global__ function from __global__ function}}
   g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}}
   // expected-note@67 {{'g' declared here}}
 }