[PATCH] D12453: [CUDA] Allow function overloads based on host/device attributes.

[Artem Belevich via cfe-commits]

tra marked 2 inline comments as done.

================
Comment at: lib/Sema/SemaCUDA.cpp:106
@@ +105,3 @@
+
+  // (a) Can't call global from global until we support dynamic execution.
+  if (CalleeTarget == CFT_Global &&
----------------
eliben wrote:
> Not just global from global. global from device too, right? As for global from HD, the CUDA guide forbids it
I'll update the comment.

As for HD->G, nvcc happily compiles following code:

```
__global__ void kernel() {}
__host__ __device__ void foo() {
#if !defined(__CUDA_ARCH__)
  kernel<<<0,0>>>();
#endif
}

```

Nvcc does produce an error for HD->G call during device compilation (the error actually complains about D->G or G->G calling). This patch matches nvcc behavior.

================
Comment at: lib/Sema/SemaChecking.cpp:529
@@ -528,3 +528,3 @@
   // of the arch we are compiling for.
-  if (BuiltinID >= Builtin::FirstTSBuiltin) {
+  if (Context.BuiltinInfo.isTSBuiltin(BuiltinID)) {
     switch (Context.getTargetInfo().getTriple().getArch()) {
----------------
eliben wrote:
> Is this part related to this patch?
It's part of D12122 which broke some of your team's tests and got rolled back.
It's a prerequisite for overloads to work (otherwise anything that uses a builtin would violate calling convention either during device or during host compilation) and it also needs to be hidden behind some option so it does not break your tests again.

I think I can commit it separately after the overload patch. Overloading will not work with builtins until then, but I don't think it's a big deal as there are no users yet.



http://reviews.llvm.org/D12453