Может ли float3 использовать объединение памяти CUDA?

Вкратце: понятия float3 не существует на том уровне, где происходит слияние. Таким образом, вопрос о том, будет ли объединен float3, на самом деле не правильный вопрос. По крайней мере, это не тот вопрос, на который вообще можно ответить. Вопрос, на который должен ответить могу, будет звучать так: «Будут ли загрузки/хранения, сгенерированные этим конкретным ядром, использующим float3 именно таким образом, в конечном итоге объединяться?» К сожалению, даже на этот вопрос можно действительно ответить, только взглянув на машинный код и, самое главное, профилируя…

Все современные архитектуры CUDA поддерживают загрузку и сохранение 1-, 2-, 4-, 8- и 16-байтной глобальной памяти. Здесь важно понимать, что это не означает, что, например, гипотетическая загрузка/сохранение 12-байтов будет происходить через какой-то другой механизм. Это означает, что глобальная память может быть доступна через 1-, 2-, 4-, 8- или 16-байтовые операции загрузки и сохранения. Вот и все; период. Существуют нет способы доступа к глобальной памяти, отличные от этих 1-, 2-, 4-, 8- или 16-байтовых операций загрузки и сохранения. В частности, нет 12-байтных загрузок и хранилищ.

float3 — это абстракция, существующая на уровне языка CUDA C++. Аппаратное обеспечение не имеет ни малейшего представления о том, что такое float3. Когда дело доходит до глобальной памяти, аппаратное обеспечение понимает только то, что вы можете загружать или хранить 1, 2, 4, 8 или 16 байт одновременно. CUDA C++ float3состоит из три числа с плавающей запятой. float (в CUDA) имеет ширину 4 байта. Таким образом, доступ к элементу float3 обычно просто сопоставляется с 4-байтовой загрузкой/сохранением. Доступ ко всем элементам float3 обычно приводит к трем загрузкам/сохранениям по 4 байта. Например:

__global__ void test(float3* dest)
{
    dest[threadIdx.x] = { 1.0f, 2.0f, 3.0f };
}

Если вы посмотрите на сборку PTX сгенерирует компилятор для этого ядра, вы увидите, что присвоение { 1.0f, 2.0f, 3.0f } нашему float3 скомпилировано до трех 4-байтных хранилищ:

    mov.u32         %r2, 1077936128;
    st.global.u32   [%rd4+8], %r2;
    mov.u32         %r3, 1073741824;
    st.global.u32   [%rd4+4], %r3;
    mov.u32         %r4, 1065353216;
    st.global.u32   [%rd4], %r4;

Это обычные загрузки/хранилища, как и любые другие, в них нет ничего особенного. И эти отдельные загрузки/хранения подвержены потенциальному объединению, как и любая другая загрузка/хранение. В этом конкретном примере шаблон доступа к памяти будет выглядеть так:

1st store:  xx xx t1 xx xx t2 xx xx t3 xx xx t4 xx xx t5 xx xx t6 …
2nd store:  xx t1 xx xx t2 xx xx t3 xx xx t4 xx xx t5 xx xx t6 xx …
3rd store:  t1 xx xx t2 xx xx t3 xx xx t4 xx xx t5 xx xx t6 xx xx …

Где tя — это я-й поток вашего варпа, а xx обозначает пропущенный 4-байтовый адрес. Как видите, между сохранениями, выполняемыми нашими потоками, есть 8-байтовые промежутки. Однако все еще существует довольно много 4-байтовых хранилищ, которые попадают в одну и ту же 128-байтную строку кэша. Таким образом, шаблон доступа по-прежнему допускает объединение немного (в любой текущей архитектуре), просто он далек от идеала. Но некоторые лучше, чем ничего. См. документация по CUDA для более подробной информации об этом.

Обратите внимание, что все это на самом деле зависит исключительно от того, к каким шаблонам доступа к памяти приводит сгенерированный машинный код. Независимо от того, может ли быть объединен доступ к памяти, и если да, то в какой степени, это не имеет ничего общего с использованием конкретного типа данных на уровне C++. Чтобы проиллюстрировать это положение, рассмотрим следующий пример:

struct Stuff
{
    float3 p;
    int blub;
};

__global__ void test(Stuff* dest)
{
    dest[threadIdx.x].p = { 1.0f, 2.0f, 3.0f };
    dest[threadIdx.x].blub = 42;
}

Глядя на сборку PTX, мы видим, что компилятор транслировал этот код C++ в четыре отдельных хранилища по 4 байта. Пока без сюрпризов. Немного изменим этот код

struct alignas(16) Stuff
{
    float3 p;
    int blub;
};

__global__ void test(Stuff* dest)
{
    dest[threadIdx.x].p = { 1.0f, 2.0f, 3.0f };
    dest[threadIdx.x].blub = 42;
}

и заметьте, что вдруг компилятор превратил все это в один 16-байтовый магазин. Зная, что объект Stuff гарантированно всегда находится на границе 16 байт и что по правилам языка C++ отдельные модификации элементов структуры здесь не могут наблюдаться в каком-либо определенном порядке другим потоком, компилятор может объединить все эти назначения в одно 16-байтовое хранилище, что в конечном итоге приводит к такому шаблону доступа, как

t1 t1 t1 t1 t2 t2 t2 t2 t3 t3 t3 t3 t4 t4 t4 t4 …

Другой пример:

__global__ void test(float3* dest)
{
    auto i = threadIdx.x % 3;
    auto m = i == 0 ? &float3::x : i == 1 ? &float3::y : &float3::z;
    dest[threadIdx.x / 3].*m = i;
}

Здесь мы снова записываем в массив float3. Тем не менее, каждый поток будет выполнять хранит ровно одну запись для одного из членов float3, а последовательные потоки будут сохранять по последовательным 4-байтовым адресам, что приводит к идеально объединенному доступу к памяти:

t1 t2 t3 t4 t5 t6 t7 t8 t9 t10 t11 t12 t13 t14 t15 …

Опять же, тот факт, что наш код C++ в какой-то момент использовал float3, сам по себе совершенно не имеет значения. Важно то, что мы на самом деле делаем, какие загрузки/сохранения генерируются и как в результате выглядит шаблон доступа…