Cuda:固定内存零拷贝问题

Question

我试过这个链接中的代码是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");
}

Answer 1

首先,要分配ZeroCopy内存,您必须将cudaHostAllocMappedflag 指定为参数cudaHostAlloc.

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

仍然是pinnedHostPointer将用于只从主机侧访问映射存储器.要从设备访问相同的内存,您必须获取设备端指针指向内存,如下所示:

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

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

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

此外,您必须将内核执行与主机同步以读取更新的值.只需cudaDeviceSynchronize在内核调用后添加.

链接问题中的代码正常工作,因为提出问题的人在64位操作系统上运行代码,并且启用了计算能力2.0和TCC的GPU.此配置自动启用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;
}