hliao added a comment. In D62603#1521788 <https://reviews.llvm.org/D62603#1521788>, @tra wrote:
> >> 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. that should assume that variable is not declared with `static`. that's also the motivation of this patch. 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
