内核访问主机内存

Question

我正在努力更好地掌握 cuda 中的内存管理。我刚刚发生的事情是严重缺乏理解。内核如何访问据我所知应该在主机内存中的值。

调用 vectorAdd() 时，它会在设备上运行该函数。但只有元素存储在设备内存中。向量的长度存储在主机上。内核如何不因尝试访问 foo.length 错误而退出，这应该在主机上。

#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>

typedef struct{
    float *elements;
    int length;
}vector;

__global__ void vectorAdd(vector foo, vector bar){
    int idx = threadIdx.x + blockDim.x * blockId.x.x;
    if(idx < foo.length){                      //this is the part that I do not understand
        foo.elements[idx] += bar.elements[idx];
    }
}

int main(void){
    vector foo, bar;
    foo.length = bar.length = 50;
    cudaMalloc(&(foo.elements), sizeof(float)*50);
    cudaMalloc(&(bar.elements), sizeof(float)*50);
    //these vectors are empty, so adding is just a 0.0 += 0.0
    int blocks_per_grid = 10;
    int threads_per_block = 5;
    vectorAdd<<<blocks_per_grid, threads_per_block>>>(foo, bar);
    return 0;
}

Answer 1

在 C 和 C++ 中，使参数可用于函数调用主体的典型机制是 pass-by-value。基本思想是制作一个单独的参数副本，供函数使用。

CUDA claims compliance to C++ (subject to various limitations), and it therefore provides a mechanism for pass-by-value. On a kernel call, the CUDA compiler and runtime will make copies of the arguments, for use by the function (kernel). In the case of a kernel call, these copies are placed in a particular area of __constant__ memory 位于 GPU 和 GPU 内存中 space，因此 "accessible" 到设备代码。

因此，在您的示例中，作为参数 vector foo, vector bar 的参数传递的 整个结构 被复制到 GPU 设备内存（特别是常量内存） CUDA 运行时。编译器以这种方式构造 CUDA 设备代码，以便根据需要直接从常量内存访问这些参数。

由于这些结构同时包含 elements 指针和标量 length，这两个 项在 CUDA 设备代码中都是可访问的，编译器将对它们的结构引用（例如 foo.length），以便从常量内存中检索所需的数量。

因此在您的示例中内核没有访问主机内存。按值传递机制使数量可用于 GPU 常量内存中的设备代码。

内核访问主机内存

kernels accessing host memory

kernel

cuda