Молодогвардейцев 454015 Россия, Челябинская область, город Челябинск 89085842764
MindHalls logo

Вычисление глобального индекса нити в CUDA

Доброго времени суток! Я уже несколько раз затрагивал тему параллельных вычислений, но все время это были вычисления на старом добром центральном процессоре. А сейчас я вкратце расскажу о технологии CUDA от фирмы NVIDIA, которая расшифровывается как Compute Unified Device Architecture, и служит для того, чтобы выполнять параллельные вычисления на процессоре графического устройства, а именно — видеокарточки NVIDIA.

Grid, Block и Thread в CUDA

Для того, чтобы эффективно программировать с помощью технологии CUDA, в первую очередь нужно освоиться и научиться свободно оперировать основными понятиями вычислительной модели. Их три: грид(grid), блок(block) и нить(thread), последний еще называют потоком, это синонимы.

Grid является самым верхним уровнем абстракции в модели, он трехмерный и состоит из блоков. Например, в гриде размером 2x2x2 содержится 8 блоков. В свою очередь, каждый блок содержит в себе трехмерный массив нитей(threads), которые и являются непосредственными исполнителями вычислений. К примеру, если каждый из 8 блоков содержит 2x2x1 нитей, то всего на устройстве будет 32 рабочих потока, готовых приступить к любым вычислениям, причем работать они будут параллельно.

Такой размах в количестве вычислительных потоков дает большую свободу действий, ограничивающуюся только вашей фантазией. Ну и еще некоторыми ограничениями, индивидуальными для каждой видеокарты. Посмотреть характеристики вашей карты можно вызвав функцию cudaGetDeviceProperties(), которая заполнит структуру типа cudaDeviceProp.

Пример того, как можно посмотреть характеристики CUDA устройств в компьютере

const int kb = 1024;
const int mb = kb * kb;

//Получить количество устройств
int devCount;
cudaGetDeviceCount(&devCount);
cout << "CUDA Devices: " << endl << endl;

//Для каждого устройства вывести его характеристики
for (int i = 0; i < devCount; ++i) {
 cudaDeviceProp props;
 cudaGetDeviceProperties(&props, i);
 
 cout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl;
 cout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl;
 cout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl;
 cout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl;
 cout << " Block registers: " << props.regsPerBlock << endl << endl;

 cout << " Warp size: " << props.warpSize << endl;
 cout << " Threads per block: " << props.maxThreadsPerBlock << endl;
 cout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << endl;
 cout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl;
 cout << endl;
}

Вычисление индексов блока и нити

Для того, чтобы написать эффективный алгоритм, например, с конфигурацией Grid:2x2x2 и Block:2x2x1 в которой 32 нити, нужно точно знать глобальный индекс каждой рабочей нити. Глобальным индексом в этом случае будет число от 0 до 31.

Посмотрим исходный код функции, которая вычислит индексы и выведет их на экран

__global__ void showIndex() {
 //Индекс текущего блока в гриде
 int blockIndex = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.y*gridDim.x;
 //Индекс треда внутри текущего блока
 int ThreadIndex = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.y*blockDim.x;

 //глобальный индекс нити
 int GlobalThreadIndex = blockIndex*blockDim.x*blockDim.y*blockDim.z + ThreadIndex;
 
 printf("block: (%d,%d,%d) threads: (%d,%d,%d) blockIndex = %d ThreadIndex = %d GlobalThreadIndex= %d\n ", blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y, threadIdx.z, blockIndex, ThreadIndex, GlobalThreadIndex);
}

Запускать это чудо можно на любой конфигурации, не только на той, о которой я говорил ранее. Кстати, о запуске, на всякий случай приведу пример того, как вызвать эту функцию из main.

Исходный код вызова функции подсчета индексов

int main(int argc, char* argv[]) {
 dim3 dimGrid(2, 2, 2); //Размерность грида
 dim3 dimBlock(2, 2, 1); //Размерность блока 

 //Вызов функции с нужной конфигурацией
 showIndex <<< dimGrid, dimBlock >>>();

 return 0;
}

Результат работы функции я привел на скриншоте

Глобальный индекс нити в CUDA

Заключение

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