Профилировщик CUDA сообщает о неэффективном доступе к глобальной памяти

У меня есть простое ядро ​​CUDA, которое, как я думал, эффективно обращается к глобальной памяти. Однако профилировщик Nvidia сообщает, что я выполняю неэффективный доступ к глобальной памяти. Мой код ядра:

__global__ void update_particles_kernel
(
    float4 *pos, 
    float4 *vel, 
    float4 *acc, 
    float dt, 
    int numParticles
)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
int offset = 0;

while(index + offset < numParticles)
{
    vel[index + offset].x += dt*acc[index + offset].x;   // line 247
    vel[index + offset].y += dt*acc[index + offset].y;
    vel[index + offset].z += dt*acc[index + offset].z;

    pos[index + offset].x += dt*vel[index + offset].x;   // line 251
    pos[index + offset].y += dt*vel[index + offset].y;
    pos[index + offset].z += dt*vel[index + offset].z;

    offset += blockDim.x * gridDim.x;
}

В частности, профайлер сообщает следующее:

введите здесь описание изображения

Из лучшего практическое руководство говорит:

"Для устройств с вычислительными возможностями 2.x требования можно сформулировать довольно просто: одновременный доступ потоков варпа будет объединяться в количество транзакций, равное количеству строк кэша, необходимых для обслуживания всех потоков варпа. По умолчанию все обращения кешируются через L1, который представляет собой 128-байтовые строки. Для разрозненных шаблонов доступа, чтобы уменьшить перевыборку, иногда может быть полезно кэшировать только в L2, который кеширует более короткие 32-байтовые сегменты ( см. Руководство по программированию CUDA C).

Для устройств с вычислительными возможностями 3.x доступ к глобальной памяти кэшируется только в L2; L1 зарезервирован для доступа к локальной памяти. Некоторые устройства с вычислительными возможностями 3.5, 3.7 или 5.2 также допускают кэширование глобальных переменных в L1."

Теперь, основываясь на этой информации, я ожидаю, что в моем ядре потребуется 16 доступов для обслуживания 32-поточного деформирования, потому что float4 составляет 16 байт, а на моей карте (770 м вычислительных возможностей 3.0) чтение из кеша L2 выполняется фрагментами по 32 байта ( 16 байт * 32 потока / 32 байта строк кэша = 16 обращений). Действительно, как вы можете видеть, профилировщик сообщает, что я выполняю 16 доступов. Чего я не понимаю, так это того, почему профилировщик сообщает, что идеальный доступ будет включать 8 транзакций L2 за доступ к строке 247 и только 4 транзакции L2 за доступ к остальным строкам. Может кто-нибудь объяснить, что мне здесь не хватает?


caching memory cuda profiler
person James    schedule 25.02.2017    source источник

Ответы (1)


arrow_upward
4
arrow_downward

У меня есть простое ядро ​​CUDA, которое, как я думал, эффективно обращается к глобальной памяти. Однако профилировщик Nvidia сообщает, что я выполняю неэффективный доступ к глобальной памяти.

Например, ваш массив float4 vel хранится в памяти следующим образом:

0.x 0.y 0.z 0.w 1.x 1.y 1.z 1.w 2.x 2.y 2.z 2.w 3.x 3.y 3.z 3.w ...
  ^               ^               ^               ^             ...
  thread0         thread1         thread2         thread3

Итак, когда вы делаете это:

vel[index + offset].x += ...;   // line 247

вы получаете доступ (храните) в местах (.x), которые я отметил выше. Пробелы между каждой отметкой ^ указывают на неэффективный шаблон доступа, на который указывает профилировщик. (Не имеет значения, что в самой следующей строке кода вы сохраняете в .y местах.)

Есть как минимум 2 решения, одним из которых будет классическая реорганизация ваших данных AoS -> SoA с соответствующими корректировками кода. Это хорошо задокументировано (например, здесь в теге cuda и в других местах) с точки зрения того, что это означает и как это сделать, поэтому я позволю вам посмотреть это.

Другим типичным решением является загрузка количества float4 на поток, когда вам это нужно, и сохранение количества float4 на поток, когда вам нужно. Ваш код может быть тривиально переработан для этого, что должно улучшить результаты профилирования:

//preceding code need not change
while(index + offset < numParticles)
{
    float4 my_vel = vel[index + offset];
    float4 my_acc = acc[index + offset];
    my_vel.x += dt*my_acc.x;   
    my_vel.y += dt*my_acc.y;
    my_vel.z += dt*my_acc.z;
    vel[index + offset] = my_vel;

    float4 my_pos = pos[index + offset];
    my_pos.x += dt*my_vel.x; 
    my_pos.y += dt*my_vel.y;
    my_pos.z += dt*my_vel.z;
    pos[index + offset] = my_pos;

    offset += blockDim.x * gridDim.x;
}

Даже если вы можете подумать, что этот код «менее эффективен», чем ваш код, потому что ваш код «похоже» загружает и сохраняет только .x, .y, .z, тогда как мой «кажется» также загружает и сохраняет .w, на самом деле по сути, нет никакой разницы из-за того, как графический процессор загружает и сохраняет в/из глобальной памяти. Хотя ваш код не затрагивает .w, в процессе доступа к соседним элементам GPU загрузит .w элементов из глобальной памяти, а также (в конечном итоге) сохранит .w элементов обратно в глобальную память.

Чего я не понимаю, так это почему профилировщик сообщает, что идеальный доступ будет включать 8 транзакций L2 на каждый доступ для строки 247.

Для строки 247 исходного кода вы получаете доступ к одному количеству float на поток для операции загрузки acc.x и одному количеству float на поток для операции загрузки vel.x. Количество float на поток само по себе должно требовать 128 байт для деформации, что соответствует 4 32-байтным кэш-линиям L2. Две загрузки вместе потребуют 8 загрузок кэшлайна L2. Это идеальный случай, предполагающий, что величины хорошо упакованы (SoA). Но это не то, что у вас есть (у вас есть AoS).

person Robert Crovella    schedule 25.02.2017
comment
Пробелы между каждым знаком ^ указывают на неэффективный шаблон доступа, на который указывает профилировщик. (Не имеет значения, что в самой следующей строке кода вы сохраняете в .y-местоположениях). Спасибо. Это была ключевая часть информации, которую мне не хватало. Я повторно запустил ваше решение, и теперь проблем нет! знак равно - person James; 25.02.2017