从我在这里阅读的一些评论中,出于某种原因,对于像CUDA这样的并行实现,最好有Structure of Arrays
(SoA
)超过Array of Structures
(AoS
)?如果那是真的,任何人都可以解释为什么? 在此先感谢! AOS相对的SOA以获得最佳性能阵列与CUDA中结构阵列的结构
回答
选择通常取决于访问模式。然而,这不仅限于CUDA-类似的考虑适用于性能可受存储器访问模式显着影响的任何架构,例如,您拥有缓存的地方,或者连续访问内存时性能更好的地方(例如CUDA中的合并内存访问)。
E.g.为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成分的将是连续的并且通常包含在相同的缓存行中。对于CUDA来说,这也意味着内存读取/写入合并。
但是,如果你要处理的颜色平面分别然后SoA的可能是首选,例如如果你想通过一些比例因子来缩放所有的R值,那么SoA意味着所有的R分量都是连续的。
一个进一步的考虑是填充和/对齐。对于上面的RGB示例,AoS布局中的每个元素都对齐到3个字节的倍数,这对CUDA,SIMD等可能不方便 - 在某些情况下,甚至可能需要在结构中填充以使对齐更方便(例如添加一个虚拟uint8_t元素以确保4个字节的对齐)。然而,在SoA情况下,这些平面是字节对齐的,对于某些算法/体系结构可以更方便。
对于大多数图像处理类型的应用程序,AoS方案更为常见,但对于其他应用程序或特定图像处理任务,情况可能并非总是如此。当没有明显的选择时,我会推荐AoS作为默认选择。
也this answer见AOS v的SOA更广泛的讨论。
感谢您的回复!嗯,所以它达到了我想象的实现访问模式。因此,如果我们有一个线程访问AoS中的10个连续元素并假设我们有100个总元素,则SoA(每个都有这10个元素)是比让每个线程从AoS访问10个连续元素更好的选择,看法? – BugShotGG
对于RGB情况,如果每个线程一起访问R,G,B组件,那么AoS可能更好,而如果每个线程只访问一个颜色平面,则SoA可能会更好。但是,您需要查看所有线程的总访问模式,因为外部内存访问通常是瓶颈。 –
SOA是SIMD处理effectly好。 对于几个原因,但基本上它更有效地加载连续4个浮点寄存器中。喜欢的东西:
float v [4] = {0};
__m128 reg = _mm_load_ps(v);
比使用:
struct vec { float x; float, y; ....} ;
vec v = {0, 0, 0, 0};
,并通过访问所有成员创建__m128
数据:
__m128 reg = _mm_set_ps(v.x, ....);
,如果你的阵列是16字节对齐的数据加载/存储速度更快,有些操作可以直接在内存中执行。
你可以通过也许可以和“将它变成一个数组”: __m128 reg = _mm_load_ps(&v.x); 这将假定是16字节对齐,在成员之间0字节 – pilkch
我只想提供一个简单的例子,显示一个结构数组(SoA)如何比结构数组(AoS)更好地执行。
在这个例子中,我考虑三种不同版本的同一代码:
- SoA的(V1)
- 直阵列(V2)
- AOS(V3)
特别是,版本2
考虑使用直阵列。版本2
和3
的时序对于该示例是相同的,并且结果比版本1
更好。我怀疑,一般来说,直接数组可能更可取,尽管以可读性为代价,因为例如可以通过const __restrict__
为这种情况启用来自统一缓存的加载。
#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)
很好的例子。 – ZeroCool
- 1. IDL - 结构阵列到结构阵列
- 2. 与结构阵列结构的劈
- 3. 阵列结构
- 4. 阵列结构
- 5. 阵列结构的
- 6. 2D阵列与结构(C++)
- 7. 结构阵列与指针
- 8. Swift 2结构与阵列
- 9. 列表结构的阵列的阵列
- 10. 在结构阵列中分配结构阵列
- 11. 阵列结构或阵列的结构的
- 12. 从结构阵列
- 13. Javascript阵列结构
- 14. Mongodb阵列结构
- 15. 在阵列结构
- 16. C中的结构阵列
- 17. 结构阵列排序阵列[数组] [结构]
- 18. 阵列中结构C#
- 19. 结构指针到结构阵列
- 20. C结构分配到结构阵列
- 21. __local结构阵列的OpenCL
- 22. 阵列的结构和sizeof
- 23. 结构的动态阵列
- 24. 阵列结构CLI的
- 25. 阵列结构转型
- 26. 递归阵列结构
- 27. PHP:更改阵列结构
- 28. 转换阵列结构
- 29. 比较2结构阵列
- 30. 全局结构阵列
我已经发布阵列的结构和结构的阵列之间的性能的比较作为回答这个讯息:[在CUDA中对结构数组进行排序](http://stackoverflow.com/questions/23541503/sorting-arrays-of-structures-in-cuda/23645954#23645954)。 – JackOLantern