CUDA: в сокращении деформации и изменчивом ключевом слове
После прочтения вопроса и его ответа из следующего
ССЫЛКА НА САЙТ
У меня все еще есть вопрос, который остается в моей памяти. Из моего опыта в C / C ++; Я понимаю, что с помощью volatile имеет свои недостатки. В ответах также указывается, что в случае CUDA оптимизации могут заменить общий массив регистрами для хранения данных, если volatile Ключевое слово не используется.
Я хочу знать, с какими проблемами производительности можно столкнуться при расчете (сумме) сокращения. например
__device__ void sum(volatile int *s_data, int tid)
{
if (tid < 16)
{
s_data[tid] += s_data[tid + 16];
s_data[tid] += s_data[tid + 8];
s_data[tid] += s_data[tid + 4];
s_data[tid] += s_data[tid + 2];
s_data[tid] += s_data[tid + 1];
}
}
Я использую в уменьшении деформации. Так как все потоки с в Warp синхронизированы, поэтому я считаю, что нет необходимости использовать syncthreads() построить.
Я хочу знать, удалит ли ключевое слово volatile испортить мою сумму (из-за оптимизации cuda)? Могу ли я использовать такое сокращение без volatile ключевое слово.
Поскольку я использую эту функцию сокращения несколько раз, volatile Ключевое слово может привести к снижению производительности?
Решение
Удаление ключевого слова volatile из этого кода мог сломайте этот код на Fermi и Kepler GPUS. Эти графические процессоры не имеют инструкций для непосредственной работы с общей памятью. Вместо этого компилятор должен передать пару load / store в регистр и из него.
В этом контексте ключевое слово volatile заставляет компилятор соблюдать цикл загрузки-эксплуатации-хранения, а не выполнять оптимизацию, которая бы сохраняла значение s_data[tid] в реестре. Сохранение суммы, накапливающейся в регистре, нарушило бы неявную синхронизацию памяти, необходимую для правильного суммирования общей памяти на уровне деформации.
Другие решения
Других решений пока нет …