Многие видели моё введение в современные технологии высокопроизводительных вычислений и оценки производительности, теперь я продолжу тему более подробным рассказом о технологии 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:
Его вы можете использовать во всех своих проектах, только вместо "../../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-части с помощью
__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 раз быстрее :-) )
В первую очередь следует прочитать документацию вместе с SDK (NVIDIA_CUDA_Programming_Guide, CudaReferenceManual, ptx_isa), после этого можно спросить на официальном форуме — там даже девелоперы nVidia часто отписываются, да и вообще много умных людей. По русски можно спросить у меня на форуме например, где отвечу я :-) Также много людей обитает на gpgpu.ru.
Надеюсь это введение поможет людям, решившим попробовать программирование для видеокарт. Если есть проблемы/вопросы — буду рад помочь. Ну а в переди нас ждет введение в Brook+ и SIMD x86
Оригинал тут http://3.14.by/ru/read/cuda-crash-course-vvedenie
Для тех кто не смотрел предыдущие серии: 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
Оригинал тут http://3.14.by/ru/read/cuda-crash-course-vvedenie
