Понимание размеров сетки CUDA, размеров блоков и организации потоков (простое объяснение) [закрыто]

161

Как организованы потоки для выполнения графическим процессором?

cibercitizen1
источник
Руководство по программированию CUDA должно быть хорошим началом для этого. Я также рекомендовал бы проверить введение CUDA отсюда .
Том

Ответы:

287

аппаратные средства

Если устройство с графическим процессором имеет, например, 4 многопроцессорных модуля, и они могут запускать 768 потоков каждый: тогда в данный момент в действительности параллельно будет работать не более 4 * 768 потоков (если вы запланировали больше потоков, они будут ожидать их очередь).

Программное обеспечение

темы организованы в блоки. Блок выполняется многопроцессорным устройством. Потоки блока могут быть идентифицированы (проиндексированы) с использованием индексов 1Dimension (x), 2Dimensions (x, y) или 3Dim (x, y, z), но в любом случае x y z <= 768 для нашего примера (применяются другие ограничения) к x, y, z, см. руководство и возможности вашего устройства).

Очевидно, что если вам нужно больше этих 4 * 768 потоков, вам нужно больше 4 блоков. Блоки также могут быть проиндексированы 1D, 2D или 3D. Существует очередь блоков, ожидающих входа в GPU (поскольку в нашем примере GPU имеет 4 мультипроцессора и одновременно выполняется только 4 блока).

Теперь простой случай: обработка изображения 512x512

Предположим, мы хотим, чтобы один поток обрабатывал один пиксель (i, j).

Мы можем использовать блоки по 64 потока каждый. Тогда нам нужно 512 * 512/64 = 4096 блоков (таким образом, чтобы иметь 512x512 потоков = 4096 * 64)

Распространено организовать (чтобы упростить индексацию изображения) потоки в 2D-блоках, имеющих blockDim = 8 x 8 (64 потока на блок). Я предпочитаю называть это потоками PerBlock.

dim3 threadsPerBlock(8, 8);  // 64 threads

и 2D gridDim = 64 x 64 блоков (требуется 4096 блоков). Я предпочитаю называть это numBlocks.

dim3 numBlocks(imageWidth/threadsPerBlock.x,  /* for instance 512/8 = 64*/
              imageHeight/threadsPerBlock.y); 

Ядро запускается так:

myKernel <<<numBlocks,threadsPerBlock>>>( /* params for the kernel function */ );       

Наконец, будет что-то вроде «очереди из 4096 блоков», где блок ожидает назначения одного из мультипроцессоров графического процессора для выполнения 64 потоков.

В ядре пиксель (i, j), обрабатываемый потоком, рассчитывается следующим образом:

uint i = (blockIdx.x * blockDim.x) + threadIdx.x;
uint j = (blockIdx.y * blockDim.y) + threadIdx.y;
cibercitizen1
источник
11
Если каждый блок может выполнять 768 потоков, зачем использовать только 64? Если вы используете максимальный предел 768, у вас будет меньше блоков и, следовательно, лучшая производительность.
Ализа
10
@Aliza: блоки логичны , ограничение в 768 потоков для каждого физического процессора. Вы используете блоки в соответствии со спецификациями вашей задачи, чтобы распределить работу между потоками. Маловероятно, что вы всегда можете использовать блоки из 768 потоков для каждой вашей проблемы. Представьте, что вам нужно обработать изображение 64х64 (4096 пикселей). 4096/768 = 5,333333 блоков?
cibercitizen1
1
Блок логичен, но каждый блок назначен ядру. если блоков больше, чем ядро, блоки ставятся в очередь, пока ядра не станут свободными. В вашем примере вы можете использовать 6 блоков, и лишние потоки ничего не делают (2/3 потоков в 6-м блоке).
Ализа
3
@ cibercitizen1 - Я думаю, что точка зрения Ализы хороша: если возможно, нужно использовать как можно больше потоков на блок. Если есть ограничение, которое требует меньше потоков, лучше объяснить, почему это может иметь место во втором примере (но все же сначала объясните более простой и более желательный случай).
6
@ thouis Да, возможно. Но дело в том, что объем памяти, необходимый каждому потоку, зависит от приложения. Например, в моей последней программе каждый поток вызывает функцию оптимизации наименьших квадратов, требующую «много» памяти. Настолько, что блоки не могут быть больше, чем потоки 4х4. Несмотря на это, полученное ускорение было впечатляющим по сравнению с последовательной версией.
cibercitizen1
9

Предположим, 9800GT GPU:

  • имеет 14 мультипроцессоров (SM)
  • каждый SM имеет 8 потоковых процессоров (потоковые процессоры AKA, SP или ядра)
  • позволяет до 512 потоков на блок
  • Значение warpsize равно 32 (что означает, что каждый из потоковых процессоров 14x8 = 112 может планировать до 32 потоков)

https://www.tutorialspoint.com/cuda/cuda_threads.htm

Блок не может иметь больше активных потоков, чем 512, поэтому __syncthreadsможет синхронизировать только ограниченное количество потоков. т.е. если вы выполняете следующее с 600 потоками:

func1();
__syncthreads();
func2();
__syncthreads();

тогда ядро ​​должно работать дважды, и порядок выполнения будет:

  1. func1 выполняется для первых 512 потоков
  2. func2 выполняется для первых 512 потоков
  3. func1 выполняется для остальных потоков
  4. func2 выполняется для остальных потоков

Примечание:

Суть в том, __syncthreadsчто это операция всего блока, и она не синхронизирует все потоки.


Я не уверен в точном количестве потоков, которые __syncthreadsможно синхронизировать, так как вы можете создать блок с более чем 512 потоками и позволить варпу обрабатывать планирование. Насколько я понимаю, точнее сказать: func1 выполняется по крайней мере для первых 512 потоков.

До того, как я отредактировал этот ответ (еще в 2010 году), я измерял 14x8x32 потоков, синхронизированных с помощью __syncthreads.

Я был бы очень признателен, если бы кто-нибудь еще раз проверил это для получения более точной информации.

Бижан
источник
Что произойдет, если func2 () зависит от результатов func1 (). Я думаю, что это неправильно
Крис
@ Крис Я написал это семь лет назад, но если я правильно помню, я провел тест на это и пришел к выводу, что ядра с большим количеством потоков, чем у gpu, ведут себя таким образом. Если вам доведется проверить этот случай и достичь другого результата, мне придется удалить этот пост.
Бижан
Извините, я думаю, что это неправильно, также, что GPU может одновременно запускать только 112 потоков.
Стивен Лу
@ StevenLu ты пробовал это? также я не думаю, что 112 параллельных потоков имеют какой-либо смысл для графического процессора. 112 - количество потоковых процессоров. Я едва могу вспомнить CUDA сейчас :)
Бижан
1
@StevenLu здесь не проблема максимального количества потоков, это __syncthreadsоперация всего блока, и тот факт, что он фактически не синхронизирует все потоки, создает неудобства для учащихся CUDA. Поэтому я обновил свой ответ на основе информации, которую вы мне дали. Я очень ценю это.
Бижан