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

通过如下代码测试：

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

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

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

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

推测原因是：

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

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

然后查了以下资料：

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

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

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

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

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

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

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

__device__ void print(int *ptr) { for (int ii = 0; ii < 5; ii++) printf(" value=%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(" source 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(" error = %d", error); printf(" index Ptr = %p", gpuPtr[index]); }

for (int ii = 0; ii < 5; ii++) { gpuPtr[index][ii] = index * 10 + ii; //printf(" value=%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(); }