Мы обсудили мотивацию и преимущества параллельной обработки в части 1; В этой части я пытаюсь объяснить организацию потоков внутри графического процессора и то, как кодировать графический процессор. Обратите внимание, что все, что обсуждается в этой серии статей, касается графических процессоров Nvidia. Мы собираемся использовать CUDA API от Nvidia, который используется для программирования графических процессоров Nvidia. Для простоты Nvidia расширила традиционный язык C и добавила специальные функции для программирования CUDA. Функции могут быть включены с помощью заголовочного файла cuda.h. Расширение .cu используется для программ cuda. Nvidia предоставляет компилятор NVCC (компилятор Nvidia C) для компиляции программ cuda. NVCC зависит от внешнего компилятора C, такого как GCC, для его работы. Прежде чем приступить к программированию cuda, давайте выучим несколько терминов.
Хост: процессор называется хостом.
Устройство: графический процессор называется устройством. Никакая программа не может полностью работать на устройстве; каждый исходный файл cuda обычно содержит как код хоста, так и код устройства.
Ядро: функция, которая выполняется параллельно на устройстве.
Основы архитектуры Cuda
Графический процессор с поддержкой cuda организован в массив потоковых мультипроцессоров (SM). Количество SM может варьироваться от одного графического процессора к другому. Каждый SM имеет несколько потоковых процессоров, которые совместно используют логику управления и кэш инструкций. У каждого ядра есть свои регистры. Каждый графический процессор поставляется с собственной графической памятью DRAM с удвоенной скоростью передачи данных (которую обычно называют глобальной памятью), которая совместно используется всеми модулями SM.
Моя Nvidia GEFORCE GTX 1050 имеет 5 модулей SM, каждый из которых имеет 128 ядер Cuda, 2000 МБ глобальной памяти и 49152 байта общей памяти на блок (мы рассмотрим общую память и блоки позже).
Организация потоков внутри графического процессора
Напомним, что графический процессор - это процессор с множеством ядер, способный выполнять множество потоков параллельно. Функция ядра при запуске на графическом процессоре запускает сетку потоков. Каждая сетка состоит из блоков, которые, в свою очередь, состоят из потоков, каждый поток запускает копию предоставленного ядра.
Количество блоков и количество потоков в блоке (таким образом, общее количество потоков) можно контролировать с помощью программисту, и найти оптимальные значения для этих двух часто бывает сложно и зависит от целевого графического процессора. Потоки, а также блоки могут быть организованы в одномерном, двухмерном или трехмерном виде (по выбору программиста) в зависимости от входных данных. Позже вы поймете необходимость многомерной организации.
Блоки назначаются SM, и количество блоков, назначенных в данный момент времени, превышает то, что может быть обработано (для использования скрытия задержки)
Я уверен, что вы, должно быть, думаете, выполняйте ли все потоки (выполняемые в a time) выполняются точно параллельно, то есть в заданный момент времени все ядра обрабатывают одну и ту же инструкцию. Ответ на ваше удивление - нет. Даже потоки в блоке могут быть не совсем параллельными. Мы также ничего не можем сказать о порядке выполнения потоков.
Деформация как единица планирования потоков:
Потоки внутри графического процессора сгруппированы в деформации, каждый поток в деформации выполняется точно параллельно, они следуют модели SIMD (Single Instruction Multiple Data), то есть команда выбирается и затем выполняется для всех потоков в деформации. В современном графическом процессоре размер деформации обычно равен 32; если количество оставшихся потоков меньше 32, то для завершения цикла используются пустые потоки. Теперь мы также можем понять, что лучше иметь количество потоков в блоке, кратное 32, чтобы ни одно ядро не тратилось впустую. Возможно, вы думаете о том, зачем было организовывать блоки. Поверьте мне; вы очень оцените блоки после того, как поймете использование разделяемой памяти. Теперь мы готовы написать нашу первую программу cuda.
Cuda Программа для вычисления векторной суммы
Мы легко видим параллелизм в вычислении векторной суммы; одна и та же функция (сумма) применяется к разным частям вектора в соответствии с моделью SIMD.
Программа cuda имеет следующие три типа функций; специальные ключевые слова используются для обозначения различных функций:
Итак, схема нашей программы, которая вычисляет векторную сумму двух массивов A и B и сохраняет ее в C, выглядит следующим образом:
Заголовочный файл cuda.h предоставляет нам функции для выделения, удаления глобальной памяти устройства и перемещения данных между устройством и глобальной памятью хоста. Это следующие функции:
cudaMalloc (void ** ptr, int size): выделить глобальную память устройства, принимает 2 параметра, адрес указателя void и размер выделяемой памяти. Базовый адрес выделенной памяти сохраняется в указателе (адрес которого вы указали).
cudaFree (void * ptr): освободить память устройства.
cudaMemcpy (destination_ptr, source_ptr, size, type): скопировать данные туда и обратно между двумя глобальными ячейками памяти. Тип параметра определяет поток, он может принимать значения cudaMemcpyHostToDevice и cudaMemcpyDeviceToHost.
Каждый поток в сетке имеет несколько предопределенных переменных: threadIdx, blockIdx, blockDim, gridDim, каждая из которых имеет тип dim3. Dim3 - это целочисленный вектор с тремя измерениями (или элементами): x, y и z.
Каждый блок можно однозначно идентифицировать в сетке по его координатам: blockId.x, blockId.y, blockId.z. Каждый поток можно однозначно идентифицировать в блоке по его координатам: threadId.x, threadId.y, threadId.z. Используя эти два и {blockDim.x, blockDim.y, blockDim.z}, мы можем однозначно идентифицировать каждый поток в сетке. Помните, что каждый поток запускает одно и то же ядро, поэтому эта уникальная идентификация также используется для выборки соответствующих данных, которые будут обрабатываться потоком, например, для одномерных данных поток может загрузить свои соответствующие данные, используя смещение blockIdx.x * blockDim.x + threadId.x с базовым адресом.
Размеры блока и сетки называются конфигурацией исполнения и могут быть установлены при запуске ядра. Мы используем две переменные dim3 для задания конфигурации. Если блоки / поток одномерные, то мы можем просто предоставить положительное целое число, размер x, для других измерений по умолчанию установлено значение 1. Следующий код
сетка dim3 (256,2,1);
блок dim3 (16,4,1);
фукернал ‹-------------------------------- сетка, блок ››› (para1, para2)
запускает ядро на 16 * 4 * 1 блоках, каждый из которых состоит из 256 * 2 * 1 потоков.
Завершенные функции vecAddKernal и vecAdd:
Итак, в следующий раз, когда вы добавите два массива, вы собираетесь использовать свой графический процессор? На самом деле не стоит; накладные расходы на перенос массивов из глобальной памяти хоста в глобальную память устройства велики, а операции, выполняемые с массивом, меньше, что делает использование GPU в этом случае невыгодным; но графический процессор определенно будет работать лучше в случае умножения матриц.
