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

Posted 2023-02-21

【中文标题】将主机内存复制到 cuda __device__ 变量

【英文标题】：copying host memory to cuda __device__ variable

【发布时间】：2014-12-07 03:29:24

【问题描述】：

我尝试使用谷歌找到解决问题的方法，但失败了。有很多 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*));

【问题讨论】：

如果您的结构仅由 POD 类型组成，那么您的版本 #2 几乎是正确的。只需将最后一个 memcpy 的大小更改为正确的大小（它只是您要复制的指针），它应该可以工作。

不仅尺寸，还需要参考 :) 我会在一秒钟内更新工作版本。

【参考方案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);

【讨论】：

对不起，我以为我把所有相关的东西都放在那里了。但你是对的，结构和材料也很重要。

所以没有办法让全局__device__内存区域具有动态大小？

使用cudaMalloc 进行动态分配。然后你可以cudaMemcpyToSymbolcudaMalloc返回的指针，或者将它作为内核参数传递。

感谢你们俩。我现在想通了，稍后会用解决方案更新问题。