Теория: CUDA с аппаратной точки зрения
Если вы регулярно читаете Tom's Hardware Guide, то архитектура последних GPU от nVidia вам знакома. Если нет, мы рекомендуем ознакомиться со статьёй "
Как можно видеть по иллюстрации выше, ядро шейдеров nVidia состоит из нескольких кластеров текстурных процессоров (Texture Processor Cluster, TPC). Видеокарта 8800 GTX, например, использовала восемь кластеров, 8800 GTS - шесть и т.д. Каждый кластер, по сути, состоит из текстурного блока и двух потоковых мультипроцессоров (streaming multiprocessor). Последние включают начало конвейера (front end), выполняющее чтение и декодирование инструкций, а также отсылку их на выполнение, и конец конвейера (back end), состоящий из восьми вычислительных устройств и двух суперфункциональных устройств SFU (Super Function Unit), где инструкции выполняются по принципу SIMD, то есть одна инструкция применяется ко всем потокам в варпе. nVidia называет такой способ выполнения SIMT (single instruction multiple threads, одна инструкция, много потоков). Важно отметить, что конец конвейера работает на частоте в два раза превосходящей его начало. На практике это означает, что данная часть выглядит в два раза "шире", чем она есть на самом деле (то есть как 16-канальный блок SIMD вместо восьмиканального). Потоковые мультипроцессоры работают следующим образом: каждый такт начало конвейера выбирает варп, готовый к выполнению, и запускает выполнение инструкции. Чтобы инструкция применилась ко всем 32 потокам в варпе, концу конвейера потребуется четыре такта, но поскольку он работает на удвоенной частоте по сравнению с началом, потребуется только два такта (с точки зрения начала конвейера). Поэтому, чтобы начало конвейера не простаивало такт, а аппаратное обеспечение было максимально загружено, в идеальном случае можно чередовать инструкции каждый такт - классическая инструкция в один такт и инструкция для SFU - в другой.
Каждый мультипроцессор обладает определённым набором ресурсов, в которых стоит разобраться. Есть небольшая область памяти под названием "Общая память/Shared Memory", по 16 кбайт на мультипроцессор. Это отнюдь не кэш-память: программист может использовать её по своему усмотрению. То есть, перед нами что-то близкое к Local Store у SPU на процессорах Cell. Данная деталь весьма любопытная, поскольку она подчёркивает, что CUDA - это комбинация программных и аппаратных технологий. Данная область памяти не используется для пиксельных шейдеров, что nVidia остроумно подчёркивает "нам не нравится, когда пиксели разговаривают друг с другом".
Данная область памяти открывает возможность обмена информацией между потоками в одном блоке. Важно подчеркнуть это ограничение: все потоки в блоке гарантированно выполняются одним мультипроцессором. Напротив, привязка блоков к разным мультипроцессорам вообще не оговаривается, и два потока из разных блоков не могут обмениваться информацией между собой во время выполнения. То есть пользоваться общей памятью не так и просто. Впрочем, общая память всё же оправданна за исключением случаев, когда несколько потоков попытаются обратиться к одному банку памяти, вызывая конфликт. В остальных ситуациях доступ к общей памяти такой же быстрый, как и к регистрам.
Общая память - не единственная, к которой могут обращаться мультипроцессоры. Они могут использовать видеопамять, но с меньшей пропускной способностью и большими задержками. Поэтому, чтобы снизить частоту обращения к этой памяти, nVidia оснастила мультипроцессоры кэшем (примерно 8 кбайт на мультипроцессор), хранящим константы и текстуры.
Мультипроцессор имеет 8 192 регистра, которые общие для всех потоков всех блоков, активных на мультипроцессоре. Число активных блоков на мультипроцессор не может превышать восьми, а число активных варпов ограничено 24 (768 потоков). Поэтому 8800 GTX может обрабатывать до 12 288 потоков в один момент времени. Все эти ограничения стоило упомянуть, поскольку они позволяют оптимизировать алгоритм в зависимости от доступных ресурсов.
Оптимизация программы CUDA, таким образом, состоит в получении оптимального баланса между количеством блоков и их размером. Больше потоков на блок будут полезны для снижения задержек работы с памятью, но и число регистров, доступных на поток, уменьшается. Более того, блок из 512 потоков будет неэффективен, поскольку на мультипроцессоре может быть активным только один блок, что приведёт к потере 256 потоков. Поэтому nVidia рекомендует использовать блоки по 128 или 256 потоков, что даёт оптимальный компромисс между снижением задержек и числом регистров для большинства ядер/kernel.
Теория: CUDA с программной точки зрения
С программной точки зрения CUDA состоит из набора расширений к языку C, что напоминает BrookGPU, а также нескольких специфических вызовов API. Среди расширений присутствуют спецификаторы типа, относящиеся к функциям и переменным. Важно запомнить ключевое слово __global__, которое, будучи приведённым перед функцией, показывает, что последняя относится к ядру/kernel - эту функцию будет вызывать CPU, а выполняться она будет на GPU. Префикс __device__ указывает, что функция будет выполняться на GPU (который, кстати, CUDA и называет "устройство/device") но она может быть вызвана только с GPU (иными словами, с другой функции __device__ или с функции __global__). Наконец, префикс __host__ опционален, он обозначает функцию, которая вызывается CPU и выполняется CPU - другими словами, обычную функцию.
Есть ряд ограничений, связанных с функциями __device__ и __global__: они не могут быть рекурсивными (то есть вызывать самих себя), и не могут иметь переменное число аргументов. Наконец, поскольку функции __device__ располагаются в пространстве памяти GPU, вполне логично, что получить их адрес не удастся. Переменные тоже имеют ряд квалификаторов, которые указывают на область памяти, где они будут храниться. Переменная с префиксом __shared__ означает, что она будет храниться в общей памяти потокового мультипроцессора. Вызов функции __global__ немного отличается. Дело в том, при вызове нужно задать конфигурацию выполнения - более конкретно, размер сетки/grid, к которой будет применено ядро/kernel, а также размер каждого блока. Возьмём, например, ядро со следующей подписью.
__global__ void Func(float* parameter); |
Оно будет вызываться в виде
Func<<< Dg, Db >>> (parameter); |
где Dg является размером сетки, а Db - размером блока. Две этих переменных относятся к новому типу вектора, появившегося с CUDA.
API CUDA содержит функции для работы с памятью в VRAM: cudaMalloc для выделения памяти, cudaFree для освобождения и cudaMemcpy для копирования памяти между RAM и VRAM и наоборот.
Мы закончим данный обзор весьма интересным способом, которым компилируется программа CUDA: компиляция выполняется в несколько этапов. Сначала извлекается код, относящийся к CPU, который передаётся стандартному компилятору. Код, предназначенный для GPU, сначала преобразовывается в промежуточный язык PTX. Он подобен ассемблеру и позволяет изучать код в поисках потенциальных неэффективных участков. Наконец, последняя фаза заключается в трансляции промежуточного языка в специфические команды GPU и создании двоичного файла.