c++ - 为什么cuda内核可以访问主机内存？

Question

我在cuda内核中直接访问host mem，没有发现错误，这是为什么呢？

我试图通过 documentation 变得更聪明.

Allocates size bytes of host memory that is page-locked and accessible to the device. The driver tracks the virtual memory ranges allocated with this function and automatically accelerates calls to functions such as cudaMemcpy*(). Since the memory can be accessed directly by the device, it can be read or written with much higher bandwidth than pageable memory obtained with functions such as malloc().

为什么很多cuda程序在cudaMallocHost后面加上cudaMemcpy？

#include <stdio.h>
#include <assert.h>

#define N 64

// cuda kernel access host mem a/b
__global__ void gpu(int *a, int *b, int *c_gpu) {
    int r = blockDim.x * blockIdx.x + threadIdx.x;
    int c = blockDim.y * blockIdx.y + threadIdx.y;

    if (r < N && c < N) {
        c_gpu[r * N + c] = a[r * N + c] + b[r * N + c];
    }
}

// cpu function
void cpu(int *a, int *b, int *c_cpu) {
    for (int r = 0; r < N; r++) {
        for (int c = 0; c < N; c++) {
            c_cpu[r * N + c] = a[r * N + c] + b[r * N + c];
        }
    }
}

int main() {
    int *a, *b, *c_cpu, *c_gpu, *c_gpu_cpu;
    size_t size = N * N * sizeof(int);

    cudaMallocHost(&a, size);
    cudaMallocHost(&b, size);
    cudaMallocHost(&c_cpu, size);
    cudaMallocHost(&c_gpu_cpu, size);
    cudaMalloc(&c_gpu, size);

    for (int r = 0; r < N; r++) {
        for (int c = 0; c < N; c++) {
            a[r * N + c] = r;
            b[r * N + c] = c;
            c_gpu_cpu[r * N + c] = 0;
            c_cpu[r * N + c] = 0;
        }
    }

    cpu(a, b, c_cpu);
    dim3 threads(16, 16, 1);
    dim3 blocks((N + threads.x - 1) / threads.x, (N + threads.y - 1) / threads.y, 1);

    gpu<<<blocks, threads>>>(a, b, c_gpu); // access cpu host mem
    cudaError_t err = cudaGetLastError();  
    if (err != cudaSuccess) {
        printf("Error: %s\n", cudaGetErrorString(err));
    }
    cudaDeviceSynchronize();

    cudaFreeHost(a);
    cudaFreeHost(b);
    cudaFreeHost(c_cpu);
    cudaFreeHost(c_gpu_cpu);
    cudaFree(c_gpu);
}

Answer 1

Why do many cuda programs add cudaMemcpy after cudaMallocHost?

因为很多CUDA程序是在统一内存系统出现之前编写的，而那个时候cudaMallocHost分配了页锁内存。该页面锁定内存仍然需要 API 调用才能进行复制。 “...直接由设备访问” 意味着 GPU 可以使用 DMA 通过 PCI Express 总线读取和写入内存，而主机内存管理器不需要做任何事情，这要快得多比传统的可分页主机内存。

随着 GPU 和主机架构以及操作系统的发展，GPU 可以通过某些 GPU 硬件和驱动程序魔法直接访问某些系统上的主机内存。但这不是普遍的。如果您发布的代码在您的系统上正常运行，那么您发布的代码将不会在每个 CUDA 系统上正常运行(您的错误检查有缺陷，它实际上可能根本无法运行)。