Page MenuHomePhabricator

[CUDA] Make all host-side shadows of device-side variables undef.
ClosedPublic

Authored by tra on Dec 13 2018, 11:43 AM.

Details

Summary

The host-side code can't (and should not) access the values that may
only exist on the device side. E.g. address of a device function
does not exist on the host side as we don't generate the code for it there

This is similar to what nvcc does and allows us to compile code which has device-side variables with initializers that only exist on device side:

__device__ void func() {}
__device__ funcptr_t Funcs[] = {func};

Diff Detail

Repository
rL LLVM

Event Timeline

tra created this revision.Dec 13 2018, 11:43 AM
tra retitled this revision from [CUDA] Make all host-side shadows of device-side variables undefined. to [CUDA] Make all host-side shadows of device-side variables undef..Dec 13 2018, 11:44 AM
jlebar accepted this revision.Dec 13 2018, 12:59 PM
This revision is now accepted and ready to land.Dec 13 2018, 12:59 PM
yaxunl accepted this revision.Dec 13 2018, 1:39 PM

LGTM. Thanks!

This revision was automatically updated to reflect the committed changes.