Зачем беспокоиться о CUDA Warps?

Ответ 1

Ситуация довольно сложная, чем то, что вы описываете.

Блоки ALU (ядра), нагрузки/хранения (LD/ST) и специальные функциональные блоки (SFU) (зеленый на изображении) являются конвейерными устройствами. Они сохраняют результаты многих вычислений или операций в одно и то же время на разных этапах завершения. Таким образом, за один цикл они могут принять новую операцию и предоставить результаты другой операции, которая была запущена давным-давно (около 20 циклов для ALU, если я правильно помню). Таким образом, у одного SM теоретически есть ресурсы для обработки 48 * 20 циклов = 960 операций ALU в то же время, что составляет 960/32 потока на warp = 30 перекосов. Кроме того, он может обрабатывать операции LD/ST и операции SFU при любых задержках и пропускной способности.

Планировщики warp (желтый на изображении) могут планировать 2 * 32 потока на warp = 64 потока для трубопроводов за цикл. Таким образом, количество результатов, которые могут быть получены за такт. Таким образом, учитывая, что существует множество вычислительных ресурсов, 48 основных, 16 LD/ST, 8 SFU, каждый из которых имеет разные задержки, одновременно обрабатывается комбинация перекосов. В любом заданном цикле планировщики warp пытаются "соединить" два перекоса для планирования, чтобы максимизировать использование SM.

Планировщики warp могут создавать перекосы либо из разных блоков, либо из разных мест в одном блоке, если инструкции независимы. Таким образом, перекосы из нескольких блоков могут обрабатываться одновременно.

Добавляя к сложности, деформации, выполняющие инструкции, для которых имеется менее 32 ресурсов, должны выдаваться несколько раз для всех потоков, подлежащих обслуживанию. Например, имеется 8 SFU, поэтому это означает, что деформация, содержащая инструкцию, которая требует SFU, должна быть запланирована 4 раза.

Это описание упрощается. Существуют и другие ограничения, которые также вступают в игру, которые определяют, как графический процессор планирует работу. Вы можете найти дополнительную информацию, выполнив поиск в Интернете для "архитектуры ферми".

Итак, придя к вашему актуальному вопросу,

зачем беспокоиться о Warps?

Зная количество потоков в деформации и принимая это во внимание, становится важным, когда вы пытаетесь максимизировать производительность вашего алгоритма. Если вы не соблюдаете эти правила, вы теряете производительность:

• В вызове ядра <<<Blocks, Threads>>> попробуйте выбрать несколько потоков, которые равномерно делятся на количество потоков в основе. Если вы этого не сделаете, вы получите запуск блока, который содержит неактивные потоки.

• В вашем ядре попробуйте каждый поток в warp следовать одному и тому же пути кода. Если вы этого не сделаете, вы получите то, что называется декомпозицией варпа. Это происходит из-за того, что GPU должен запускать весь warp через каждый из расходящихся путей кода.

• В вашем ядре попытайтесь использовать каждый поток в загрузке основы и хранить данные в определенных шаблонах. Например, имеют потоки в warp, которые используют последовательные 32-битные слова в глобальной памяти.

Ответ 2

Связаны ли потоки, сгруппированные в Warps, 1 - 32, 33 - 64...?

Да, модель программирования гарантирует, что потоки сгруппированы в основы в определенном порядке.

В качестве простого примера оптимизации расходящихся путей кода можно использовать разделение всех потоков в блоке в группах из 32 потоков? Например: switch (threadIdx.s/32) {case 0:/* 1 warp */ break; case 1:/* 2 warp */break;/* Etc */}

Точно:)

Сколько байтов необходимо прочитать за один раз для одиночной Warp: 4 байта * 32 нитей, 8 байт * 32 нитей или 16 байт * 32 темы? Насколько мне известно, одна транзакция в глобальную память за один раз принимает 128 байтов.

Да, транзакции в глобальной памяти составляют 128 байт. Таким образом, если каждый поток читает 32-битное слово из последовательных адресов (они, вероятно, также должны быть выровнены по 128 байт), все потоки в warp могут обслуживаться одной транзакцией (4 байта * 32 потока = 128 байт). Если каждый поток читает больше байтов или если адреса не последовательны, необходимо выпустить больше транзакций (с отдельными транзакциями для каждой отдельной 128-байтовой строки, которая была затронута).

Это описано в Руководстве по программированию CUDA 4.2, раздел F.4.2, "Глобальная память". Там также есть рекламный ролик, в котором говорится, что ситуация отличается от данных, которые кэшируются только в L2, поскольку кэш L2 имеет 32-байтовые строки кэша. Я не знаю, как организовать кэширование данных только в L2 или сколько транзакций заканчивается.