OpenCL. Как начать

[ProPupil]

Для c# есть целых две библиотеки:Cloo и OpenClTemplate. Как раз в ваших приведённых ссылках есть ссылка на статью, посвящённую этим двум библиотекам, нос оглашусь с Вами, что на русском все-таки маловато, хоть нагло язычная документация и достаточно чёткая. По крайней мере, у меня не возникало проблем с её пониманием.

[rznELVIS]

тут для меня проблема по другому чуть выглядит. когда осваиваешь OpenCL через подобную Net-библиотеку, то иногда начинаешь сомневаться в чем проблема (при наличие ошибки) в библиотеки или в коде на OpenCL. а в в работе с OpenCL ошибки можно получить даже при корректном kernel. Пример некорректное значение global и local work-item size. И по-этом чтобы не бояться эффекта «сломанного телефона», то я и предлагаю получить подтверждение очевидного, поработав с OpenCL через родной API, а потом уж гоу и в .Net.

[HomoLuden]

И по-этом чтобы не бояться эффекта «сломанного телефона»

aka «Дырявая» или «Утекающая Абстракция»

[alexac]

Пожалуйста уберите try/catch вокруг чтения файла. Он там бесполезен. Вы используете только функции libc, которые не бросают исключения. Единственная ошибка, которую в этом коде еще можно было бы получить — ошибка выделения памяти, что в текущем виде все равно приведет к segmentation fault, а не к исключению. Тот кусок стоило написать как-то так:

const char fileName[] = "../forTest.cl";
size_t source_size;
std::unique_ptr<char[]> source;

try {
  std::ifstream stream;
  stream.exceptions(std::ios_base::bad_bit | std::ios_base::fail_bit | std::iOS_base::eof_bit);
  stream.open(fileName, std::ios_base::in);
  stream.seekg(0, std::ios_base::end);
  size_t source_size = stream.tellg();
  stream.seekg(0, std::ios_base::beg);
  source.reset(new char[source_size]);
  stream.read(source.get(), source_size);
} catch (const std::exception& e) {
  std::cerr << "Can't read kernel: " << e.what() << std::endl;
  exit(1);
}

[rznELVIS]

Учту. спасибо

[vilgeforce]

«Главное это проверить, что устройство функционирует нормально.» — ни разу не гарантия, что у вас будут нормально перечисляться устройства, как показывает практика.

clGetPlatformIDs у вас вызывается, но нет выделения памяти под результаты и освобождения ее потом. И такая же проблема с другими функциями.

[rznELVIS]

Тестировал на 5 компьютерах/ноутах. Если драйвера стоят актуальные то проблем не бывало. Слышал о возможных проблемах с устройствами, но сам не встречал. А хотелось писать о том что видел сам.

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

[rznELVIS]

Спасибо. приму к сведению. Работал только со старым OpenGL. По-этому и вспомнил его для хорошей метафоры. к тому же совсем недавно тыкал OpenGL через библиотеку Tao (студентам показывал), в шейдеры не лазил.

[stgatilov]

Я думаю, при освоении OpenCL лучше перейти на официальные C++-обёртки как можно быстрее. Использование C-шного API — лишняя трата ценного времени: тонны параметров, void*, обработка ошибок и пр.

P.S. Khronos как-то не так написан…

[rznELVIS]

Буду рад ссылочкам, на понравившиеся вам фреймворки.

Khronos поправил)) спасибо

[stgatilov]

Это не фреймворк, это именно обёртки над C API, официальные от Khronos.
Все команды в обёртках являются эквивалентами команд C-шного API, и наоборот, для каждой команды C-шного API есть эквивалент в обёртках. Параметры почти такие же, названия примерно одинаковые. То есть отдёльно «учить» C++ обёртки не нужно: зная C API, можно легко начать пользоваться обёртками.

Среди бонусов:

• всюду нормальная типизация вместо void * и размеров в байтах
• использование std::vector или cl::vector, cl::string вместо char*
• автоматическое отслеживание времени жизни объектов (типа умных указателей)
• автоматическая обработка ошибок через исключения (нужно включить)

Хедер
Спецификация

Пока есть только для OpenCL 1.2, новые версии запаздывают.

[BelBES]

А есть какие-то бенчмарки на тему того, насколько производительность отличается у кода нормально написанного через ANSI C API и при помощи оберток? Ведь скорей всего из оберток генерируется менее оптимальный код.

[stgatilov]

В спецификации (ссылка выше) об этом написано: «The C++ API corresponds closely to the underlying C API and introduces no additional execution overhead».

