CUDA 在 __global__ 与 __device__中分配内存问题

通过如下代码测试：

发现在__global__或__device__ 环境下可正常调用cudaMalloc分配gpu显存，并且可正常访问

__global__函数执行完成后，在新的global函数中，仍然可以访问在gpu环境下分配的显存

但是在执行cudaMemcpy函数对在__global__或__device__ 环境下分配的显存进行拷贝时，提示参数错误。（执行cudaMemcpy 拷贝到host环境下显存地址的指针值都是正确的。

与在__global__或__device__函数内部打印的值相同）

推测原因是：

在__global__或__device__函数中分配的cudaMalloc显存地址得到的指针，是直接指向显存的。

在__host__  函数中分配的cudaMalloc显存地址得到的指针，是将gpu地址进行了一次映射，映射到内存地址上，指向cudaMemcpy从显存拷贝到内存时，通过访问内存地址，得到地址中存储的显存地址，然后再把显存内的数据拷贝到内存中。

然后查了以下资料：

CUDA动态分配内存的问题-CSDN社区

看有人解答说核函数内无法动态分配内存，基本确认问题，暂时就不花时间分析了。

但准确点说，应该是核函数内分配的内存，在cpu中无法访问。

另外在核函数分配的显存大小也与在host中分配显存有明显差异：

在host中，循环调用cudaMalloc，可以分配达到甚至超过显存容量的显存空间（超出部分分配至内存：类似用硬盘虚拟内存）

而在__global__中，大约可分配4~6 MB空间（连续分配则可分配空间大一些）

所以怀疑在内核环境下cudaMalloc分配的显存与host环境下分配的显存不在相同区域。

下面的链接的问题与我的问题类似

请问CUDA核函数中可以申请多大的内存啊-CSDN社区

__device__ void print(int *ptr)
{
    for (int ii = 0; ii < 5; ii++)
        printf("\nvalue=%d", ptr[ii]);
}

__global__ void DynamicMalloc(int **gpuPtr, int *gpuPtr0)
{
    int index = blockDim.x * blockIdx.x + threadIdx.x;
    cudaError_t error;

    if (index >= 10)
        return;
    if (index == 0)
    {
        gpuPtr[index] = gpuPtr0;
    }
    else
    {
        printf("\nsource Ptr = %p", gpuPtr[index]);
        error = cudaMalloc(&gpuPtr[index], sizeof(int) * 5);
        //deviceMalloc((void **)&gpuPtr[index], sizeof(int) * 5);
        //gpuPtr[index] = (int *)malloc(sizeof(int ) * 5);
        printf("\n error = %d", error);
        printf("\nindex Ptr = %p", gpuPtr[index]);
    }

    for (int ii = 0; ii < 5; ii++)
    {
        gpuPtr[index][ii] = index * 10 + ii;
        //printf("\nvalue=%d", gpuPtr[index][ii]);
    }
    //print(gpuPtr[index]);
}

__global__ void DynamicPrint(int *gpuPtr)
{
    int index = blockDim.x * blockIdx.x + threadIdx.x;

    if (index >= 1)
        return;

    print(gpuPtr);
}

void gpuMallocTest()
{
    dim3 grid(1),block(512);
    int **gpuPtr = 0, *gpuPtr0 = 0, **cpuPtr = 0, cpuPtrs[10][10] = { 0 };
    cudaError_t error ;

    cudaMalloc(&gpuPtr, sizeof(int *) * 10);
    cudaMalloc(&gpuPtr0, sizeof(int) * 5);

    cpuPtr = (int **)malloc(sizeof(int *) * 10);
    DynamicMalloc << <grid, block >> > (gpuPtr, gpuPtr0);
    cudaDeviceSynchronize();
    cudaMemcpy(cpuPtr, gpuPtr, sizeof(int *) * 10, cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();

    for (int ii = 0; ii < 10; ii++)
    {
        error = cudaMemcpy(cpuPtrs[ii], cpuPtr[ii], sizeof(int ) * 5, cudaMemcpyDeviceToHost);

        DynamicPrint << <grid, block >> > (cpuPtr[ii]);
    }
    cudaMemcpy(cpuPtrs[0], gpuPtr0, sizeof(int) * 5, cudaMemcpyDeviceToHost);

    cudaDeviceSynchronize();
}