[PATCH] CUDA: Add option to allow host device functions to call host functions
Reid Kleckner
rnk at google.com
Tue Feb 24 13:04:53 PST 2015
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/
More information about the cfe-commits
mailing list