为什么cuda内核可以访问主机内存？答案

【问题标题】：why cuda kernel can access host memory?为什么cuda内核可以访问主机内存？
【发布时间】：2021-11-02 16:28:55
【问题描述】：

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

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

分配页面锁定且可供设备访问的主机内存的 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);
}

【问题讨论】：

标签： c++ cuda

【解决方案1】：

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

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

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

【讨论】：