CUDA，如何在CUDA内核中实现动态结构体数组

Question

我正在尝试实现一个保存数据数组的结构我想实现动态数组，例如:

struct myStruct {
  float3 *data0, *data1;
};

__global__ void kernel(myStruct input) {
  unsigned int N = 2;
  while(someStatements) {
    data0 = new float3[N];
    // do somethings
    N *= 2;
  }
}

如何在 CUDA 内核中执行类似的操作？

Answer 1

如果您打算在具有最新版本 CUDA 的计算能力 2.x 或 3,x 设备上运行此代码，则您的内核代码几乎是正确的。 Fermi 和 Kepler 硬件上的 CUDA 4.x 和 5.0 支持 C++ new 运算符。请注意，使用 new 或 malloc 分配的内存是在设备上的运行时堆上分配的。它具有创建上下文的生命周期，但您目前无法从 CUDA 主机 API 直接访问它(因此通过 cudaMemcpy 或类似方式)。

我将您的结构和内核变成了一个简单的示例代码，您可以自己尝试一下，看看它是如何工作的:

#include <cstdio>

struct myStruct {
    float *data;
};

__device__ 
void fill(float * x, unsigned int n)
{
    for(int i=0; i<n; i++) x[i] = (float)i;
}

__global__ 
void kernel(myStruct *input, const unsigned int imax)
{
    for(unsigned int i=0,N=1; i<imax; i++, N*=2) {
        float * p = new float[N];
        fill(p, N);
        input[i].data = p;
    }
}

__global__
void kernel2(myStruct *input, float *output, const unsigned int imax)
{
    for(unsigned int i=0,N=1; i<imax; i++, N*=2) {
        output[i] = input[i].data[N-1];
    }
}

inline void gpuAssert(cudaError_t code, char * file, int line, bool Abort=true)
{
    if (code != 0) {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code),file,line);
        if (Abort) exit(code);
    }       
}
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

int main(void)
{

    const unsigned int nvals = 16;
    struct myStruct * _s;
    float * _f, * f;

    gpuErrchk( cudaMalloc((void **)&_s, sizeof(struct myStruct) * size_t(nvals)) );
    size_t sz = sizeof(float) * size_t(nvals);
    gpuErrchk( cudaMalloc((void **)&_f, sz) );
    f = new float[nvals];

    kernel<<<1,1>>>(_s, nvals);
    gpuErrchk( cudaPeekAtLastError() );

    kernel2<<<1,1>>>(_s, _f, nvals);
    gpuErrchk( cudaPeekAtLastError() );
    gpuErrchk( cudaMemcpy(f, _f, sz, cudaMemcpyDeviceToHost) );
    gpuErrchk( cudaDeviceReset() );

    for(int i=0; i<nvals; i++) {
        fprintf(stdout, "%d %f\n", i, f[i]);
    }

    return 0;
}

需要注意的几点:

1. 此代码只能在 Fermi 或 Kepler GPU 上使用 CUDA 4.x 或 5.0 进行编译和运行
2. 您必须将 GPU 的正确架构传递给 nvcc 来编译它(例如，我使用 nvcc -arch=sm_30 -Xptxas="-v"-o dynstruct dynstruct.cu 进行编译Linux 上的 GTX 670)
3. 示例代码使用“收集”内核将数据从运行时堆中的结构复制到主机 API 可以访问的分配，以便可以打印结果。这是针对我之前提到的关于 cudaMemcpy 无法直接从运行时堆内存中的地址进行复制的限制的解决方案。我希望这个问题可以在 CUDA 5.0 中得到解决，但最新的候选版本仍然有这个限制。