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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits