将主机内存复制到 cuda __device__ 变量

Question

我尝试使用谷歌找到解决问题的方法，但失败了。有很多 sn-ps 并不完全适合我的情况，尽管我认为这是一个非常标准的情况。

我必须将几个不同的数据数组传输到 cuda。它们都是具有动态大小的简单结构数组。因为我不想将所有内容都放入 cuda 内核调用中，所以我想，__device__ 变量应该正是我所需要的。

这就是我尝试将主机数据复制到__device__ 变量的方式：

// MaterialDescription.h
struct MaterialDescription {
    unsigned char type;
    unsigned char diffuseR, diffuseG, diffuseB;
    __device__ __forceinline__ float4 diffuseColour() const {  return make_float4((float) diffuseR / 255.f, (float) diffuseG / 255.f, (float) diffuseB / 255.f, 0); }
};

// kernel.h
__device__ MaterialDescription* g_materials;
__global__ void deferredRenderKernel() {
     something = g_materials[indexDependingOnData].diffuseColour();
}

//Cuda.cu
const std::vector<MaterialDescription>& materials = getData();

// version 1
cudaMemcpyToSymbol(g_materials, &materials.front(), sizeof(MaterialDescription) * materialCount);

// version 2
MaterialDescription* ptr;
cudaMalloc((void**)&ptr, sizeof(MaterialDescription) * materialCount);
cudaMemcpy(ptr, &materials.front(), sizeof(MaterialDescription) * materialCount, cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(g_materials, ptr, sizeof(MaterialDescription) * materialCount);

// version 3
cudaMalloc((void**)&g_materials, sizeof(MaterialDescription) * materialCount);
cudaMemcpyToSymbol(g_materials, &materials.front(), sizeof(MaterialDescription) * materialCount);

deferredRenderKernel<<<numBlocks, threadsPerBlock>>();

但是，唯一有效的版本包括内核参数

// kernel.h
__device__ MaterialDescription* g_materials;
__global__
void deferredRenderKernel(MaterialDescription* ptr) {
    g_materials = ptr;
    something = g_materials[indexDependingOnData].diffuseColour();
}

//Cuda.cu
// version 4, the only one working. but i pass again via kernel param
// in the worst case i'll stick to this, at least i wouldn't have to pass the
// parameters into device functions
MaterialDescription* ptr;
cudaMalloc((void**)&ptr, sizeof(MaterialDescription) * materialCount);
cudaMemcpy(ptr, &materials.front(), sizeof(MaterialDescription) * materialCount, cudaMemcpyHostToDevice);

deferredRenderKernel<<<numBlocks, threadsPerBlock>>(ptr);

编辑： 这个版本（由 Robert Crovella 提出）也可以工作，但内存不是动态分配的。

 // kernel.h
 __device__ MaterialDescription g_materials[VIENNA_MAX_MATERIAL_COUNT];
__global__
void deferredRenderKernel() {
    something = g_materials[indexDependingOnData].diffuseColour();
}

// cuda.h
// version 1
cudaMemcpyToSymbol(g_materials, &materials.front(), sizeof(MaterialDescription) * materialCount);

其他变量和结构同上。

编辑：

解决方案

它终于按我想要的方式工作了。

MaterialDescription.h

struct MaterialDescription {
    unsigned char type;
    unsigned char diffuseR, diffuseG, diffuseB;
    __device__ __forceinline__ float4 diffuseColour() const {  return make_float4((float) diffuseR / 255.f, (float) diffuseG / 255.f, (float) diffuseB / 255.f, 0); }
};

kernel.h

__device__ MaterialDescription* g_materials;
__global__ void deferredRenderKernel() {
    something = g_materials[indexDependingOnData].diffuseColour();
}

Cuda.cu

const std::vector<MaterialDescription>& materials = getData();
MaterialDescription* dynamicArea;

// allocate memory on the device for our data
cudaMalloc((void**)&dynamicArea, sizeof(MaterialDescription) * materialCount); 

// copy our data into the allocated memory
cudaMemcpy(dynamicArea, &materials.front(), sizeof(MaterialDescription) * materialCount, cudaMemcpyHostToDevice);

// copy the pointer to our data into the global __device__ variable.
cudaMemcpyToSymbol(g_materials, &dynamicArea, sizeof(MaterialDescription*));

Answer 1

如果你在问这样的问题时能给出一个完整的例子，那就太好了。看看你对MaterialDescription 和materials 的定义会很有用。查看what SO expects 以了解“为什么我的代码不起作用？”类型的问题

这只保存一个指针的存储空间：

__device__ MaterialDescription* g_materials;

您不能将整个结构/对象复制到指针上。

当您像这样分配设备变量时，它是静态分配，这意味着需要在编译时知道大小。因此，如果您在编译时知道大小（或最大大小），则可以执行以下操作：

__device__ MaterialDescription g_materials[MAX_SIZE];

// this assumes materialCount <= MAX_SIZE
cudaMemcpyToSymbol(g_materials, &(materials.front()), sizeof(MaterialDescription) * materialCount);