Зачем знать о CUDA Warps?

Question

У меня GeForce GTX460 SE, так что это: 6 SM x 48 ядер CUDA = 288 ядер CUDA. Известно, что в одной Деформации содержится 32 потока, и что в одном блоке одновременно (одновременно) может быть выполнен только один Деформация. То есть в одном мультипроцессоре (СМ) одновременно может выполняться только один Блок, один Деформация и только 32 потока, даже если доступно 48 ядер?

И кроме того, в качестве примера для распределения конкретных потоков и блоков могут быть использованы threadIdx.x и blockIdx.x. Для их выделения используйте ядро ​​<<< Blocks, Threads >>> (). Но как распределить определенное количество варп и распределить их, и если это невозможно, зачем беспокоиться о варпах?

Answer 1

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

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

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

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

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

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

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

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

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

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

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

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

Answer 2

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

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

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

Именно так:)

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

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

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