Доброго времени суток! Я уже несколько раз затрагивал тему параллельных вычислений, но все время это были вычисления на старом добром центральном процессоре. А сейчас я вкратце расскажу о технологии 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;
}
Результат работы функции я привел на скриншоте
Заключение
Мы смогли получить глобальные индексы всех нитей на произвольной конфигурации запуска графического процессора. Это позволит эффективно реализовывать различные параллельные алгоритмы, например, умножение матриц, в котором за вычисление одного элемента в результирующей матрице будет отвечать одна нить. Но это в следующий раз, а на сегодня у меня все, спасибо за внимание!