复制一个包含指针的结构体到CUDA设备

Question

我正在开发一个项目，需要让我的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;
}

我的问题是：这种方式正确吗？

这看起来需要大量的工作，而且我提醒你，这只是一个非常简单的结构体。如果我的结构体包含许多指针或具有指针本身的结构体，则分配和复制的代码将会相当冗长和混乱。

- Thorkil Holm-Jacobsen

---

步骤7和9是多余的，但除此之外基本上就是这样。正如下面的答案所说，最好避免在GPU上使用复杂的指针数据结构。GPU的性能会变差，并且API并不是为此设计的。 - talonmies

但是为了正确地将结构体复制到设备上，我必须将h_a的指针设置为d_arr(步骤4)。因此，当我将数据复制回来时，我还必须将h_a中的指针设置为刚刚复制到的数组。我同意在上面的示例中步骤7是多余的，因为结构体中没有其他信息，但如果有的话，那么这一步就不会是多余的。或者我完全错了吗？ - Thorkil Holm-Jacobsen

这只是一个完全虚构的例子，所以它是一个大多数无关紧要的观点。但是想象一下，如果您想要在循环中运行内核（比如说它是迭代方案的一部分，并且您需要将数据返回到主机以检查收敛性），那么步骤7既多余又错误。理想情况下，您应该有三个结构的副本——一个带有主机数据的主机结构、设备结构的主机副本和设备结构。在您的代码中，h_a 应该是/是其中第二个... - talonmies

这个能动态地做到吗？如果您不知道主机上数组的大小会是多少呢？ - rank1

谢谢tahatmat提供了这种在主机和设备内存之间复制结构的模式。然而，我认为值得一提的是第二种方式，它似乎更加一致，并有助于避免实现第9步。函数cudaMemcpy()的特定之处实际上允许在主机代码中解引用设备指针，如下所示：您可以跳过第4步，在将h_a复制到d_a的第5步之后，手动将每个设备指针地址复制到d_a中，例如：cudaMemcpy(&(d_a->arr), &(d_arr), sizeof(int*), cudaMemcpyHostToDevice)。同样，“d_a->arr”是合法的。 - vitrums

Answer 1

编辑：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;

- harrism

Answer 2

如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 destructor\n");
    //  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: %i\n", 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指向的内存位置，则无法在主退出时再次释放。因此，主机和设备执行的析构函数必须不同。这可以通过上面代码中注释的析构函数来实现。

- Vitality

Answer 3

结构体数组在CUDA中是一场噩梦。您将不得不将每个指针复制到新的结构体中，以便设备可以使用。也许您可以使用结构体数组？如果不行，我发现唯一的方法是像您所做的那样攻击它，这种方式并不美观。

编辑： 由于我无法在顶部帖子上发表评论：步骤9是多余的，因为您可以将步骤8和9更改为

// 8. Copy pointer from device to host.
cudaMemcpy(h->arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost);

- martiert

---

首先，这个回答是危险的，因为它违反了有关并行计算中AOS/SOA标准智慧。数组结构（SOA）比结构数组（AOS）在所有并行计算中都更可取，包括具有SSE/AVX指令集的多核CPU。原因是SOA在线程间保持引用局部性（例如，同时运行的相邻线程访问d_a.arr的相邻元素）。带有指针的结构体与数组结构不同。其次，您可以通过按值传递结构来简化此代码。 - harrism

@harrism 为什么在CUDA中不推荐使用结构体数组？我不理解这一点，你能给我一个例子或链接吗？谢谢。 - BugShotGG

@GeoPapas 这里 是一个讨论 SOA vs. AOS 并附有示例的问题/答案。 - Robert Crovella

@RobertCrovella 谢谢您的回复，但我已经在这里提出了一个问题[Here]（https://dev59.com/K2Mm5IYBdhLWcg3wDbin），答案非常清楚。 :) - BugShotGG