CUDA автоматически преобразовывает массивы float4 в структуру массивов?

Question

У меня есть следующий фрагмент кода:

#include <stdio.h>

struct Nonsense {
    float3 group;
    float other;
};

__global__ void coalesced(float4* float4Array, Nonsense* nonsenseArray) {
    float4 someCoordinate = float4Array[threadIdx.x];
    someCoordinate.x = 5;
    float4Array[threadIdx.x] = someCoordinate;

    Nonsense nonsenseValue = nonsenseArray[threadIdx.x];
    nonsenseValue.other = 3;
    nonsenseArray[threadIdx.x] = nonsenseValue;
}

int main() {
    float4* float4Array;
    cudaMalloc(&float4Array, 32 * sizeof(float4));
    cudaMemset(float4Array, 32 * sizeof(float4), 0);

    Nonsense* nonsenseArray;
    cudaMalloc(&nonsenseArray, 32 * sizeof(Nonsense));
    cudaMemset(nonsenseArray, 32 * sizeof(Nonsense), 0);

    coalesced<<<1, 32>>>(float4Array, nonsenseArray);
    cudaDeviceSynchronize();
    return 0;
}

Когда я запускаю это через профилировщик Nvidia в Nsight и смотрю на шаблон доступа к глобальной памяти, у float4Array идеально сочетаются операции чтения и записи. Между тем, массив Nonsense имеет плохие шаблоны доступа (из-за того, что он является массивом структур).

Преобразует ли NVCC автоматически массив float4, который концептуально представляет собой массив структур, в структуру массива для улучшения структуры доступа к памяти?

Answer 1

Нет, он не преобразует его в структуру массивов. Я думаю, что если вы тщательно обдумаете это, вы придете к выводу, что компилятору практически невозможно реорганизовать данные таким образом. В конце концов, то, что передается, это указатель.

Существует только один массив, и элементы этого массива по-прежнему имеют элементы struct в том же порядке:

float address (i.e. index):      0      1      2      3      4      5 ...
array element             : a[0].x a[0].y a[0].z a[0].w a[1].x a[1].y ...

Однако массив float4 дает лучший шаблон, поскольку компилятор генерирует одиночную 16-байтовую загрузку на поток . Это иногда называют «векторной загрузкой», потому что мы загружаем вектор (float4 в данном случае) на поток. Следовательно, смежные потоки все еще читают смежные данные, и у вас есть идеальное поведение объединения. В приведенном выше примере поток 0 будет читать a[0].x, a[0].y, a[0].z и a[0].w, поток 1 будет читать a[1].x, a[1].y и т. Д. Все это будет происходить в одном запрос (т. е. инструкция SASS), но может быть разделен на несколько транзакций . Разделение запроса на несколько транзакций не приводит к потере эффективности (в данном случае).

В случае структуры Nonsense компилятор не распознает, что эта структура также может быть загружена аналогичным образом, поэтому под капотом он должен генерировать 3 или 4 загрузки на поток:

• одна 8-байтовая загрузка (или две 4-байтовые загрузки) для загрузки первых двух слов float3 group
• одна 4-байтовая загрузка для загрузки последнего слова float3 group
• одна 4-байтовая загрузка для загрузки float other

Если вы отобразите вышеупомянутые нагрузки на поток, возможно, используя приведенную выше диаграмму, вы увидите, что каждая загрузка включает шаг (неиспользуемые элементы между элементами, загруженными на поток), что приводит к снижению эффективности.

Используя тщательное приведение типов или определение объединения в вашей структуре, вы можете заставить компилятор загрузить вашу Nonsense структуру за одну загрузку.

Этот ответ также охватывает некоторые идеи, связанные с преобразованием AoS -> SoA и связанным с этим повышением эффективности.

Этот ответ содержит информацию о векторной нагрузке.