[PATCH] D62603: [CUDA][HIP] Skip setting `externally_initialized` for static device variables.

[Artem Belevich via Phabricator via cfe-commits]

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