将包含指针的结构复制到CUDA设备
我正在开发一个项目，我需要我的CUDA设备对包含指针的结构进行计算。
typedef struct StructA {
int* arr;
} StructA;
当我为结构分配内存然后将其复制到设备时，它只会复制结构而不是指针的内容。 现在我正在通过首先分配指针来解决这个问题，然后将主体结构设置为使用该新指针（驻留在GPU上）。 以下代码示例使用上面的结构描述了这种方法：
#define N 10
int main() {
int h_arr[N] = {1,2,3,4,5,6,7,8,9,10};
StructA *h_a = (StructA*)malloc(sizeof(StructA));
StructA *d_a;
int *d_arr;
// 1. Allocate device struct.
cudaMalloc((void**) &d_a, sizeof(StructA));
// 2. Allocate device pointer.
cudaMalloc((void**) &(d_arr), sizeof(int)*N);
// 3. Copy pointer content from host to device.
cudaMemcpy(d_arr, h_arr, sizeof(int)*N, cudaMemcpyHostToDevice);
// 4. Point to device pointer in host struct.
h_a->arr = d_arr;
// 5. Copy struct from host to device.
cudaMemcpy(d_a, h_a, sizeof(StructA), cudaMemcpyHostToDevice);
// 6. Call kernel.
kernel<<<N,1>>>(d_a);
// 7. Copy struct from device to host.
cudaMemcpy(h_a, d_a, sizeof(StructA), cudaMemcpyDeviceToHost);
// 8. Copy pointer from device to host.
cudaMemcpy(h_arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost);
// 9. Point to host pointer in host struct.
h_a->arr = h_arr;
}
我的问题是： 这是做到这一点的方法吗？
这似乎是一个非常多的工作，我提醒你，这是一个非常简单的结构。 如果我的结构包含很多指针或指针自身的结构，分配和复制的代码将会相当广泛和混乱。
编辑： CUDA 6引入了统一内存，这使得“深层复制”问题变得更加容易。 看到这个职位了解更多详情。
不要忘记，你可以通过价值传递结构到内核。 此代码有效：
// pass struct by value (may not be efficient for complex structures)
__global__ void kernel2(StructA in)
{
in.arr[threadIdx.x] *= 2;
}
这样做意味着您只需将阵列复制到设备，而不是结构：
int h_arr[N] = {1,2,3,4,5,6,7,8,9,10};
StructA h_a;
int *d_arr;
// 1. Allocate device array.
cudaMalloc((void**) &(d_arr), sizeof(int)*N);
// 2. Copy array contents from host to device.
cudaMemcpy(d_arr, h_arr, sizeof(int)*N, cudaMemcpyHostToDevice);
// 3. Point to device pointer in host struct.
h_a.arr = d_arr;
// 4. Call kernel with host struct as argument
kernel2<<<N,1>>>(h_a);
// 5. Copy pointer from device to host.
cudaMemcpy(h_arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost);
// 6. Point to host pointer in host struct
// (or do something else with it if this is not needed)
h_a.arr = h_arr;
正如Mark Harris指出的那样，结构可以通过值传递给CUDA内核。 但是，由于在内核退出时调用析构函数，因此应该注意设置适当的析构函数。
考虑下面的例子
#include <stdio.h>
#include "Utilities.cuh"
#define NUMBLOCKS 512
#define NUMTHREADS 512 * 2
/***************/
/* TEST STRUCT */
/***************/
struct Lock {
int *d_state;
// --- Constructor
Lock(void) {
int h_state = 0; // --- Host side lock state initializer
gpuErrchk(cudaMalloc((void **)&d_state, sizeof(int))); // --- Allocate device side lock state
gpuErrchk(cudaMemcpy(d_state, &h_state, sizeof(int), cudaMemcpyHostToDevice)); // --- Initialize device side lock state
}
// --- Destructor (wrong version)
//~Lock(void) {
// printf("Calling destructorn");
// gpuErrchk(cudaFree(d_state));
//}
// --- Destructor (correct version)
// __host__ __device__ ~Lock(void) {
//#if !defined(__CUDACC__)
// gpuErrchk(cudaFree(d_state));
//#else
//
//#endif
// }
// --- Lock function
__device__ void lock(void) { while (atomicCAS(d_state, 0, 1) != 0); }
// --- Unlock function
__device__ void unlock(void) { atomicExch(d_state, 0); }
};
/**********************************/
/* BLOCK COUNTER KERNEL WITH LOCK */
/**********************************/
__global__ void blockCounterLocked(Lock lock, int *nblocks) {
if (threadIdx.x == 0) {
lock.lock();
*nblocks = *nblocks + 1;
lock.unlock();
}
}
/********/
/* MAIN */
/********/
int main(){
int h_counting, *d_counting;
Lock lock;
gpuErrchk(cudaMalloc(&d_counting, sizeof(int)));
// --- Locked case
h_counting = 0;
gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice));
blockCounterLocked << <NUMBLOCKS, NUMTHREADS >> >(lock, d_counting);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost));
printf("Counting in the locked case: %in", h_counting);
gpuErrchk(cudaFree(d_counting));
}
与未注释的析构函数（不要太在意代码实际上做什么）。 如果您运行该代码，您将收到以下输出
Calling destructor
Counting in the locked case: 512
Calling destructor
GPUassert: invalid device pointer D:/Project/passStructToKernel/passClassToKernel/Utilities.cu 37
然后有两个对析构函数的调用，一次在内核出口，一次在主出口。 该错误消息与这样一个事实有关，如果d_state所指向的内存位置在内核出口处被释放，则它们不能在主出口处被释放。 因此，对于主机和设备执行，析构函数必须不同。 这是通过上面代码中注释的析构函数完成的。
数组的结构是cuda的噩梦。 您必须将每个指针复制到设备可以使用的新结构。 也许你可以使用一系列结构？ 如果不是，我发现的唯一方法就是以你的方式来攻击它，而这绝对不是。
编辑：因为我不能给顶部帖子发表评论：步骤9是多余的，因为你可以改变步骤8和9
// 8. Copy pointer from device to host.
cudaMemcpy(h->arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost);