Блоки, темы, размер

Было много дискуссий о том, как выбрать #blocks & blockSize, но я все еще что-то упустил. Многие из моих проблем касаются этого вопроса: как блоки / деформации CUDA отображаются на ядрах CUDA? (Для упрощения обсуждения достаточно памяти perThread и perBlock. Ограничения памяти здесь не проблема.)

kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal);

1) Чтобы СМ был максимально загружен, я должен установить nThreads к кратному warpSize, Правда?

2) SM может одновременно выполнять только одно ядро. То есть все HWcore этого SM выполняют только kernelA. (Не некоторые HWcore работают под управлением kernelA, в то время как другие запускают kernelB.) Так что, если у меня есть только один поток, я "трачу" другие HWcores. Правда?

3) Если варп-планировщик выдает работу в единицах warpSize (32 потока), и каждый SM имеет 32 HWcores, тогда SM будет полностью использован. Что происходит, когда SM имеет 48 HWcores? Как сохранить все 48 ядер полностью загруженными, когда планировщик запускает работу по 32 блока? (Если предыдущий абзац верен, разве не лучше, если планировщик выдаст работу в единицах размера HWcore?)

4) Похоже, что варп-планировщик ставит в очередь 2 задачи одновременно. Так что, когда текущее ядро ​​останавливается или блокируется, 2-е ядро ​​заменяется. (Это не ясно, но я предполагаю, что очередь здесь имеет глубину более 2 ядер.) Это правильно?

5) Если мой HW имеет верхний предел 512 потоков на блок (nThreadsMax), это не означает, что ядро ​​с 512 потоками будет работать быстрее всего на одном блоке. (Опять же, mem не проблема.) Есть большая вероятность, что я получу лучшую производительность, если распределю ядро ​​из 512 потоков по многим блокам, а не по одному. Блок выполняется на одном или нескольких SM. Правда?

5а) Я думаю, чем меньше, тем лучше, но имеет ли значение, насколько я маленький nBlocks? Вопрос в том, как выбрать значение nBlocks это прилично? (Не обязательно оптимально.) Есть ли математический подход к выбору nBlocksили это просто пробная версия.

Источник

2 ответа

Решение

Позвольте мне попытаться ответить на ваши вопросы один за другим.

  1. Это правильно.
  2. Что вы подразумеваете именно под "HWcores"? Первая часть вашего утверждения верна.

  3. Согласно Техническому документу NVIDIA Fermi Compute Architecture: "SM планирует потоки в группах из 32 параллельных потоков, называемых деформациями. Каждый SM имеет два планировщика деформации и два блока диспетчеризации команд, что позволяет одновременно создавать и выполнять два деформации. Планировщик двойного деформирования Fermi выбирает две деформации и выдает одну инструкцию из каждой деформации группе из шестнадцати ядер, шестнадцати модулей загрузки / хранения или четырех SFU. Поскольку деформации выполняются независимо, планировщику Fermi не нужно проверять зависимости внутри потока инструкций ".

    Кроме того, в официальном документе NVIDIA Keppler Architecture говорится: "Четырехканальный планировщик деформации Kepler выбирает четыре деформации, и в каждом цикле могут отправляться две независимые инструкции".

    Таким образом, "избыточные" ядра используются для планирования более чем одной деформации за раз.

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

  5. Не совсем верно: каждый блок привязан к одному SM, так как там находится его общая память.
  6. Это сложная проблема и зависит от того, как реализовано ваше ядро. Возможно, вы захотите взглянуть на веб-семинар nVidia " Лучшая производительность при низком уровне занятости " Василия Волкова, который объясняет некоторые из наиболее важных вопросов. Тем не менее, прежде всего, я бы предложил вам выбрать количество потоков для улучшения занятости, используя калькулятор занятости CUDA.
Источник

1) да.

2) Устройства CC 2.0 - 3.0 могут одновременно выполнять до 16 сетей. Каждый SM ограничен 8 блоками, поэтому для достижения полного параллелизма устройство должно иметь как минимум 2 SM.

3) Да, планировщики деформации выбирают и выдают деформации одновременно. Забудьте о концепции ядер CUDA, они не имеют значения. Чтобы скрыть задержку, вам нужно иметь высокий уровень параллелизма на уровне обучения или высокую занятость. Рекомендуется иметь>25% для CC 1.x и>50% для CC >= 2.0. Как правило, для CC 3.0 требуется более высокая загруженность, чем для устройств 2.0, из-за удвоения количества планировщиков, но только с увеличением деформаций на SM на 33%. Эксперимент Nsight VSE Issue Efficiency - лучший способ определить, достаточно ли у вас перекосов, чтобы скрыть задержки инструкций и памяти. К сожалению, Visual Profiler не имеет этой метрики.

4) алгоритм планировщика деформации не задокументирован; однако он не учитывает, из какой сетки возник блок потока. Для устройств CC 2.x и 3.0 распределитель работы CUDA распределяет все блоки из сетки перед распределением блоков из следующей сетки; однако это не гарантируется моделью программирования.

5) Для того, чтобы SM был занят, у вас должно быть достаточно блоков для заполнения устройства. После этого вы хотите убедиться, что у вас достаточно перекосов, чтобы достичь разумного уровня занятости. Существуют как плюсы, так и минусы использования больших блоков нитей. Большие блоки потоков обычно используют меньше кэша команд и занимают меньше места в кэше; тем не менее, большие потоки блокируются в syncthreads (SM может стать менее эффективным, так как на выбор остается меньше перекосов) и имеют тенденцию к тому, чтобы инструкции выполнялись на аналогичных исполнительных блоках. Я рекомендую попробовать 128 или 256 потоков на блок потоков для начала. Есть веские причины для больших и меньших блоков резьбы. 5a) Используйте калькулятор занятости. Выбор слишком большого размера блока потока часто приводит к ограничению регистров. Если вы выберете слишком малый размер блока потока, вы обнаружите, что вы ограничены общей памятью или пределом 8 блоков на SM.

Источник
Другие вопросы по тегам