Сваричевский Михаил – CUDA: С места в карьер

Многие видели моё введение в современные технологии высокопроизводительных вычислений и оценки производительности, теперь я продолжу тему более подробным рассказом о технологии CUDA. Для тех кто не смотрел предыдущие серии: CUDA позволяет писать и запускать на видеокартах nVidia(8xxx и выше) программы написанные на С++ со специальными расширениями. На правильных задачах достигается значительное превосходство по производительности на $ по сравнению с обычными CPU. Достижимая производительность – 1 трлн и выше операций в секунду на GTX295.

NB: Статья – краткое введение, покрыть все ньюансы программирования под CUDA в одной статье вряд ли возможно :-)

О железе

CUDA работает на видеокартых начиная с 8400GS и выше. Разные видеокарты имеют разые возможности. В целом, если вы видите что в видеокарте например 128 SP(Streaming Processor) – это значит что там 8 SIMD MP (multiprocessor), каждый из которых делает одновременно 16 операций. На один MP есть 16кб shared memory, 8192 штуки 4-хбайтных регистров (В картах серии GTX2xx значения больше). Также есть 64кб констант общие для всех MP, они кешируются, при непопадании в кеш – достаточно большая задержка (400-600 тактов). Есть глобальная память видеокарты, доступ туда не кешируется, и текстуры (кешируется, кеш оптимизирован для 2D выборок). Для использования нескольких видеокарт нужно во первый отключать SLI в дровах, а во вторых – на каждую видеокарту запускать по потоку, и вызывать cudaSetDevice().

С чего начать?

Самый быстрый способ научиться программировать на CUDA – это взять какой-нибуть пример из SDK, запустить его, и затем модифицировать, пока работает(собственно так я и делал, когда писал свой BarsWF ) :-) Для начала идем на http://www.nvidia.com/object/cuda_get.html и качаем SDK и Toolkit под вашу операционную систему нужной битности. (к сожалению например 32-х битный SDK и 64-хбитный toolkit мешать нельзя). Полезно обновить драйвер видеокарты до последней версии (т.к. CUDA быстро развивается, всегда полезно иметь последние дрова, и вам и пользователям ваших программ).Сдесь я буду рассматривать разработку под Windows в Visual Studio (2005, недавно с 2008 тоже стало можно). Для примера возьмем пример SDK Mandelbrot. Самое важное – это .cu файл, обратите внимание на его Custom Build Rule:
$(CUDA_BIN_PATH)\nvcc.exe -ccbin "$(VCInstallDir)bin" -c -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT -I"$(CUDA_INC_PATH)" -I./ -I../../common/inc -I"$(DXSDK_DIR)\Include" -o $(ConfigurationName)\Mandelbrot_sm10.obj Mandelbrot_sm10.cu

Его вы можете использовать во всех своих проектах, только вместо «../../common/inc» можно указать абсолютный путь (или переменную окружения). nvcc – это и есть великий и ужасный компилатор CUDA. На выходе он генерирует объектный файл, в котором уже включена откомпилированная программа для видеокарты. Обратите внимение на описание интерфейса в Mandelbrot_kernel.h – тут руками приходится описывать kernel-ы которые мы собираемся вызывать из основной С++ программы (впрочем их обычно не много, так что это не страшно). После того как вам удалось запустить пример SDK, можно рассмотреть, чем же CUDA программа отличается от обычной.

NB: Если вы добавите параметр -keep то после компиляции сможете найти много занимательных промежуточных файлов.

Определение функций

Перед функциями в .cu файле могут стоять следующие «модификаторы»:

__device__ – это означает, что функция выполняется только на видеокарте. Из программы выполняющейся на обычном процессоре (host) её вызвать нельзя.
__global__ – Эта функция – начало вашего вычислительного ядра. Выполняется на видеокарте, но запускается только с хоста.
__host__ – Выполняется и запускается только с хоста (т.е. обычная функция C++). Если вы перед функцией укажите например __host__ и __device__ – то будут скомпилированы 2 версии функции (не все комбинации допустимы).

