[PATCH] CUDA: Add option to allow host device functions to call host functions

[Jacques Pienaar]

================
Comment at: include/clang/Basic/DiagnosticSemaKinds.td:6070
@@ -6069,1 +6069,3 @@
   "function %1 in %select{__device__|__global__|__host__|__host__ __device__}2 function">;
+def warn_host_device_function_calling_host_function : Warning<
+  "calling __host__ function %0 from __host__ __device__ function %1 can lead to runtime errors">,
----------------
eliben wrote:
> Update this name to the new one for consistency
Done.

================
Comment at: test/CodeGenCUDA/host-device-calls-host.cu:21
@@ +20,3 @@
+extern "C"
+__host__ __device__ void hd_function() {
+  host_function();
----------------
rnk wrote:
> jpienaar wrote:
> > rnk wrote:
> > > I think this is a more interesting test case:
> > >   extern "C" {
> > >   void host_function() {}
> > >   __host__ __device__ void hd_function(bool b) { if (b) host_function(); }
> > >   __device__ void device_function() { hd_function(false); }
> > >   }
> > > 
> > > It actually tests emission of the bogus call, even though it can never occur in practice. What should clang do for that?
> > Both nvcc and clang (with this patch) accepts this and the resulting code executes without errors. clang should still warn that this can cause a runtime failure as there is no call-site analysis performed.
> Exciting. =D I think throwing this example into the IRgen test suite is nice because it's a good representative edge case.
Done.

================
Comment at: test/SemaCUDA/function-target-hd.cu:3
@@ +2,3 @@
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
----------------
eliben wrote:
> Much better thanks. Just a small nit: add a comment that explains the various permutations of CUDA_ARCH and TEST_WARN_HD to make this test file more readable a year from now
Done.

http://reviews.llvm.org/D7841

EMAIL PREFERENCES
  http://reviews.llvm.org/settings/panel/emailpreferences/