Блог

Насколько я понимаю, только доступ к памяти на 4 байта, 8 байт или 16 байт на поток может обеспечить объединение глобальной памяти CUDA. После этого часто используемый float3 является 612-байтовый тип и исключен для объединения. Я прав?

tl; dr: концепция a float3 не существует на уровне, где происходит объединение. Таким образом, вопрос о том, будет ли a объединен или нет 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 состоит из трех поплавков. A float (в CUDA) имеет ширину 4 байта. Таким образом, доступ к элементу a float3 будет, как правило, просто сопоставляться с 4-байтовой загрузкой / хранилищем. Доступ ко всем элементам a 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 …
  

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

 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;
}
  

 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;
}
  

 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 ? amp;float3::x : i == 1 ? amp;float3::y : amp;float3::z;
    dest[threadIdx.x / 3].*m = i;
}
  

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