Вопрос про производительность вряд ли возникнет у человека, который использует C++ wrappers и мельком видел их внутренности.
Вариант API никак не влияет на код собственно вычислительных программ OpenCL, так что ядра работают абсолютно с одинаковой скоростью. Копирование данных запускает те же самые команды C-шного API, то есть работает столько же времени. Все понятия, такие как асинхронные вызовы например, одинаковы и в C, и в C++.
Чаще всего разница только в том, что вместо прямого вызова C-шной функции вызывается метод класса C++, который внутри вызывает ту же C-шную функцию. Скорее всего он инлайнится, а даже если и нет, overhead от лишнего вызова будет ничтожно мал по сравнению с временем работы самого вызова в драйвере. Ну может быть какой-то overhead есть из-за того, что используется std::vector<cl::device> и std::vector<std::string> вместо вручную выделенных кусков памяти. Но это опять же ничтожно малое время по сравнению с временем инициализации девайса и компиляции программы. Обработка ошибок тоже может отнимать время, но чтобы это стало заметным, надо написать невероятно неэффективную программу с огромным количеством API-вызовов.
По сути C++ обёртка — это просто более удобное использование C-шного API.

[RomanArzumanyan]

API написан на С неспроста. Как только начнёте оптимизировать производительность (использовать события и их списки, разбивать обьекты памяти для одновременного использования, и т. п.), наверняка вернётесь к низкоуровневому API. ИМХО, официальные С++ обёртки добавляют удобства только в момент инициализации.

[stgatilov]

Мне кажется, API написан на C в первую очередь потому, что его больше не на чем писать. C-шные функции легко использовать из любого языка программирования в любой ситуации. Если делать C++ API первостепенным, то начнётся куча проблем (радости с именованием символов, версии и настройки STL). К тому же, драйверы как-то принято писать на чистом C, а не на C++ со всякими std::vector и std::string=)

Конечно, мне пока что не были необходимы ни синхронизация по событиям, ни тем более разделение буфера на части. Однако мне кажется, что как минимум синхронизация по событиям в C++ обёртках есть, и нет проблем её использовать. В любом случае, такие проблемы уже совсем не для новичков.

К тому же, никто не мешает использовать C++ обёртки почти всюду, а в случае возникновения трудностей перейти на C API в одном конкретном проблемном месте. В этом ещё одно отличиё обёрток от фреймворка: обёртки и исходный API можно легко смешивать (с фреймворкам дело обычно сложнее).
Например, мне пришлось писать извлечение бинарного кода программы на C API из-за бага обёрток. Надо сказать, получилось 25 уродливых строчек с кучей звёздочек, sizeof, и самопальных макросов SAFE_CALL. Если бы не баг в обёртках (который, кстати, никто не мешает исправить самому локально), было бы строчек 2-5.

[RomanArzumanyan]

C-шные функции легко использовать из любого языка программирования в любой ситуации

С этим не поспоришь.

Пытался пару раз использовать С++ обёртку — мне не понравилось то, что она ничуть не более «интеллектуальная», чем оригинал: например, при чтении буффера нужно всё так же указавать размер, смещение и т. п. — как будто бы класс объекта памяти не должен инкапсулировать в себе эти данные. А если уж указываешь (почти) всё руками, то, ИМХО, невелика разница — делать это из С или из плюсов.

[stgatilov]

Да, C++ обёртка не решает волшебным образом все проблемы именно потому, что она почти полностью соответствует C-шному API.
Чтобы не мучиться постоянно, придётся пописать какие-то классы поверх неё (например типизированный буфер). Будет проще, чем поверх C-шного API.

[Salabar]

По-моему, базовый ОпенЦЛ это аналог JDBC в СУБД. Писать программы но можно, но в определенный момент есть риск немножко поехать, так что для начала следует написать себе небольшой фреймворк под задачу. То, что можно писать на один параметр благодаря плюсам это классно, но это не избавит от десятков clSetKernelArg подряд. С другой стороны, если в качестве будущего OpenCL разработки собираются позиционировать этот ад: www.khronos.org/registry/sycl/specs/sycl-1.2.pdf, то это даже не такое и зло.

[stgatilov]

но это не избавит от десятков clSetKernelArg подряд

Десяток clSetKernelArg подряд — это C API.
В C++ обёртках ядро принято вызывать примерно как обычную функцию.
Вот пример:

  //устанавливаем параметры запуска (кол-во потоков, очередь)
  auto launcher = kernel.bind(queue, AlignUp(paramsCnt, 64), 64); 
  //вызываем ядро, передавая параметры как в обычную функцию
  launcher(paramsBuffer, paramsCnt, resultsBuffer);

