Может ли float3 наслаждаться объединением памяти CUDA? - PullRequest
0 голосов
/ 15 апреля 2019

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

1 Ответ

2 голосов
/ 15 апреля 2019

tl; dr: понятие 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 состоит из трех поплавков.A 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 i - i -й поток вашей деформации и 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, сам по себе совершенно не имеет значения. Важно то, что мы на самом деле делаем, какие нагрузки / хранилища генерируются и как выглядит шаблон доступа в результате & hellip;

Добро пожаловать на сайт PullRequest, где вы можете задавать вопросы и получать ответы от других членов сообщества.
...