OpenCL. Практика



    Здравствуй, уважаемое хабрасообщество.

    В предыдущих статьях мы рассмотрели OpenCL в целом, потом подробно вникли в суть стандарта и разобрали на каких идеях базируется эта технология.
    OpenCL. Что это такое и зачем он нужен? (если есть CUDA)
    OpenCL. Подробности технологии
    Теперь настало время пощупать эту технологию живьем.

    Приготовления


    Итак, для работы нам понадобится: спецификация стандарта, SDK (AMD или NVidia) и, опционально, литература по OpenCL, например, отсюда.
    Если вы устанавливаете Nvidia Computing SDK – вы автоматически получите все нужные документы. Кроме того бонусом Вы получите много интересных примеров программ (30 штук в последнем релизе SDK). Благодаря этим примерам легко научиться правильно использовать OpenCL, использовать несколько OpenCL устройств одновременно, пользоваться связкой OpenCL-OpenGL (это взаимодействие оговорено стандартом!) итд.

    Компилятор OpenCL встроен в драйвер, поэтому выбор IDE для разработки никак не ограничен, посему не буду описывать процесс настройки какой-то определенной IDE. Все что Вам надо сделать – это прописать пути до заголовков и библиотек, которые установит SDK.

    Поехали.


    Напишем простую программу для суммирования двух векторов. Такая программа есть в примерах SDK для CUDA и OpenCL, но наша программа будет немного отличаться (уберем проверки кодов ошибок на каждом шаге, и немного упростим программу, оставив только самую суть).

    GPU часть


    Начнем с самого интересного, и, пожалуй, самого простого в данном примере – с кода, который будет исполнен на GPU.

    Синтаксис OpenCL для написания kernel'ов сам по себе не представляет собой ничего особенного и слабо отличается от синтаксиса той же CUDA – это все старый добрый С с небольшими модификациями.
    Создадим файл vectorAdd.cl – здесь будет располагаться наш kernel.
    __kernel void VectorAdd(__global const float* a, __global const float* b, __global float* c, int iNumElements)
    {
      // get index into global data array
      int iGID = get_global_id(0);

      // bound check (equivalent to the limit on a 'for' loop for standard/serial C code
      if (iGID >= iNumElements)
      { 
        return;
      }
      
      // add the vector elements
      c[iGID] = a[iGID] + b[iGID];
    }


    * This source code was highlighted with Source Code Highlighter.


    Получили свой номер в глобальном пространстве индексов и сложили элементы векторов с соответствующим индексом, а если наш номер больше размера вектора – ничего не делам.

    Все выглядит легко и просто: kernel – простая функция, объявление которой предваряется ключевым словом __kernel (два подчеркивания) а дальше все как в С – возвращаемый тип, название функции, параметры (при определении параметров так же необходимо указать модификаторы __global, __local, __private).
    kernel пишется на языке С. Существует ряд расширений (кроме синтаксиса) и ограничений. Кратко об ограничениях можно прочитать тут. Более полно в стандарте, так же полезен может быть OpenCL Programming Guide
    Расширениями языка являются: тип данных «изображение» 2d и 3d, типы данных вектор интов, флоатов итп. размерности от 2 до 4.
    При объявлении переменных надо указывать область памяти, где они должны располагаться: __global, __local, __private. Если область памяти не указана – будет использована private-память.

    Если в kernel необходимо использовать другие функции, скрытые для CPU то можно определить их в том же файле, но без указания модификатора __kernel.

    Хостовая часть


    Простейший kernel мы создали. Теперь давайте разберемся, как запустить этот kernel на видеокарте.

    Хостовая часть программы тоже будет проста и ограничится запуском kernel.

    Функции для работы с kernel предоставляет OpenCL API. Это С-функции. Тут можно скачать C++-bindings и документацию к ним.

    Мануал по всем функциям API находится в том же документе, где описывается стандарт OpenCL.


    Для работы любого kernel необходим контекст, в котором он будет исполняться. Создадим объект «контекст».
    cl_context cxGPUContext = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, &ciErr1);

    * This source code was highlighted with Source Code Highlighter.

    первый параметр – список свойств контекста и их значений. NULL означает использование дефолтных implementation-defined свойств.
    Далее объясняем системе, что собираемся работать с GPU – это означает, что устройство с которым мы будем работать, может быть так же использовано и для 3d API, например OpenGL. (список возможных значений этого параметра есть так же в спецификации стандарта)
    Следующие два параметра нужны для регистрации call-back-функции, которая будет вызвана OpenCL в случае появления ошибок в контексте.
    Последний параметр – для возврата кода ошибки. Может быть NULL.

    На самом деле это не единственный способ создания контекста. Но статья и не претендует на описание всех функций OpenCL API. Просто такой создания контекста способ нам более удобен.

    Далее выберем устройство (у меня в системе оно всего одно, но на будущее пусть в нашей программе используется устройство с максимальным числом FLOPS).
    cl_device_id cdDevice = oclGetMaxFlopsDev(cxGPUContext);

    * This source code was highlighted with Source Code Highlighter.

    Отмечу, что в различных примерах из SDK весь процесс инициализации порой различается, может быть это сделано намеренно, дабы показать, что существует не один способ выполнять данные действия и заставить разработчика покопаться в спецификациях. К примеру, тут мы выбрали устройство с максимальными FLOPS, но мы могли бы воспользоваться функцией clGetContextInfo для получения списка всех устройств, ассоциированных с контекстом (см. оригинальный пример VectorAdd).


    Выбрали и инициализировали устройство.
    Теперь свяжем с нашим устройством очередь команд.
    cl_command_queue cqCommandQue = clCreateCommandQueue(cxGPUContext, cdDevice, 0, &ciErrNum);

    * This source code was highlighted with Source Code Highlighter.


    Из интересных параметров только 0. На самом деле это список свойств очереди команд: можно ли выполнять команды не последовательно и разрешено ли профилирование команд.

    Для работы с устройством все готово, мы можем отправлять команды в очередь для исполнения.
    Создадим объекты памяти, через которые будут связаны области памяти на устройстве и хосте.
    cl_mem cmDevSrcA = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &ciErr1);
    cl_mem cmDevSrcB = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &ciErr2);
    cl_mem cmDevDst = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &ciErr2);


    * This source code was highlighted with Source Code Highlighter.


    указываем тип доступа к памяти объектов (для устройства), размер области памяти и область памяти хоста с которой связан объект памяти (тут NULL).
    Инициализировать входные данные можно было бы и сразу, если предпоследним параметром передать указатель на область памяти хоста, которую надо скопировать на устройство.
    Но мы сделаем это позже, перед самым запуском kernel, чтобы не занимать место на устройстве раньше времени.

    Все подготовительные работы завершены, теперь мы примемся за сам kernel. Как Вы помните, компилятор OpenCL встроен прямо в платформу. По этому причине сборка OpenCL-kernel должна осуществляться во время исполнения (собирать kernel можно как из исходников так и из бинарников).

    Приступим.


    Создадим kernel из того файла, который мы написали ранее.
    char *source = oclLoadProgSource(source_path, "", &program_length);

    * This source code was highlighted with Source Code Highlighter.

    Получили исходник программы в строке char*. Source_path – полный путь до файла vectorAdd.cl, далее следует «преамбула» — обычно это header или список define'ов. Последний параметр – размер выходной строки.

    ВНИМАНИЕ!!! oclLoadProgSource — не является функцией OpenCL API, а находятся в вспомогательной библиотеке, поставляемой вместе с Nvidia Computing SDK.

    Создаем объект программы из полученных исходников, последующие функции – это OpenCL API.
    cl_program cpProgram = clCreateProgramWithSource(cxGPUContext, 1,(const char **)&source, &program_length, &ciErrNum);

    * This source code was highlighted with Source Code Highlighter.

    программа может состоять из нескольких файлов, каждый из которых необходимо загрузить в отдельную строку char*, массив таких строк мы передаем для создания программы. Второй параметр тут означает размер этого массива. В нашем случае – 1.
    После массива строк передается массив длин этих строк.
    Все остальные параметры не заслуживают внимания.

    Слепили программу из кучи файлов, теперь давайте ее соберем (компиляция и линковка)
    сlBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL);

    * This source code was highlighted with Source Code Highlighter.

    Собирает программу для выбранных устройств из списка устройств связанных с контекстом (контекст не передаем, так как он связан с объектом программы).
    Тут можно указать для каких из устройств связанных с контекстом надо собрать программу. Так же можно установить опции компилятора и настроить call-back-функцию для возможности асинхронной компиляции.
    У нас пока всего одно устройство, и мы можем позволить себе синхронную компиляцию. И опциями компилятора пользоваться пока не станем.
    Все это подробно описано в спецификации стандарта.После сборки с нашим объектом программы уже связан готовый исполняемый файл. Теперь из функции (а это ведь пока просто функция) с идентификатором __kernel надо создать kernel.
    cl_kernel ckKernel = clCreateKernel(cpProgram, "VectorAdd", &ciErr1);

    * This source code was highlighted with Source Code Highlighter.

    Все готово, все приготовления завершены и настал момент истины: запуск kernel на исполнение.
    НО! Предварительно надо установить входные параметры, с которыми будет вызван наш kernel.
    clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void*)&cmDevSrcA);
    clSetKernelArg(ckKernel, 1, sizeof(cl_mem), (void*)&cmDevSrcB);
    clSetKernelArg(ckKernel, 2, sizeof(cl_mem), (void*)&cmDevDst);
    clSetKernelArg(ckKernel, 3, sizeof(cl_int), (void*)&iNumElements);


    * This source code was highlighted with Source Code Highlighter.

    указываем порядковый номер параметра, размер и объект памяти.

    Вот теперь точно все. Начинается работа с очередью:
    Скопируем (асинхронно; за это отвечает третий аргумент) данные на устройство.
    clEnqueueWriteBuffer(cqCommandQue, cmDevSrcA, CL_FALSE, 0, sizeof(cl_float) * szGlobalWorkSize, srcA, 0, NULL, NULL);
    clEnqueueWriteBuffer(cqCommandQue, cmDevSrcB, CL_FALSE, 0, sizeof(cl_float) * szGlobalWorkSize, srcB, 0, NULL, NULL);


    * This source code was highlighted with Source Code Highlighter.

    Самое главное – поставить kernel на исполнение.
    clEnqueueNDRangeKernel(cqCommandQue, ckKernel, 1, NULL, &szGlobalWorkSize, NULL, 0, NULL, NULL);

    * This source code was highlighted with Source Code Highlighter.

    С первыми двумя параметрами все понятно.
    Третий параметр – размерность пространства индексов. Вектор – одномерный.
    За ним следует аргумент, который означает размер сдвига в пространстве индексов и в текущей версии стандарта должен быть всегда NULL.
    szGlobalWorkSize указывает размер пространства индексов — это общее количество work-item'ов, которые будут выполняться.
    Размер группы оставляем на усмотрение драйверу (NULL).
    Следующие два парметра используются для синхронизации при использовании out-of-order исполнения команд. Это список событий, которые должны завершиться перед запуском этой команды (сначала идет размер списка, потом сам список).
    Через последний параметр возвращается объект-событие, сигнализирующее о завершении команды.

    Осталось только прочитать результат. Сделаем это синхронно:
    clEnqueueReadBuffer(cqCommandQue, cmDevDst, CL_TRUE, 0, sizeof(cl_float) * szGlobalWorkSize, dst, 0, NULL, NULL);

    * This source code was highlighted with Source Code Highlighter.

    Теперь осталось очистить память, удалив созданные объекты памяти и программ. Это не сложно, и легко найти в любом примере из SDK, поэтмоу я не буду приводить здесь этот код.

    Заключение


    Простейшая программа готова.

    Дальше можно изучать спецификации вглубь, изучать особенности конкретных реализаций OpenCL, пробовать различные опции оптимизации в компиляторах, читать документы, выпускаемые поставщиками этих реализаций и, конечно же, писать более сложные программы.

    Существует множество нюансов и подводных камней при написании OpenCL программ, так же на данный момент доступен профилировщик OpenCL приложений раработанный Nvidia, предоставляющий ряд интересных возможностей. Но это все уже выходит за рамки данной статьи и, если уважаемые читатели проявят интерес, про различные тонкости и особенности OpenCL приложений и OpenCL для Nvidia GPU можно будет написать отдельную статью.

    Cписок полезных ссылок


    www.nvidia.com/object/cuda_opencl.html — тут можно зарегистрироваться для получения доступа к Nvidia Computing SDK и скачать полезные документы
    www.khronos.org/registry/cl — страница на сайде Khronos Group. Спецификации, заголовочные файлы итд.
    developer.amd.com/GPU/ATISTREAMSDKBETAPROGRAM/Pages/default.aspx — AMD Stream SDK с поддержкой OpenCL.
    Настоятельно рекомендую ознакомиться с двумя документами Nvidia: OpenCL Programming Guide и OpenCL Best Practices Guide.

    Поделиться публикацией

    Похожие публикации

    Комментарии 44
      0
      Спасибо. Жалко видеокарта у меня старая, обязательно пощупал бы.

      >про различные тонкости и особенности OpenCL приложений и OpenCL для Nvidia GPU можно будет написать отдельную статью

      Опять у Nvidia/Ati всё по-разному работает?
        +1
        Работает, я надеюсь, все же одинаково. В целом.
        Но ведь на производительность сильно влияет конкретная архитектура устройства. И вот тут уже могут вылезти особенности и даже ошибки.
        Пример: на моей видеокарте NVidia 8192 регистра, то есть если задать размер группы 512 элементов (максимум для моей видеокарты), то на каждый из них по 16 регистров. Если 16 регистров недостаточно для выполнения kernel — полетит ошибка во время исполенния kernel.
        То есть если ужать kernel в 16 регистров не выйдет — прийдется задействовать карту не на полную мощьность.

        Для чипов AMD, возможно, такое ограничение будет более сильное или слабое, а может его и вовсе не будет, все зависит от архитектуры.

        Такие вещи специфичны для конкретного устройства и знание некоторых может помочь ускорить работу приложения, а, иногда, и сберечь нервы)
          0
          Если не секрет, как вы посчитали регистры? Взяли из ТТХ?
            +1
            Да, это из ТТХ.
            В 260-х и старше — регистров больше…
            Рамер группы 512 — не обязательно так много, можно и меньше, если много работы с памятью нет, падение будет крайне незначительное.
              0
              У меня сейчас валяются N260GTX и 4870 ZT.
                0
                Комментарий оборвался. Так вот, у меня сейчас валяются 260-ка и 4870, если есть идеи, как можно погонять тестами монстров — готов выслушать (или в ПМ). Величайшим достижением для вас будет запуск аппаратной Phys-X на 4ххх или 5ххх, как уже один раз делали на 3870. Все пользователи АТИ на вас молиться будуь.
                  0
                  можно было бы, например, сравнить производительность драйверов OpenCL Nvidia и AMD. При примерно равных условиях: какой-нибудь простой пример из SDK запустить с одинаковыми global size и work-group size и посмотреть что быстрее… ну это из совсем тривиального…
                  Только сдается мне одновременно они не заработают — конфликтовать драйверы начнут друг с другом. Хотя это только догадки, конечно.
                    0
                    Стоят и не конфликтуют. Просто нужны дрвоа нвидии версии не выше 186, иначе при наличии АТИ в системе видеокарта переходит в состояние «оффлайн».
                      0
                      Не конфликтуют именно драйверы OpenCL? Ну то есть то, что карты разных производителей в одной системе работать могут это факт известный, а вот как там с использованием OpenCL при этом? Хотя по идее должны работать… ведь для того все и задумывалось)

                      Если будете проводить тесты — напишите пожалуйста, что из этого вышло — очень интересно)
                        0
                        Не конфликтуют, проблемы с n260gtx, капризная, с радеоном работает через раз.
                0
                Число регистров итп. для конкретной карты Nvidia можно посмотреть в Nvidia OpenCL Programming Guide (Appendix A) — на него есть ссылка в статье.
                Там написано какая карта имеет какую версию Compute Capability, там же и расшифровка, что в себя включает каждая версия Compute Capability.
            0
            Прочитал эту и предыдущие статьи, но так и не взял в толк вот чего…

            Как OpenCL соотносится с Java?

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

            OpenCL — это замена Java? Или что? Их можно вообще сравнивать?
              +1
              OpenCL и Java — разные инструменты для решения разных задач. OpenCL никогда не будет иметь такого набора библиотек.
                –5
                Спасибо, капитан.
                  +5
                  Человек постарался избавить вас от феерической каши в голове, а вы решили, что самый умный.
                    –3
                    Человек написал бессмысленный комментарий ни о чем. Я и так понимаю, что это разные инструмены. Вопрос в том, в чем же эта разница заключается.
                      0
                      Человек и написал что они настолько разные что указать различия невозможно, это как различия между подводной лодкой и букетом роз.
                        0
                        Я Вам легко напишу связь и различие между подводной лодкой и букетом роз. Помимо бессмысленного и ничего не объясняющего «это разные вещи» можно вполне внятно сказать, например, «на подводной лодке можно перевезти букет роз». Сразу становится понятным какие у них взаимоотношения и как одно может взаимодействовать с другим.
                          0
                          Тогда из этого утверждения у малознакомого с темой человека вырисуется определение: «Подводная лодка — средство для транпортировки цветов по морским каналам».
                0
                тут смысл не столько в общей среде выполнения, сколько в унифицированном доступе к ресурсам GPU
                  –1
                  А Java не имеет унифицированного доступа к ресурсам GPU?

                  Если имеет — то зачем OpenCL? Если не имеет — то можно ли в Java использовать OpenCL в качестве прослойки между Java и GPU?
                    0
                    вообще-то не имеет
                      0
                      тут также как и с OpenGL, насколько я помню можно
                        0
                        Есть Java обертка для OpenCL code.google.com/p/nativelibs4java/wiki/OpenCL
                          0
                          Java = Виртуальная машина, OpenCL = доступ к железу через среду разработки на достаточно низком уровне. Так понятно?
                            0
                            Java работает поверх OpenCL?
                              0
                              Java может использовать драйвер OpenCL.
                              Чуть выше давали ссылку на java-обертку для OpenCL API.
                              Насколько я понимаю, это не освобождает от написания kernel на С-подобном языке, а оборачивает только вызовы OpenCL API, описанные в статье.
                                0
                                Вот, вот теперь понятно. Спасибо.
                        +1
                        OpenCL — это компилятор плюс библиотека функций для общения с целевым устройством. Они разные для каждого устройства. О бинарниках, способных выполняться на любых устройствах, о виртуальной машине на GPU и даже о привычной работе с памятью речи (пока) не идет.
                        OpenCL можно сравнивать разве что с обычным C, но никак не с Java.
                          0
                          Получается, OpenCL — это просто такой драйвер к железке, который заставляет эту железку выполнять нецелевые функции?

                          А в чем тогда заключается стандарт, если все-равно для каждого устройства своя библиотека?
                            +1
                            Стандартизован язык, на котором можно писать собственно вычислительные функции, а также набор и поведение функций OpenCL API. Один и тот же исходник можно (теоретически, конечно) скомпилировать и запустить на любой платформе, для которой есть подходящий драйвер и компилятор.
                            А почему функции нецелевые-то?
                              0
                              Ну, подбор пароля к архиву на GPU — это определенно нецелевая функция.
                                0
                                Сегодня нецелевая, а завтра выйдет реклама «покупайте наши графические ускорители, чтобы быстрее ломать пароли!» и станет целевая ;).
                                  0
                                  Уже почти так и есть. Новая архитектура Nbidia Fermi (которая в железе пока в числе 7 уникальных чипов существует) была анонсирована как отлично подходящая под вычисления.

                                  Позже рекламщики Nvidia срочно бросились исправлять перегибы и рассказывать, что для 3D графики архитектура тоже очень подходит и покупать ее для игр тоже имеет смысл.
                                    0
                                    Ферми пока не существует на потоке, у руля (временно) ATI 5870, по бумажным характеристикам 300-ое поколение нвидии — монстры, очень хочется узнать, что же будет в деле…
                              0
                              OpenCL можно сравнить с CUDA только OpenCL для всех а Cuda только для Nvidia но там с совместимостью по круче…
                          0
                          все пора менять названия видеокартам, на другое.
                            0
                            Переименовать на nVidia Fermi?
                              0
                              Интересно, что они подразумевают под поддержкой С++ в этой архитектуре?
                              С++ kernel'ы в CUDA?
                                0
                                Боюсь, что да. Спасибо унифицированной памяти.
                                0
                                нет на, что-о типо йаСчиталко!
                              0
                              > ВНИМАНИЕ!!! oclLoadProgSource — не является функцией OpenCL API, а находятся в вспомогательной библиотеке, поставляемой вместе с Nvidia Computing SDK.

                              oclGetMaxFlopsDev — тоже оттуда
                                0
                                К сожалению по ссылке не удалось скачать NVidia OpenCL SDK. Login требует, кнопки «зарегистрироваться» не лицезрел.
                                • НЛО прилетело и опубликовало эту надпись здесь

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

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