如何为数组结构分配内存？

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));

我想看看AOS和SOA在CUDA下的性能差异。我还有没有其他方法可以为SOA分配内存（如下图）？用我上面的方法是一个好的做法吗？

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

Answer 1

但我似乎想不出给结构体指针分配内存的方法。abc *foo ...... 我还有其他方法可以为SOA分配内存吗（如下）？

我不知道你的困难在哪里。只要你没有一个数组结构的数组，为什么不简单的使用。

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

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

AoS与SoA的问题是取决于应用的，关于这个话题，在SO上的CUDA应用有很多优秀的问题回答（如 本回答). 底线是，当一个翘曲中的所有线程访问一个连续的内存块时，就会发生凝聚式内存访问。因此，如果对每个字段的访问可以被聚合，那么在使用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);

}

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

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

证实了SoA是一个更好的选择。