Из некоторых комментариев, которые я прочитал здесь, по какой-то причине предпочтительнее иметь Structure of Arrays
(SoA
) над Array of Structures
(AoS
) для параллельных реализаций, таких как CUDA? Если это правда, может кто-нибудь объяснить, почему?
Спасибо заранее!
Структура массивов по сравнению с массивом структур в CUDA
Ответ 1
Выбор 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 и др. - в некоторых случаях, возможно, даже требуется заполнить внутри структуры, чтобы сделать выравнивание более удобным (например, добавьте элемент dummy uint8_t для обеспечения 4-байтового выравнивания). Однако в случае SoA плоскости выровнены по байт, что может быть более удобным для определенных алгоритмов/архитектур.
Для большинства приложений типа обработки изображений сценарий AoS гораздо более распространен, но для других приложений или для конкретных задач обработки изображений это не всегда так. Когда нет очевидного выбора, я бы рекомендовал AoS в качестве выбора по умолчанию.
См. также этот ответ для более общего обсуждения AoS v SoA.
Ответ 2
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-байтовой выгруженной загрузки/хранения данных быстрее, а некоторые операционные системы могут выполняться непосредственно в памяти.
Ответ 3
Я просто хочу привести простой пример, показывающий, как Struct of Arrays (SoA) работает лучше, чем Array of Structs (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)