为什么CUDA内核可以访问主机内存？

Question

我在cuda内核中直接访问主机内存，没有发现错误，为什么会这样？

我试图从文档中获取更多信息。

分配大小为size的主机内存，该内存是页面锁定的，并且可被设备访问。驱动程序跟踪使用此函数分配的虚拟内存范围，并自动加速对诸如cudaMemcpy*（）之类的函数的调用。由于设备可以直接访问内存，因此与使用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);
}

- coordinate

Answer 1

为什么许多CUDA程序在cudaMallocHost之后添加cudaMemcpy呢？

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

随着GPU和主机架构以及操作系统的发展，通过一些GPU硬件和驱动程序的魔法，在某些系统上GPU可以直接访问主机内存。不过这并不是普遍适用的。您发布的代码将不能在每个CUDA系统上正确运行，即使它在您的系统上正常工作（您的错误检查不足，可能根本没有工作）。

- talonmies