Как GPU группирует потоки в деформации / волновые фронты?

Question

Насколько я понимаю, деформация - это группа потоков, которые определяются во время выполнения через планировщик задач. Одной из важнейших составляющих производительности CUDA является расхождение потоков внутри деформации, существует ли способ сделать правильное предположение о том, как оборудование будет создавать деформации в блоке потока?

Например, я запустил ядро ​​с 1024 потоками в блоке потоков, как устроены перекосы, могу ли я сказать (или, по крайней мере, сделать правильное предположение) из индекса потока?

Так как, делая это, можно минимизировать расхождение потоков в пределах данной деформации.

Answer 1

Расположение потоков внутри основы зависит от реализации, но у меня всегда было такое же поведение:

Основа состоит из 32 потоков, но диспетчер основы будет выдавать 1 инструкцию для остановки основы каждый раз (16 потоков)

• Если вы используете 1D-блоки (допустимо только измерение threadIdx.x), то диспетчер деформации выдаст 1 инструкцию для threadIdx.x = (0..15) (16..31)... и т. Д.

• Если вы используете 2D-блоки (измерения threadIdx.x и threadIdx.y действительны), то планировщик деформации будет пытаться выполнить следующую процедуру:

threadIdx.y = 0 threadIdx.x = (0..15) (16..31)... и т. д.

поэтому потоки с последовательным компонентом threadIdx.x будут выполнять одну и ту же инструкцию в группах по 16.

Answer 2

Деформация состоит из 32 потоков, которые будут выполняться одновременно. В любой момент времени на GPU будет выполняться пакет из 32, и это называется деформацией.

Я не нашел нигде, где говорится, что вы можете контролировать, что будет выполнять варп дальше, единственное, что вы знаете, это то, что он состоит из 32 потоков и что блок потоков всегда должен быть кратным этому числу.

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

Это также относится к операциям с памятью и задержке:

Когда потоки в деформации выдают операцию памяти устройства, эта инструкция займет очень много времени, возможно, сотни тактов, из-за большой задержки памяти. Основные архитектуры добавили бы иерархию кэш-памяти, чтобы уменьшить задержку, и Fermi включает некоторые аппаратные кеши, но в основном графические процессоры предназначены для потоковых или пропускных вычислений, где кэш-память неэффективна. Вместо этого эти графические процессоры допускают задержку памяти, используя высокую степень многопоточности. Tesla поддерживает до 32 активных деформаций на каждом мультипроцессоре, а Fermi поддерживает до 48. Когда одна деформация останавливается в операции с памятью, мультипроцессор выбирает другую готовую деформацию и переключается на нее. Таким образом, ядра могут быть производительными до тех пор, пока имеется достаточный параллелизм, чтобы держать их занятыми.

источник

Что касается деления потоковых блоков на основы, я обнаружил следующее:

если блок 2D или 3D, потоки упорядочены по первому измерению, затем второму, затем третьему, а затем разбиваются на деформации 32

источник