Host CUDA device global variable access

Question

I was wondering if there is an official source, why the following works:

#include <iostream>

struct Array{
    int el[10000];
};

__device__ Array devAr;

void test(Array& ar = devAr){
    for(int i=0; i<10000; i++)
        ar.el[i] = i;
    std::cout << ar.el[0] + ar.el[9999] << std::endl;
}

int main(){
    test();
}

You get the warning "the __device__ variable" devAr "cannot be directly read in the host function" if you try to access devAr directly, but there is no such warning by reference (for good reason). But in both cases, you can access the variable from the host. So there seems to be a host instance of this variable.

What I need to know: can I take this for granted?

Another test chart showing pointer values:

#include <iostream>
#include <cstdio>

__device__ int devAr[2];

__global__ void foo(){
    printf("Device: %p\n", &devAr);
    devAr[0] = 1337;
}

int main()
{
    devAr[0] = 4;
    std::cout << devAr[0] << std::endl;
    void* ad;
    cudaGetSymbolAddress(&ad, devAr);
    std::cout << ad << " " << &devAr << std::endl;
    foo<<<1,1>>>();
    cudaDeviceSynchronize();
    int arHost[2];
    cudaMemcpyFromSymbol(arHost, devAr, sizeof(arHost), 0);
    std::cout << "values: " << arHost[0] << std::endl;
}

Output:

4
0x500bc0000 0x66153c
Device: 0x500bc0000
values: 1337

Answer 1

, , , :

a __device__ devAr -

, :

#include <iostream>

__device__ int devAr[1];

int main()
{
    devAr[0] = 4;
    std::cout << devAr[0] << std::endl;
}

:

• __device__ int devAr[1]; devAr (, ).
• devAr , , . devAr .
• , devAr[0] = 4; 4 - .

, , , :

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

using namespace std;

__device__ int devAr[1];

__global__ void foo()
{
    printf("dev: %d \n", devAr[0]);
    devAr[0] = 5;
    printf("dev: %d \n", devAr[0]);
}

int main()
{
    cout << "host: " << devAr[0] << endl;
    devAr[0] = 4;
    cout << "host: " << devAr[0] << endl;

    foo << <1, 1 >> >();
    cudaDeviceSynchronize();
    cout << "host: " << devAr[0] << endl;
}

:

host: 0
host: 4
dev: 0
dev: 5
host: 4

UPDATE:

, , SO, , :

• cudaMemcpy() vs cudaMemcpyFromSymbol():

  (__device__, __constant__, ) , , , . CUDA . API - __constant__ __device__. API- ..

• CUDA:

  *PNT - __device__, , . (, .) , , (void**)&PNT, , . - , cudaGetSympolAddress() , cudaMemcpyToSymbol(), @talonmies.

• CUDA:

  , A B . , . . , cudaGetSymbolAddress .

• cudaMemcpyToSymbol cudaMemcpy, (cudaMemcpyToSymbol):

  API CUDA , GPU, API. , __device__ .

• cudaMemcpyFromSymbol __device__:

  , -:... , , , . , cudaGetSymbolAddress

, 3- :

• __device__ int devAr[1]; " " devAr (. 1 3).
• devAr - API-, cudaGetSymbolAddress ( , , ), devAr.

- " ", , CUDA, , . , , (.. devAr), , - devAr , API API.