Cuda：固定内存零复制问题

我尝试了这个链接中的代码CUDA 固定内存是零拷贝吗？询问者声称该程序对他来说效果很好 但在我的身上效果不一样 如果我在内核中操作它们，这些值不会改变。

基本上我的问题是，我的 GPU 内存不够，但我想做需要更多内存的计算。我的程序使用 RAM 内存或主机内存，并且能够使用 CUDA 进行计算。链接中的程序似乎解决了我的问题，但代码没有给出该人所示的输出。

关于零复制内存的任何帮助或任何工作示例都会很有用。

谢谢

__global__ void testPinnedMemory(double * mem)
{
double currentValue = mem[threadIdx.x];
printf("Thread id: %d, memory content: %f\n", threadIdx.x, currentValue);
mem[threadIdx.x] = currentValue+10;
}

void test() 
{
const size_t THREADS = 8;
double * pinnedHostPtr;
cudaHostAlloc((void **)&pinnedHostPtr, THREADS, cudaHostAllocDefault);

//set memory values
for (size_t i = 0; i < THREADS; ++i)
    pinnedHostPtr[i] = i;

//call kernel
dim3 threadsPerBlock(THREADS);
dim3 numBlocks(1);
testPinnedMemory<<< numBlocks, threadsPerBlock>>>(pinnedHostPtr);

//read output
printf("Data after kernel execution: ");
for (int i = 0; i < THREADS; ++i)
    printf("%f ", pinnedHostPtr[i]);    
printf("\n");
}

首先要分配ZeroCopy内存，你必须指定cudaHostAllocMapped标志作为参数cudaHostAlloc.

cudaHostAlloc((void **)&pinnedHostPtr, THREADS * sizeof(double), cudaHostAllocMapped);

依然是pinnedHostPointer仅用于从主机端访问映射内存。要从设备访问相同的内存，您必须像这样获取指向内存的设备端指针：

double* dPtr;
cudaHostGetDevicePointer(&dPtr, pinnedHostPtr, 0);

将此指针作为内核参数传递。

testPinnedMemory<<< numBlocks, threadsPerBlock>>>(dPtr);

此外，您必须将内核执行与主机同步才能读取更新的值。只需添加cudaDeviceSynchronize在内核调用之后。

链接问题中的代码正在运行，因为提出问题的人正在 64 位操作系统上运行代码，并且 GPU 具有计算能力 2.0 并启用了 TCC。此配置自动启用统一虚拟寻址GPU 的功能，其中设备将主机 + 设备内存视为单个大内存，而不是单独的内存，并且使用分配的主机指针cudaHostAlloc可以直接传递给内核。

对于您的情况，最终代码将如下所示：

#include <cstdio>

__global__ void testPinnedMemory(double * mem)
{
    double currentValue = mem[threadIdx.x];
    printf("Thread id: %d, memory content: %f\n", threadIdx.x, currentValue);
    mem[threadIdx.x] = currentValue+10;
}

int main() 
{
    const size_t THREADS = 8;
    double * pinnedHostPtr;
    cudaHostAlloc((void **)&pinnedHostPtr, THREADS * sizeof(double), cudaHostAllocMapped);

    //set memory values
    for (size_t i = 0; i < THREADS; ++i)
        pinnedHostPtr[i] = i;

    double* dPtr;
    cudaHostGetDevicePointer(&dPtr, pinnedHostPtr, 0);

    //call kernel
    dim3 threadsPerBlock(THREADS);
    dim3 numBlocks(1);
    testPinnedMemory<<< numBlocks, threadsPerBlock>>>(dPtr);
    cudaDeviceSynchronize();

    //read output
    printf("Data after kernel execution: ");
    for (int i = 0; i < THREADS; ++i)
        printf("%f ", pinnedHostPtr[i]);    
    printf("\n");

    return 0;
}