CUDA中arrays与结构arrays的结构
从我读过的一些注释中，出于某种原因，最好是在CUDA之类的并行实现Array of Structures （ AoS ）的Array of Structures （ SoA ）？ 如果这是真的，任何人都可以解释为什么？ 提前致谢！
为了获得最佳性能，AoS与SoA的select通常取决于访问模式。 然而，这不仅仅局限于CUDA – 类似的考虑适用于性能可能受到内存访问模式显着影响的任何体系结构，例如，具有高速caching的地方或者连续内存访问（例如，CUDA中的内存访问合并）的性能更好的体系结构。
例如，对于RGB像素与单独的RGB平面：
struct { uint8_t r, g, b; } AoS[N]; struct { uint8_t r[N]; uint8_t g[N]; uint8_t b[N]; } SoA;
如果要同时访问每个像素的R / G / B组件，那么AoS通常是有意义的，因为R，G，B组件的连续读取将是连续的并且通常包含在同一个caching行中。 对于CUDA来说，这也意味着内存读/写合并。
但是，如果要分别处理颜色平面，则可能首选SoA，例如，如果要按比例因子缩放所有R值，则SoA表示所有R分量都是连续的。
另一个考虑是填充/alignment。 对于上面的RGB示例，AoS布局中的每个元素都alignment到3个字节的倍数，这对于CUDA，SIMD等可能不方便 – 在某些情况下甚至可能需要在结构中填充以使alignment更方便（例如添加一个虚拟的uint8_t元素以确保4个字节的alignment）。 然而，在SoA情况下，这些平面是字节alignment的，对于某些algorithm/体系结构可以更方便。
对于大多数image processingtypes的应用程序，AoSscheme更为常见，但对于其他应用程序或特定的image processing任务，情况可能并非总是如此。 当没有明显的select时，我会推荐AoS作为默认select。
关于AoS v SoA的更多一般性讨论，请参见这个答案 。
SoA对于SIMD处理是非常有效的。 由于几个原因，但基本上更有效的是在一个寄存器中加载4个连续的浮点数。 像这样的东西：
float v [4] = {0}; __m128 reg = _mm_load_ps( v );
比使用：
struct vec { float x; float, y; ....} ; vec v = {0, 0, 0, 0};
并通过访问所有成员创build一个__m128数据：
__m128 reg = _mm_set_ps(vx, ....);
如果你的数组是16字节alignment的数据加载/存储速度更快，并且可以直接在内存中执行某些操作。
我只想提供一个简单的例子，显示一个结构数组（SoA）比一个结构数组（AoS）更好的performance。
在这个例子中，我正在考虑相同代码的三个不同版本：
- SoA（v1）
- 直阵（v2）
- AoS（v3）
特别是，版本2考虑使用直arrays。 版本2和版本3的时间在这个例子中是一样的，结果比版本1 。 我怀疑，一般来说，直接数组可能是可取的，尽pipe以可读性为代价，因为例如，可以通过const __restrict__为这种情况启用统一caching加载。
#include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> #include <thrust\device_vector.h> #include "Utilities.cuh" #include "TimingGPU.cuh" #define BLOCKSIZE 1024 /******************************************/ /* CELL STRUCT LEADING TO ARRAY OF STRUCT */ /******************************************/ struct cellAoS { unsigned int x1; unsigned int x2; unsigned int code; bool done; }; /*******************************************/ /* CELL STRUCT LEADING TO STRUCT OF ARRAYS */ /*******************************************/ struct cellSoA { unsigned int *x1; unsigned int *x2; unsigned int *code; bool *done; }; /*******************************************/ /* KERNEL MANIPULATING THE ARRAY OF STRUCT */ /*******************************************/ __global__ void AoSvsSoA_v1(cellAoS *d_cells, const int N) { const int tid = threadIdx.x + blockIdx.x * blockDim.x; if (tid < N) { cellAoS tempCell = d_cells[tid]; tempCell.x1 = tempCell.x1 + 10; tempCell.x2 = tempCell.x2 + 10; d_cells[tid] = tempCell; } } /******************************/ /* KERNEL MANIPULATING ARRAYS */ /******************************/ __global__ void AoSvsSoA_v2(unsigned int * __restrict__ d_x1, unsigned int * __restrict__ d_x2, const int N) { const int tid = threadIdx.x + blockIdx.x * blockDim.x; if (tid < N) { d_x1[tid] = d_x1[tid] + 10; d_x2[tid] = d_x2[tid] + 10; } } /********************************************/ /* KERNEL MANIPULATING THE STRUCT OF ARRAYS */ /********************************************/ __global__ void AoSvsSoA_v3(cellSoA cell, const int N) { const int tid = threadIdx.x + blockIdx.x * blockDim.x; if (tid < N) { cell.x1[tid] = cell.x1[tid] + 10; cell.x2[tid] = cell.x2[tid] + 10; } } /********/ /* MAIN */ /********/ int main() { const int N = 2048 * 2048 * 4; TimingGPU timerGPU; thrust::host_vector<cellAoS> h_cells(N); thrust::device_vector<cellAoS> d_cells(N); thrust::host_vector<unsigned int> h_x1(N); thrust::host_vector<unsigned int> h_x2(N); thrust::device_vector<unsigned int> d_x1(N); thrust::device_vector<unsigned int> d_x2(N); for (int k = 0; k < N; k++) { h_cells[k].x1 = k + 1; h_cells[k].x2 = k + 2; h_cells[k].code = k + 3; h_cells[k].done = true; h_x1[k] = k + 1; h_x2[k] = k + 2; } d_cells = h_cells; d_x1 = h_x1; d_x2 = h_x2; cellSoA cell; cell.x1 = thrust::raw_pointer_cast(d_x1.data()); cell.x2 = thrust::raw_pointer_cast(d_x2.data()); cell.code = NULL; cell.done = NULL; timerGPU.StartCounter(); AoSvsSoA_v1 << <iDivUp(N, BLOCKSIZE), BLOCKSIZE >> >(thrust::raw_pointer_cast(d_cells.data()), N); gpuErrchk(cudaPeekAtLastError()); gpuErrchk(cudaDeviceSynchronize()); printf("Timing AoSvsSoA_v1 = %f\n", timerGPU.GetCounter()); //timerGPU.StartCounter(); //AoSvsSoA_v2 << <iDivUp(N, BLOCKSIZE), BLOCKSIZE >> >(thrust::raw_pointer_cast(d_x1.data()), thrust::raw_pointer_cast(d_x2.data()), N); //gpuErrchk(cudaPeekAtLastError()); //gpuErrchk(cudaDeviceSynchronize()); //printf("Timing AoSvsSoA_v2 = %f\n", timerGPU.GetCounter()); timerGPU.StartCounter(); AoSvsSoA_v3 << <iDivUp(N, BLOCKSIZE), BLOCKSIZE >> >(cell, N); gpuErrchk(cudaPeekAtLastError()); gpuErrchk(cudaDeviceSynchronize()); printf("Timing AoSvsSoA_v3 = %f\n", timerGPU.GetCounter()); h_cells = d_cells; h_x1 = d_x1; h_x2 = d_x2; // --- Check results for (int k = 0; k < N; k++) { if (h_x1[k] != k + 11) { printf("h_x1[%i] not equal to %i\n", h_x1[k], k + 11); break; } if (h_x2[k] != k + 12) { printf("h_x2[%i] not equal to %i\n", h_x2[k], k + 12); break; } if (h_cells[k].x1 != k + 11) { printf("h_cells[%i].x1 not equal to %i\n", h_cells[k].x1, k + 11); break; } if (h_cells[k].x2 != k + 12) { printf("h_cells[%i].x2 not equal to %i\n", h_cells[k].x2, k + 12); break; } } }
以下是时间（在GTX960上执行的时间）：
Array of struct 9.1ms (v1 kernel) Struct of arrays 3.3ms (v3 kernel) Straight arrays 3.2ms (v2 kernel)
