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

    Оригинал тут http://3.14.by/ru/read/cuda-crash-course-vvedenie

    Комментарии 26

      +1
      Основы CUDA — steps3d.narod.ru/tutorials/cuda-tutorial.html. Может тоже будет интересно просмотреть.
      Спасибо за статью, ещё бы про ATI чего нибудь прочитать.
        0
        Про Brook+ будет через 1-2 дня
          +1
          А про OpenCL будет?
            +1
            Да, но на данный момент только теория, т.к. пока как следует ни у кого ничего не работает (в смысле на видеокартах).
          • НЛО прилетело и опубликовало эту надпись здесь
          0
          А я жду GCC+Cuda =) И буду как в рекламе, компилировать, компилировать, компилировать, компилировать…
            0
            А что его ждать, берете и компилируете, все ж работает :-)
              0
              Те уже есть GCC написанный на CUDA? и он использует ресурсы карты? ХДЕ?
                0
                CUDA под линух использует GCC. А так чтобы без изменений программы начинали работать под CUDA — этого не будет никогда :-)
                  0
                  Ну так и ждем  GCC+CUDA =)
                    +2
                    Понял Вас, вы хотите чтобы все собранные в GCC программы работали по умолчанию не на CPU, а на GPU. К сожалению этого ждать придётся долго, наверное :). Да и в случае как с шейдерами, думается в будущем будет один API объединяющий CUDA + Brook.
                      0
                      Уже известно, что это будет OpenCL и DirectX 11
                        0
                        Не правильно поняли, я хочу что бы сам gcc работал на CUDA.
                          +1
                          CUDA не предназначено для решения таких задачь, это не эффективно (т.е. скорость в десятки раз медленнее чем на CPU).
              0
              Попробовал демку Elcomsoft Wireless Auditor, на ноутбуке с Intel T9300 (2.5 ГГц) и Nvidia 9500GS дало в среднем 5000 паролей в секунду. Так что польза ощутима :) Кстати они заявляют что поддерживают и ATI HD Series.
                0
                А как дела у AMD/ATI с подобным?
                  0
                  «Про Brook+ будет через 1-2 дня „
                    0
                    А про OpenCL напишите чтонить?
                      0
                      Читайте камменты :-)
                    0
                    У них Brook, ждите в следующей серии. Если коротко — там ассемблер, разработка менее удобна.
                      0
                      Нет, там С. Ассемблер в CAL/CTM
                    0
                    Карточки отличаются только объемом памяти\количеством процессоров? Функционал у всех одинаковый? Для девятой\двухсотой серии какие-то особые отличия есть?
                      0
                      В 9-й в основном поддержка атомарных операций добавилась.
                      В 200-й — 64-х битные вычисления (в 10 раз медленнее), больше регистров
                      +4
                      Статей такого типа обычно очень не хватает, когда начинаешь изучать что-то новое и сложное. Спасибо!
                      Как правило есть либо обрывки документации, либо монументальный труд, к концу прочтения которого уже поздно применять новую технологию :) Считаю, что Хабр нужен как раз для таких обзорных и одновременно детальных статей. Плюс в карму и за публикацию.
                        0
                        Проблема со специализированными статьями в том, что они интересны небольшому проценту хаброюзеров :-( Например моя статья про ноотропы набрала вдвое больше обоих статей про высокопроизводительные вычисления в сумме :-) Неблагодарная тема. Но тем не менее, серию я закончу :-)
                        0
                        У вас опечатка в:
                        __device__ — означает что переменная находится в глобальной памяти видеокарты (т.е. которой там 512-1024Мб и выше).
                        не __device__, а __global__.

                        Только полноправные пользователи могут оставлять комментарии. Войдите, пожалуйста.

                        Самое читаемое