如何在CUDA中为结构体数组分配内存？

Question

假设数组大小为SOME_CONSTANT。因此我有这个AOS（结构体数组）

struct abc {
  float a;
  float b;
  float c;
};

我首先定义指针，然后分配一个AOS来为其分配内存。

abc *foo = 0;
foo = (abc *)malloc(SOME_CONSTANT * sizeof(abc));

那么现在很好。我想要创建一个数组结构 (SOA)，就像这样:

struct abc {
  float *a;
  float *b;
  float *c;
};

但是我似乎想不出一种为结构体指针abc *foo分配内存的方法。

我能想到的最好方法是这样的：

struct abc {
  float a[SOME_CONSTANT];
  float b[SOME_CONSTANT];
  float c[SOME_CONSTANT];
};

接着通过执行以下步骤：

abc *foo = 0;
foo = (abc *)malloc(sizeof(abc));

我想要比较在CUDA中AOS和SOA之间的性能差异。除了下方的方法，还有其他方式可以为SOA分配内存吗？使用上述方法是否是一种良好的实践？

struct abc {
  float *a;
  float *b;
  float *c;
};

- paulplusx

---

只需分配内存并将指针分配给每个成员即可。 - talonmies

@talonmies 您的意思是我应该为三个成员数组进行三次不同的内存分配吗？

1. 首先创建结构体指针 abc *foo = 0;
2. 然后为每个成员分配内存并赋值 foo->a= (float *)malloc(SOME_CONSTANT * sizeof(float )); foo->b= (float *)malloc(SOME_CONSTANT * sizeof(float )); foo->c= (float *)malloc(SOME_CONSTANT * sizeof(float ));

- paulplusx

Answer 1

但是我似乎想不出为结构体指针abc *foo分配内存的方法......是否有其他方法可以为下面的SOA分配内存？

我不确定你遇到了什么困难。只要你没有一个由数组组成的结构体数组，为什么不直接使用：

abc *foo;
cudaMalloc((void **)&foo, SOME_CONSTANT*sizeof(abc));

使用我上面提到的方法是一个好的做法吗？

AoS与SoA的选择取决于应用场景，关于这个主题在CUDA应用程序上有很多优秀的问题/答案（例如像这个答案）。总之，当一个warp中的所有线程访问一段连续的内存块时，就会出现协同访存。因此，如果每个字段的访问可以协同，那么当使用SoA时可以期望看到更高的内存带宽。针对你给出的示例，让我们运行一个简单的测试来量化性能差异：

#include <stdio.h>
#include <stdlib.h>

#define CHECK_CUDA(call)                                            \
{                                                                   \
const cudaError_t error = call;                                     \
if (error != cudaSuccess)                                           \
{                                                                   \
printf("ERROR:: File: %s, Line: %d, ", __FILE__, __LINE__);         \
printf("code: %d, reason: %s\n", error, cudaGetErrorString(error)); \
exit(EXIT_FAILURE);                                                 \
}                                                                   \
}

const int SOME_CONSTANT = 1024 * 1000; // to be executed on 1024 threads per block on 1000 blocks

// To be used as a SoA 
struct soa_abc {
    float *a;
    float *b;
    float *c;
};

// To be used as an AoS    
struct aos_abc {
    float a;
    float b;
    float c;
};

__global__ void kernel_soa(soa_abc foo) {
    unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
    foo.a[tid] = 1.f;
    foo.b[tid] = 2.f;
    foo.c[tid] = 3.f;
}

__global__ void kernel_aos(aos_abc *bar) {
    unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
    bar[tid].a = 1.f;
    bar[tid].b = 2.f;
    bar[tid].c = 3.f;
}

int main()
{
    float milliseconds = 0;
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    // SoA
    soa_abc foo;
    CHECK_CUDA(cudaMalloc((void **)&foo.a, SOME_CONSTANT * sizeof(float)));
    CHECK_CUDA(cudaMalloc((void **)&foo.b, SOME_CONSTANT * sizeof(float)));
    CHECK_CUDA(cudaMalloc((void **)&foo.c, SOME_CONSTANT * sizeof(float)));

    cudaEventRecord(start);
    kernel_soa <<<SOME_CONSTANT/1000, 1000 >>> (foo);
    CHECK_CUDA(cudaDeviceSynchronize());
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);

    printf("Time for SoA is %f ms.\n", milliseconds);

    CHECK_CUDA(cudaFree(foo.a));
    CHECK_CUDA(cudaFree(foo.b));
    CHECK_CUDA(cudaFree(foo.c));

    // AoS
    aos_abc *bar;
    CHECK_CUDA(cudaMalloc((void **)&bar, SOME_CONSTANT*sizeof(aos_abc)));

    cudaEventRecord(start);
    kernel_aos <<<SOME_CONSTANT/1000, 1000 >>> (bar);
    CHECK_CUDA(cudaDeviceSynchronize());
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);

    printf("Time for AoS is %f ms.\n", milliseconds);

}

在Windows和CUDA 10上使用Quadro P400进行测试，结果如下：

Time for SoA is 0.492384 ms.
Time for AoS is 1.217568 ms.

这证实了SoA是更好的选择。

- If_You_Say_So