Определение данных

__device__ – означает что переменная находится в глобальной памяти видеокарты (т.е. которой там 512-1024Мб и выше). Очень медленная память по меркам вычислений на видеокарте(хоть и быстрее памяти центрального процессора в несколько раз), рекомендуется использовать как можно реже. В этой памяти данные сохраняются между вызовами разных вычислительных ядер. Данные сюда можно записывать и читать из host-части с помощью

cudaMemcpy(device_variable, host_variable, size, cudaMemcpyHostToDevice); //cudaMemcpyDeviceToHost – в обратную сторону

__constant__ – задает переменную в константной памяти. Следует обратить внимание, что значения для констант нужно загружать функцией cudaMemcpyToSymbol. Константы доступны из всех тредов, скорость работы сравнима с регистрами(когда в кеш попадает).
__shared__ – задает переменную в общей памяти блока тредов (т.е. и значение будет общее на всех). Тут нужно подходить с осторожностью – компилятор агрессивно оптимизирует доступ сюда(можно придушить модификатором volatile), можно получать race condition, нужно использовать __syncthreads(); чтобы данные гарантированно записались. Shared memory разделена на банки, и когда 2 потока одновременно пытаются обратиться к одному банку, возникает bank conflict и падает скорость.

Все локальные переменные которые вы определеили в ядре (__device__) – в регистрах, самая высокая скорость доступа.

Как поток узнает над чем ему работать

Основая идея CUDA в том, что для решения вашей задачи вы запускаете тысячи и тысячи потоков, поэтому не стоит пугаться того что тут будет дальше написано :-)
Допустим, надо сделать какую-то операцию над картинкой 200x200. Картинка разбивается на куски 10x10, и на каждый пиксел такого кусочка запускаем по потоку. Выглядить это будет так:
dim3 threads(10, 10);//размер квардатика, 10*10
dim3 grid(20, 20);//сколько квадратиков нужно чтобы покрыть все изображение

your_kernel<<<grid, threads>>>(image, 200,200);//Эта строка запустит 40'000 потоков (не одновременно, одновременно работать будет 200-2000 потоков примерно).

В отличии от Brook+ от AMD, где мы сразу определяем какому потоку над какими данными работать, в CUDA все не так: передаваеиые kernel-у параметры одинаковые для всех потоков, и поток должен сам получить данные для себя, чтобы сделать это, потоку нужно вычислить, в каком месте изображения он находится. В этом помогают магические переменные blockDim, blockIdx.
const int ix = blockDim.x * blockIdx.x + threadIdx.x;
const int iy = blockDim.y * blockIdx.y + threadIdx.y;

В ix и iy – координаты, с помощью которых можно получить исходные данные из массива image, и записать результат работы.

Оптимизация

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

Как можно меньше используйте __global__ память.
При работе с __shared__ памятью избегайте конфлктов банков (впрочем многие задачи могут быть решены без shared памяти).
Как можно меньше ветвлений в коде, где разные потоки идут по разным путям. Такой код не выполняется параллельно.
Используйте как можно меньше памяти. Чем меньше памяти вы используете, тем агрессивнее компилятор и железо смогут запускать ваш kernel (например он может взять 100 тредов, и используя в 100 больше регистров запустить одновременно на одном MP, радикально уменьшая задержки)

Не получается?

В первую очередь следует прочитать документацию вместе с SDK (NVIDIA_CUDA_Programming_Guide, CudaReferenceManual, ptx_isa), после этого можно спросить на официальном форуме – там даже девелоперы nVidia часто отписываются, да и вообще много умных людей. По русски можно спросить у меня на форуме например, где отвечу я :-) Также много людей обитает на gpgpu.ru.

Надеюсь это введение поможет людям, решившим попробовать программирование для видеокарт. Если есть проблемы/вопросы – буду рад помочь. Ну а в переди нас ждет введение в Brook+ и SIMD x86.