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

[Reid Kleckner]

lgtm

One other mechanism you can consider is a default-error warning. We mostly added this mechanism so that we could emit errors by default, but suppress them in system headers or other places. However, I think this mostly confuses end users who have to pass things like -Wno-host-calls-from-host-device in order to silence something that looks like an error, not a warning.


================
Comment at: test/CodeGenCUDA/host-device-calls-host.cu:21
@@ +20,3 @@
+extern "C"
+__host__ __device__ void hd_function() {
+  host_function();
----------------
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.

http://reviews.llvm.org/D7841

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