Здесь аргументы — объекты типа cl::Buffer и примитивные типы.
Хотя должен заметить, что обёртки писались под старый стандарт C++, поэтому количество аргументов ограничено 32 (хотя вроде не сложно это число увеличить). Бывало, в это ограничение я не укладывался =)

[Salabar]

Так и с С API никто не мешает функцию написать. Суть в том, что OpenCL сам по себе мало помогает определять зависимости между вычислениями. Низкоуровневый слишком. Это даже хорошо, можно даже пойти еще дальше и сделать что-то типа Vulkan с поддержкой OpenCL C, чтобы не было всяких нубских clCreateBuffer, но если я в курсаче на десяток ядер запутался, то как что-то серьезное на этом писать я не представляю. Я хочу сесть и подумать над либой, которая позволит составить граф вычислений декларативно, чтобы библиотека сама расставляла примитивы синхронизации, создавала промежуточные буферы, но, честно говоря, даже не представляю, с чего начинать.

[stgatilov]

Во-первых, можно для начала написать прототип с синхронными запусками ядер и копирований. В большинстве случаев будет прекрасно работать (если ядра делают много вычислений).
Потом можно попробовать что-то типа ленивых вычислений. Достаточно реализовать класс future-буфера, у которого внутри есть cl::Buffer (где будут данные) и cl::Event (когда данные в буфере будут правильными). Ну а дальше сделать обёртку над каждой командой, которая будет сама устанавливать списки ожидания для команд и записывать куда следует cl::Event готовности результата.
В целом кажется, что стоит покурить ключевые слова вроде future/promise/async.

P.S. Писать отдельное декларативное описание ядер и зависимостей между ними в каком-то особом формате — кажется излишним.

[stgatilov]

С другой стороны, если в качестве будущего OpenCL разработки собираются позиционировать этот ад: www.khronos.org/registry/sycl/specs/sycl-1.2.pdf, то это даже не такое и зло.

SYCL — замечательная задумка. Пока что она сильно молодая, чтобы всерьёз возлагать на неё надежды. Судя по всему, это будет платформенно-независимой версией CUDA runtime.
Должен заметить, что CUDA runtime на порядок удобнее в использовании, чем OpenCL, что затрудняет переход CUDA-поклонников на OpenCL (особенно учёных, которым чихать на кроссплатформенность).

Конечно, SYCL выглядит как-то далековато от простого OpenCL, всюду куча абстракций, на первый взгляд ненужных. Однако какие-то вещи оттуда существенно облегчают работу.

Например, я реализовал в нашем проекте «двойной» типизированный буфер, который существует одновременно на host и device, и синхронизируется автоматически прозрачно для пользователя. Естественно, пользователь должен корректно вызывать функции доступа к сторонам буфера, указывая каждый раз: нужен доступ «только на чтение» или «на чтение и запись».
Чуть позже я узнал о SYCL. В нём sycl::buffer является ровно тем же самым общим типизированным буфером, и для работы с ним надо ровно так же вызывать функцию с указанием типа доступа.

Также в SYCL есть более простой способ использования асинхронной работы.
Короче говоря, лично я надеюсь, что в будущем смогу писать на SYCL =)

[igor_suhorukov]

А как происходит отладка OpenCL kernel функций и C#+host api?

[Salabar]

На некоторых платформах можно из ядер вызывать printf. Для всего остального придется сначала отладить их как обычные Си-функции и надеяться, что багов с многопоточностью там не осталось.

[igor_suhorukov]

Надежда — не лучшее решение в разработке ПО)

[RomanArzumanyan]

Если вы пишете под AMD, то можно отлаживать код ядер из AMD CodeXL. Intel тоже в каком-то виде поддерживает отладку внутри ядра. В остальных случаях — уповать на наличие printf и/или сливать результаты вычислений в отладочные объекты памяти и делать postmortem.

[igor_suhorukov]

Когда-то давно писал как можно отлаживать java + OpenCL

[ALEX_k_s]

А в новых версиях OpenCL еще не появилась поддержка template C++?

[Salabar]

В 2.1 появится возможность писать ядра на подмножестве С++17 (и вообще на чем угодно, благодаря компиляции в промежуточный код SPIR-V), если вы про это.

[tapin13]

Спасибо за статью, жаль что так и нет продолжения.

Не увидел ссылки на работающий код. Так как написал свой (по статье) и выложил на гит, решил поделиться ссылкой github.com/tapin13/openCL-helloWorld