如何为 CUDA 中的数组结构分配内存？答案

【问题标题】：How to allocate memory to struct of arrays in CUDA?如何为 CUDA 中的数组结构分配内存？
【发布时间】：2020-09-03 04:27:52
【问题描述】：

假设数组大小为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;
};

【问题讨论】：

只需分配内存并将指向该内存的指针分配给每个成员。
@talonmies 你的意思是我应该为三个成员数组进行 3 次不同的内存分配吗？ 1.首先创建struct指针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 ));

标签： arrays pointers struct cuda malloc

【解决方案1】：

但我似乎想不出一种将内存分配给结构的方法指针abc *foo ......有没有其他方法可以为 SOA 分配内存（下）？

我不确定你的困难是什么。只要你没有数组结构的数组，为什么不简单地使用：

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

使用我上面的方法是个好习惯吗？

AoS 与 SoA 的问题取决于应用程序，对于 SO 上的 CUDA 应用程序（如 this answer），有许多关于此主题的优秀问题/答案。底线是当一个扭曲中的所有线程访问一个连续的内存块时，就会发生合并的内存访问。因此，如果可以合并对每个字段的访问，您可以在使用 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 是更好的选择。

【讨论】：