[PATCH] CUDA host device code with two code paths

[Jacques Pienaar]

I think your memory is correct (at least thats what I thought too). And
yes, the macro would be defined externally when targeting the device. At
that point we can remove this check for macro definition as we'd then be
able to check the flags directly. So I do see this as a temporary solution
which disrupts as little as possible.

Is the concern here with how the test is written? (i.e., the test
explicitly sets this macro which will in future be set by the compiler
itself). In which case it could be changed to

__host__ __device__ void hd1(void) {
#ifdef __CUDA_ARCH__
  hd1d();
  hd1h(); // expected-error {{no matching function}}
#else
  hd1d(); // expected-error {{no matching function}}
  hd1h();
#endif
  hd1hd();
  hd1g<<<1, 1>>>(); // expected-error {{reference to __global__
function 'hd1g' in __host__ __device__ function}}
}


On Mon, Dec 1, 2014 at 10:00 AM, Reid Kleckner <rnk at google.com> wrote:

> Can you remind me what the CUDA compilation model is currently? My memory
> was that the clang driver was eventually going to launch two -cc1 actions,
> one for device and one for host, presumably with different flags. I would
> expect that lib/Frontend/InitPreprocessor.cpp would define this macro when
> targeting the device.
>
> If we're doing a single compilation with a fat object approach, we may
> need to do something weird to get this right. =/
>
> http://reviews.llvm.org/D6457
>
>
>


-- 
Jacques Pienaar | Platforms | jpienaar at google.com | 765-430-6883
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20141201/cd74b120/attachment.html>