数组的结构与结构数组
Structure of Arrays vs Array of Structures
从我在这里读到的一些评论来看,对于像CUDA这样的并行实现,最好是Structure of Arrays
(SoA
(而不是Array of Structures
(AoS
(。如果这是真的,谁能解释为什么?
AoS与SoA以获得最佳性能通常取决于访问模式。然而,这不仅限于 CUDA - 类似的考虑适用于性能可能受到内存访问模式显着影响的任何架构,例如,您有缓存的地方或连续内存访问的性能更好(例如 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 组件的连续读取将是连续的,并且通常包含在同一缓存行中。对于 CUDA,这也意味着内存读/写合并。
但是,如果您要单独处理颜色平面,那么 SoA 可能是首选,例如,如果您想按某个比例因子缩放所有 R 值,则 SoA 意味着所有 R 分量都是连续的。
另一个考虑因素是填充/对齐。对于上面的 RGB 示例,AoS 布局中的每个元素都对齐为 3 个字节的倍数,这对于 CUDA、SIMD 等来说可能不方便 - 在某些情况下甚至可能需要在结构中填充以使对齐更方便(例如,添加一个虚拟uint8_t元素以确保 4 字节对齐(。然而,在SoA情况下,平面是字节对齐的,这对于某些算法/架构来说可能更方便。
对于大多数图像处理类型的应用程序,AoS 场景更为常见,但对于其他应用程序或特定的图像处理任务,情况可能并非总是如此。当没有明显的选择时,我会推荐 AoS 作为默认选择。
另请参阅此答案,了解有关 AoS 与 SoA 的更多一般性讨论。
我只想提供一个简单的示例,说明数组结构 (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 <thrustdevice_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 = %fn", 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 = %fn", timerGPU.GetCounter());
timerGPU.StartCounter();
AoSvsSoA_v3 << <iDivUp(N, BLOCKSIZE), BLOCKSIZE >> >(cell, N);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
printf("Timing AoSvsSoA_v3 = %fn", 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 %in", h_x1[k], k + 11);
break;
}
if (h_x2[k] != k + 12) {
printf("h_x2[%i] not equal to %in", h_x2[k], k + 12);
break;
}
if (h_cells[k].x1 != k + 11) {
printf("h_cells[%i].x1 not equal to %in", h_cells[k].x1, k + 11);
break;
}
if (h_cells[k].x2 != k + 12) {
printf("h_cells[%i].x2 not equal to %in", 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)
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};
并通过访问所有成员创建__m128
数据:
__m128 reg = _mm_set_ps(v.x, ....);
如果您的阵列是 16 字节对齐的数据加载/存储速度更快,并且某些操作可以直接在内存中执行。
- C++将文本文件中的数据读取到结构数组中
- JSON转换为nlohmann JSON-lib中的结构数组
- 需要帮助将结构数组传递给函数
- 在 64 位边界上对齐C++结构数组?
- 带有枚举方向/类型的气泡排序结构数组
- 在C++中访问结构数组中的数据
- 从指向结构数组的指针中提取元素
- 如何制作一个地图,其中的值是C++中的结构数组
- 如何在函数中显示结构数组
- 包含作为成员C++函数的结构数组
- 如何修复函数中的 fstream 文件输入以将正确的信息存储在结构数组中?
- std::transform 将一个结构数组复制到另一个结构数组
- 将非托管结构数组封送到托管数组
- 显示结构数组
- 结构数组中所有类似字段的总和
- 结构数组的构造函数错误,错误消息:没有构造函数实例与参数列表匹配
- 在 glsl opengl 中将嵌套结构数组作为统一传递
- 使用初始化列表填充C++中的多维结构数组时出现问题
- 如何将字符串字符与结构数组进行比较?
- C++结构数组