Должен ли я проверять количество потоков в коде ядра?

Question

Я новичок с CUDA, и мои коллеги всегда дизайн ядра с помощью следующей упаковки:

__global__ void myKernel(int nbThreads) 
{ 
    int threadId = blockDim.x*blockIdx.y*gridDim.x //rows preceeding current row in grid 
      + blockDim.x*blockIdx.x    //blocks preceeding current block 
      + threadIdx.x; 

    if (threadId < nbThreads) 
    { 
     statement(); 
     statement(); 
     statement(); 
    } 
} 

Они думают, что есть некоторые ситуации, когда CUDA может запустить больше потоков, чем указано ради выравнивания/сновальной, так что нам нужно проверять его каждый раз. Тем не менее, я пока не видел ни одного примера ядра в Интернете, где они действительно делают эту проверку.

Может ли CUDA запускать больше потоков, чем заданные размеры блока/сетки?

Answer 1

CUDA не будет запускать больше потоков, чем те, которые указаны размером блока/сетки.

Однако из-за гранулярности размеров блока (например, желательно иметь размеры блоков не более 32 и ограничено по размеру 1024 или 512), часто бывает трудно сопоставить сетка нитей должна быть численно равна требуемому размеру проблемы.

В этих случаях типичное поведение заключается в том, чтобы запускать больше потоков, эффективно округляя до следующего четного размера на основе гранулярности блоков и использовать код «проверки потока» в ядре, чтобы убедиться, что «дополнительные потоки », то есть те, которые находятся за пределами размера проблемы, ничего не делают.

В вашем примере, это можно было бы уточнить, написав:

__global__ void myKernel(int problem_size) 


    if (threadId < problem_size) 

который сообщается, что подразумевается, что только потоки, соответствующего размеру задачи (которые могут не соответствовать запущенному размеру сетки) делать какую-либо реальную работу ,

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

Если я хочу, чтобы каждый threadblock был кратным 32, нет ни одного количества блоков нити, которые я могу выбрать, которые дадут мне ровно 10000 потоков. Поэтому я мог бы выбрать 256 потоков в потоковом блоке и запустить 40 потоков, предоставляя мне всего 10240 потоков. Используя проверку потока, я препятствую «лишним» 240 нитям делать что-либо.