tra added a comment.
>> NVCC also allows that: https://godbolt.org/z/t78RvM
>
> BTW, that code posted looks quite weird to me, how the code could make sense
> by return a pointer of device variable? or a pointer of shadow host variable?
Magic. :-)
More practical example would be something like this:
__device__ int array[10];
__host__ func() {
cudaMemset(array, 0, sizeof(array));
}
cudaMemset is a host function and it needs to use something that exists on the
host side as the first argument.
In order to deal with this, compiler:
- creates uninitialized `int array[10]` on the host side. This allows ising
sizeof(array) on the host size.
- registers its address/size with CUDA runtime. This allows passing address of
host-side shadow array to various CUDA runtime routines. The runtime knows what
it has on device side and maps shadow's address to the real device address.
This way CUDA runtime functions can make static device-side data accessible
without having to explicitly figure out their device-side address.
Repository:
rC Clang
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D62603/new/
https://reviews.llvm.org/D62603
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
