визуализировать используя opencl что это
Использование изображений, доступных для чтения и записи, в OpenCL 2.0
До OpenCL 2.0 было невозможно проводить операции чтения и записи изображения в рамках одного и того же ядра. Можно было объявлять изображения как CL_MEM_READ_WRITE, но после передачи изображения ядру приходилось выбирать одно из двух: либо __read_only (доступ только для чтения), либо __write_only (доступ только для записи). В OpenCL 2.0 появилась возможность читать и записывать изображения в пределах одного ядра. Однако, имеется несколько особенностей, о которых мы подробно поговорим в этом посте.
Фрагмент кода 1. Можно было создать буфер изображения с помощью CL_MEM_READ_WRITE
Фрагмент кода 2. В OpenCL 2.0 появилась возможность читать и записывать изображения в пределах одного ядра
Преимущества изображений, доступных для чтения и записи
Свертка изображений менее эффективна при использовании новой функциональности чтения и записи изображений, но другие алгоритмы обработки изображений могут получить значительные преимущества за счет этой функциональности. Одним из примеров таких процессов может быть составление изображений.
В OpenCL 1.2 и более ранних версиях изображения могли иметь только квалификаторы __read_only и __write_only. В OpenCL 2.0 появился квалификатор __read_write, выходные данные могут быть скопированы во входной буфер. Это позволяет снизить количество необходимых ресурсов. Для каких-либо изменений изображения необходимо обрабатывать изображение как буфер и работать с этим буфером (см. cl_khr_image2d_from_buffer).
Текущее решение состоит в том, чтобы обрабатывать изображения как буферы и управлять буферами. Для обработки двухмерных изображений как буферов требуется затратить определенное количество ресурсов. При этом также становится невозможно использовать возможности срезания и фильтрации, доступные в read_images. Поэтому желательно использовать изображения с квалификатором read_write.
Обзор примера
В примере два растровых изображения (input1.bmp и input2.bmp) помещаются в буфер. Затем эти изображения накладываются одно на другое на основе значения альфа (это фактор веса в уравнении вычисления пикселей). Значение альфа передается в виде параметра.
Рисунок 1. Альфа = 0,84089642
Входные изображения должны быть 24- или 32-разрядными. На выходе получается 24-разрядное изображение. Входные изображения должны быть одинакового размера. Изображения были в формате ARGB, это учитывалось при их загрузке.
Рисунок 2. Альфа = 0,32453
Значение бета равно разности между единицей и значением альфа.
Эти два значения определяют «вес» изображений 1 и 2 в выходном изображении.
Для изменения яркости каждого пикселя можно использовать значение гамма. По умолчанию это значение равно нулю. Пользователь может изменить яркость готового изображения целиком.
Пример запуска программы
Рисунок 3. Запуск программы на устройстве OpenCL 2.0
Ограничения изображений, доступных для чтения и записи
Ограничители нельзя использовать с изображениями, для которых требуется синхронизация между разными рабочими группами. Для свертки изображений требуется синхронизация всех потоков. Свертка по отношению к изображениям обычно предусматривает математические действия над двумя матрицами и создание в результате третьей матрицы. В примере свертки изображений используется размытие Гаусса. В других примерах используется повышение резкости изображений, обнаружение краев и придание рельефности.
В качестве примера рассмотрим размытие Гаусса. Фильтр Гаусса — это низкопроходный фильтр, удаляющий высокочастотные значения. В результате снижается уровень детализации изображения и получается эффект размытия. Применение размытия Гаусса — то же самое, что преобразование изображения с помощью функции гауссова распределения (ее часто называют маской). Для демонстрации функциональности чтения и записи изображений пришлось применить размытие по горизонтали и по вертикали.
В OpenCL 1.2 это пришлось бы делать в два прохода. Одно ядро использовалось бы только для размытия по горизонтали, а другое — для размытия по вертикали. Результат одного размытия будет использоваться в качестве входных данных для следующего (в зависимости от того, какое размытие было первым).
Фрагмент кода 3. Ядро размытия Гаусса в OpenCL 1.2
В OpenCL 2.0 эти два ядра можно объединить в одно. Используйте ограничитель, чтобы принудительно завершать размытие по горизонтали или по вертикали перед началом следующего размытия.
Фрагмент кода 4. Ядро размытия Гаусса в OpenCL 2.0
Оказалось, что ограничители неэффективны. Использование ограничителей не гарантирует, что размытие по горизонтали будет выполнено до начала размытия по вертикали (если первым было размытие по горизонтали). В результате при нескольких запусках получались разные результаты. Ограничители можно использовать для синхронизации потоков в группе. Причина проблемы состоит в том, что происходит чтение краевых пикселей из нескольких рабочих групп, а способа синхронизации между несколькими рабочими группами нет. Первоначальное предположение о возможности реализации единого размытия Гаусса с помощью чтения и записи изображений оказалось неверным, поскольку в OpenCL невозможна синхронизация зависимостей данных между рабочими группами.
Вычисляем на видеокартах. Технология OpenCL. Часть 1b. Пишем для OpenCL
22 июня автор курса «Разработчик C++» в Яндекс.Практикуме Георгий Осипов провёл вебинар «Вычисляем на видеокартах. Технология OpenCL».
После перерыва продолжаем публикацию текстовой версии вебинара.
У программы для OpenCL есть две части: kernel-код и host-код — то, что выполняется на видеокарте, и то, что выполняется на компьютере. Кроме того, программу нужно скомпилировать и запустить. Всё это будет рассмотрено в сегодняшней статье. Начнём с самого интересного — напишем часть kernel.
В предыдущих сериях
Kernel
Kernel-код для OpenCL пишется на C с ограничениями. В нём нельзя использовать:
Но не будем отвлекаться и вернёмся к написанию kernel. Мы будем писать функцию, рисующую множество Мандельброта.
Пишем точку входа — kernel-функцию. Она всегда void и будет вызвана для каждого work-item’а
Зелёным выделено название kernel, синим — встроенные функции OpenCL, красным — функции, которые мы скоро напишем.
Это двумерная задача. И я хочу, чтобы один work-item рисовал какую-то точку множества Мандельброта. Каждый work-item выполняет код функции независимо от других, и ему прежде всего нужно понять, какую точку множества Мандельброта отрисовать. Чтобы это как-то кастомизировать, в аргументах задаём точку плоскости, которая находится в центре нашей картинки, и масштаб. Разберём аргументы подробнее:
w и h используются в условии основного if : если work-item оказался в добавочных пикселях справа или снизу, то он будет отдыхать. Конечно, из-за дивергенции он, скорее всего, будет имитировать те же действия, что и работающие work-item’ы, но на производительности это не скажется.
Вычисляем количество итераций для каждой точки экрана
Вычисляем цвет по количеству итераций
Вот такой нехитрый kernel у нас получился. Переходим к части host.
Теперь kernel-код надо запустить на GPU. План такой:
Программа, использующая OpenCL, выглядит так:
Ошибки в этом примере обрабатывать не будем
Чтобы main не разрастался, напишем отдельную функцию для выбора устройства.
В этом примере не будем умничать и возьмём первое же устройство. Вначале попробуем найти GPU. Если не получилось — ищем CPU. Да, OpenCL может выполняться и на обычном процессоре, если установлен соответствующий драйвер. В реальной программе можно получить список всех устройств и предоставить возможность выбора пользователю. Проверять железо на поддержку не нужно, потому что библиотека OpenCL выдаёт только те девайсы, на которых она будет работать.
В этом примере синим цветом выделены функции из API OpenCL. Здесь также использованы три типа данных OpenCL:
Теперь всё готово к тому, чтобы переходить к запуску kernel-кода. Но для начала нужно его скомпилировать. Сделаем это такой функцией:
Скомпилированный код программы можно сохранить и переиспользовать
Функция сперва загружает исходный код, а затем выполняет компиляцию. Это простой пример, и в нём не затронуты следующие возможности, необходимые в реальной программе:
Код можно прочитать из файла или встроить в программу как ресурс
Вызовом clSetKernelArg задаются все аргументы, которые есть у kernel-функции draw_mandelbrot. Они сохраняются внутри объекта cl_kernel. Особо можно отметить, что clSetKernelArg — это единственная непотокобезопасная функция OpenCL. Все остальные вызовы можно делать из разных тредов вашей программы без каких-либо синхронизаций. Разумеется, если вы уверены в целостности своих данных.
Полный код задания аргументов:
Не хватает только функции align :
Чтобы глобальный размер был кратен размеру рабочей группы, воспользовались такой функцией
Подберём параметры и запустим функцию, читая результат в вектор:
Осталось совсем немного — сохранить результат. Для этого выберем самый простой формат изображений — PPM. Вот функция для сохранения в него:
И не забываем убирать за собой! В программировании нужно освобождать те ресурсы, которые вы уже использовали. А это ни много ни мало kernel-функция, буфер, очередь, вся программа и контекст.
Чтобы не забывать освобождать ресурсы, можно написать обёртку на C++. Или воспользоваться готовой
Сборка и запуск
Итак, программа написана. Но сохраним интригу и прежде чем показать результат её работы, разберём сборку и запуск. Многие программисты на C++ сталкивались с проблемами сборки: то что-то не компилируется, то возникают конфликты, то странные ошибки. С OpenCL в этом плане не всё так плохо. Главная причина: библиотека OpenCL — динамическая. Так задумано её использование. Для обращения к ней нужен лишь несложный интерфейс, который умещается в нескольких файлах библиотеки clew. Clew — это маленькая библиотечка, позволяющая делать замечательные вещи. Она легко собирается, сама ищет динамическую библиотеку OpenCL и содержит все необходимые include-файлы.
Библиотека clew — не единственный поставщик cl.h. Альтернатива — тяжеловесный инструмент CUDA Toolkit или готовый пакет на системах с пакетным менеджером. При компоновке нужно добавить соответствующую библиотеку. Clew состоит всего лишь из одного компилируемого файла, и можно просто добавить его в проект.
Для запуска программы нужно, чтобы она увидела динамическую библиотеку OpenCL, а библиотека OpenCL должна увидеть видеокарту. На Windows для работы с видеокартой достаточно драйвера. Для других систем или для использования CPU, вероятно, придётся устанавливать специальные драйверы. Вот ссылка на драйвер для интеловских процессоров.
Как правило, самый универсальный способ установить всё и сразу — установить CUDA Toolkit. В него входит cl.h, динамическая библиотека OpenCL и драйвер карты Nvidia.
Ну что ж, программа готова, можно запускать!
Изображение множества Мандельброта, полученное написанной программой
Итоги
Пример сделал красивую картинку, но показал не всё. В частности, мы не рассмотрели:
Напоследок я приготовил схему, которая напомнит, как по шагам написать host-код для OpenCL. Думаю, начинающим она будет полезна.
В этой статье мы подробно разобрали написание простейшей программы для OpenCL. В следующей части рассмотрим алгоритмы на GPU и напишем более сложный kernel-код.
Полный код примера опубликован в репозитории на GitHub.
OpenCL под C# это просто
Хотя технология OpenCL появилась ещё в 2008 году, большого распространения она не получила до сих пор. Плюсы технологии несомненны: ускорение вычислений, кроссплатформенность, способность исполнять код как под GPU, так и под CPU, поддержка стандарта целым рядом компаний: Apple, AMD, Intel, nVidia и некоторыми другими. Минусов не так много, но и они есть: более медленная работа на nVidia, чем через CUDA, сложность использования. Первый из минусов влияет только при серьёзной разработке, где скорость программы важнее кроссплатформенности. Второй и является основным препятствием на пути разработчиков, делающих выбор в пользу того или иного метода разработки. Чтобы разобраться в куче хэдэров, драйверов и стандартов требуется куча времени. Но не всё так плохо. Темой этой статьи будет короткий guide по тому, как наиболее простым способом можно запустить OpenCL под C# и получить удовольствие от параллельного программирования.
Настройка драйверов
nVidia.
Intel.
В отличие от AMD и nVidia проблем с драйверами от Intel у меня не возникло ни разу. Взять их можно тут.
Установка драйверов на систему ещё не гарантирует что они будут у вас работать. Количество багов у AMD зашкаливает (проблемы возникали при установке на 2 из 3х компьютеров), у nVidia оно велико (1 из 3х). Поэтому, после установки рекомендую сначала проверить, подключилось ли OpenCL. Наиболее просто это можно сделать через программы показывающие параметры видеокарт. Я пользуюсь GpuCapsViewer, так же работают opencl-z и GPU-Z.
Если проблемы возникли… Удалите все старые драйвера, переставьте. Для AMD — убедитесь что ставите правильную версию драйверов. Драйвера для ноутбуков у них часто глючные. Если проблемы не исчезли — переустановка винды вас спасёт.
Врапперы
Так как нашей целью является максимально простое программирование на OpenCL под C#, мы не будем заниматься извращениями и подключать OpenCL хэдэры, а воспользуемся уже готовыми библиотеками, упрощающими разработку. Самой полной и безглючной версией, как мне кажется, на сегодняшний день является cloo.dll, входящая в OpenTK. Самой простой в применении, автоматизирующей множество операций является OpenCLTemplate, являющаяся надстройкой над cloo. Из минусов последней — некоторые глюки при работе с AMD, например с последней версией драйверов (11.6) могут отказаться инициализироваться устройства. Так как проект OpenSource, глюки которые у меня были я нашёл и поправил, но когда зарелизят новую версию библиотеки — не знаю. Так же есть несколько менее известных врапперов, которые можно найти на просторах интернета.
Первая программа
В качестве первой программы посчитаем через OpenCL сумму двух векторов, v1 и v2. Пример программы написанной с использованием cloo.dll:
Программа на Cloo
//Текст програмы, исполняющейся на устройстве (GPU или CPU). Именно эта программа будет выполнять паралельные
//вычисления и будет складывать вектора. Программа написанна на языке, основанном на C99 специально под OpenCL.
string vecSum = @»
__kernel void
floatVectorSum(__global float * v1,
__global float * v2)
<
int i = get_global_id(0);
v1[i] = v1[i] + v2[i];
>
» ;
//Список устройств, для которых мы будем компилировать написанную в vecSum программу
List Devs = new List ();
Devs.Add(ComputePlatform.Platforms[1].Devices[0]);
Devs.Add(ComputePlatform.Platforms[1].Devices[1]);
Devs.Add(ComputePlatform.Platforms[1].Devices[2]);
//Компиляция программы из vecSum
ComputeProgram prog = null ;
try
<
//Инициализация новой программы
ComputeKernel kernelVecSum = prog.CreateKernel( «floatVectorSum» );
А ниже та же программа, реализованная через OpenCLTemplate.DLL
private void btnOpenCL_Click( object sender, EventArgs e)
<
//Текст програмы, исполняющейся на устройстве (GPU или CPU). Именно эта программа будет выполнять паралельные
//вычисления и будет складывать вектора. Программа написанна на языке, основанном на C99 специально под OpenCL.
string vecSum = @»
__kernel void
floatVectorSum(__global float * v1,
__global float * v2)
<
int i = get_global_id(0);
v1[i] = v1[i] * v2[i];
>» ;
//Инициализация платформы. В скобках можно задавать параметры. По умолчанию инициализируются только GPU.
//OpenCLTemplate.CLCalc.InitCL(Cloo.ComputeDeviceTypes.All) позволяет инициализировать не только
//GPU но и CPU.
OpenCLTemplate.CLCalc.InitCL();
//Команда выдаёт список проинициализированных устройств.
List L = OpenCLTemplate.CLCalc.CLDevices;
//Команда устанавливает устройство с которым будет вестись работа
OpenCLTemplate.CLCalc.Program.DefaultCQ = 0;
//Компиляция программы vecSum
OpenCLTemplate.CLCalc.Program.Compile( new string [] < vecSum >);
//Присовоение названия скомпилированной программе, её загрузка.
OpenCLTemplate.CLCalc.Program.Kernel VectorSum = new OpenCLTemplate.CLCalc.Program.Kernel( «floatVectorSum» );
int n = 100;
float [] v1 = new float [n], v2 = new float [n], v3 = new float [n];
//Инициализация и присвоение векторов, которые мы будем складывать.
for ( int i = 0; i //Загружаем вектора в память устройства
OpenCLTemplate.CLCalc.Program.Variable varV1 = new OpenCLTemplate.CLCalc.Program.Variable(v1);
OpenCLTemplate.CLCalc.Program.Variable varV2 = new OpenCLTemplate.CLCalc.Program.Variable(v2);
//Объявление того, кто из векторов кем является
OpenCLTemplate.CLCalc.Program.Variable[] args = new OpenCLTemplate.CLCalc.Program.Variable[] < varV1, varV2 >;
//Сколько потоков будет запущенно
int [] workers = new int [1] < n >;
//Исполняем ядро VectorSum с аргументами args и колличеством потоков workers
VectorSum.Execute(args, workers);
//выгружаем из памяти
varV1.ReadFromDeviceTo(v3);
Как видно, второй вариант куда проще и интуитивнее.
Ссылки напоследок
В первую очередь программированию на OpenCL через С# посвящён этот сайтик — www.cmsoft.com.br. К сожалению, его ведёт мало народу, поэтому примеры часто неадекватные, а OpenCLTemplate, созданный авторами сайта весьма глючный.
Полезным местом является сайт www.opentk.com где весьма оперативно отвечают на вопросы о программирование через cloo.dll
Примерно те же ответы можно получить и на сайте sourceforge.net/projects/cloo
Стандарт программирования OpenCl, основанный на C99 описан здесь — www.khronos.org/opencl
Продолжение статьи «Введение в OpenCl» рассказывает об особенностях языка программирования, которым мы программируем видеокарту.
OpenCL в повседневных задачах
Недавно мы рассказывали про HSA и в ходе обсуждения преимуществ нового подхода к построению ПК затронули такую интересную тему, как GPGPU — вычисления общего назначения на графическом ускорителе. Сегодня видеоускорители AMD предоставляют доступ к своим ресурсам с помощью OpenCL — фреймворка, обеспечивающего сравнительно простое и понятное программированое высокопараллельной системы.
Сегодня технологии OpenCL поддерживаются всеми основными игроками на рынке: возможность предоставить программам доступ к «продвинутому» ускорению (к тому же бесплатная, т.к. OpenCL не подразумевает каких-либо отчислений и роялти) явно того стоит, а от универсальности таких API выигрывают все, кто реализует поддержку OpenCL в своих продуктах.
Подробнее о том, где сегодня можно встретить OpenCL в повседневной жизни, как он ускоряет обычный офисный софт и какие возможности открывает разработчикам сегодня и поговорим.
GPGPU, OpenCL и немного истории
Само собой, OpenCL — не единственный способ реализовывать общие вычисления на GPU. Помимо OpenCL на рынке присутствуют CUDA, OpenACC и C++AMP, но по-настоящему популярными и находящимися на слуху являются первые две технологии.
Разработкой стандарта OpenCL занимались те же люди, которые подарили миру технологии OpenGL и OpenAM: Khronos Group. Сама торговая марка OpenCL принадлежит компании Apple, но, к счастью для программистов и пользователей по всему миру, данная технология не является закрытой или привязанной к продукции «яблочной» компании. Помимо Apple в Khronos Goup входят такие гиганты рынка, как Activision Blizzard, AMD, IBM, Intel, NVidia и ещё с десяток компаний (в основном, производителей ARM-решений), которые присоединились к консорциуму позже.
В определённой мере OpenCL и CUDA идеологически и синтаксически схожи, от чего сообщество только выиграло. Программистам (в силу схожести определённых методов и подходов) проще использовать обе технологии, переходить от «закрытой» и привязанной к железу NVidia CUDA к универсальному и работающему везде (в том числе и на обычных многоядерных CPU, и на суперкомпьютерах на базе архитектуры CELL) OpenCL.
OpenCL в повседневном использовании
Сейчас вы подумаете «ага, будут говорить про игры и фотошоп». Нет, OpenCL способен ускорить вычисления не только связанные с графикой. Одним из самых популярных приложений, использующих возможности GPGPU является… кроссплатформенный офисный пакет LibreOffice. Поддержка OpenCL появилась в нём в 2014 году и применяется для ускорения расчётов в табличном менеджере Calc.
Вот наглядное видеосравнение производительности системы с AMD A10-7850K с графическим ядром R7 и Intel Core I5 с HD4600 на борту:
В синтетических тестах тот же AMD A10-7850K по гетерогенным вычислениям с применением OpenCL обгоняет почти вдвое более дорогой i5-4670K / 4690:
К слову, в науке и её прикладных направлениях есть масса задач, которые отлично перекладываются на векторные процессоры видеоускорителей и позволяют выполнять расчёты в десятки и сотни раз быстрее, чем на CPU.
Например, различные разделы линейной алгебры. Умножение векторов и матриц — это то, чем GPU занимаются каждый день, работая с графикой. В этих задачах им практически нет равных, т.к. их архитектура затачивалась под решение таких задач годами.
Сюда же относятся и быстрое преобразование Фурье, и всё, что с ним связано: решение сложных дифференциальных уравнений различными методами. Отдельно можно выделить гравитационные задачи N-тел которые применяются для расчёта аэро- и гидродинамики, моделировании жидкостей и плазмы. Сложность расчётов заключается в том, что каждая частица взаимодействует с другими, законы взаимодействия достаточно сложны, а вычисления требуется проводить параллельно. Для таких задач OpenCL и возможности GPU AMD подходят как нельзя лучше, т.к. параллельные вычисления с множеством объектов и так успешно решаются на процессорах такого типа каждый день: в пиксельный шейдерах.
Структурированные решётки часто применяются в растровой графике. Неструктурированные — в вычислениях в области гидродинамики и при различных вычислениях с элементами, чьи графы имеют разный вес. Отличия структурированных решёток от неструктурированных в количестве «соседей» каждого элемента: у структурированных оно равное, у неструктурированных — разное, но и те и те отлично ложатся на возможности OpenCL по ускорению вычислений. Сложности по переносу вычислений, в основном, математические. То есть основная задача у программиста — не только «написать» работу системы, но и разработать математическое описание, которое переложит данные на аппаратные возможности с помощью OpenCL.
Комбинаторная логика (сюда же относится и вычисление хэшей), методы Монте-Карло — то, что отлично переносится на GPU. Множество вычислительных модулей, высокая производительность в параллельных вычислениях — то, что реально ускоряет эти алгоритмы.
Что ещё можно ускорить с помощью OpenCL и мощных GPU?
Поиск обратного пути. Вычисления графов и динамическое программирование: сортировки, обнаружение коллизий (соприкосновений, пересечений), генерацию регулярных структур, различные алгоритмы выборки и поиска. С некоторыми ограничениями, но поддаются оптимизации и ускорению работы нейронных сетей и связанных с ними структур, но здесь, скорее, проблемы в том, что нейронные структуры просто «дорого» виртуализировать, выгоднее использовать FPGA-решения. Отлично показывают себя работы конечных автоматов (которые и так применяются в работе с GPU, например, когда речь идёт о компрессии / декомпресии видеосигнала или работе по поиску повторяющихся элементов).
OpenCL vs CUDA
Сравнивать напрямую производительность OpenCL и CUDA практически не имеет смысла. Во-первых, если мы будем сравнивать их на видеокартах AMD и NVidia, то в грубой гонке вычислительных возможностей победят видеоадаптеры AMD: современные ускорители NVidia имеют ряд ограничений по производительности в формате FP64, внедрённые самой NVidia для того, чтобы продавать «профессиональные» видеокарты для вычислений (серии Tesla и Titan Z). Их цена несоизмеримо выше, чем у аналогов по FLOPS на базе решений AMD и их «родственных» карт в номерной линейке NVidia, что делает сравнение достаточно сложным. Можно учитывать производительность-на-Ватт или производительность-на-доллар, но к чистому сравнению вычислительной мощности это не имеет практически никакого отношения: «FLOPS’ы любой ценой» слабо вяжутся с текущей финансовой обстановкой, а по производительности-на-доллар «старушка» 7970 GHz Edition (она же R9 280X) до сих пор является одной из самых выгодных видеокарт.
Во-вторых, можно попробовать сравнить OpenCL и CUDA на видеокарте от NVidia, но сама NVidia реализует OpenCL через CUDA на уровне драйвера, так что сравнение будет несколько нечестным по вполне понятным причинам.
С другой стороны, если брать во внимание не только производительность, то кое-какой анализ провести всё же можно.
OpenCL работает на куда большем списке железа, чем NVidia CUDA. Практически все CPU, поддерживающие набор инструкций SSE 3, видеоускорители начиная с Radeon HD5xxx и NVidia GT8600 заканчивая новейшими Fury / Fury-X и 980Ti / Titan X, APU от AMD, встроенная графика Intel — в общем, практически любое современное железо с несколькими ядрами может воспользоваться преимуществанми данной технологии.
Особенности реализации CUDA и OpenCL (а также достаточно сложная документация, т.к. параллельное программирование в целом далеко не самая лёгкая область разработки), скорее, отражены в специфических возможностях и инструментах разработки, а не в области производительности.
Например, у OpenCL имеются некоторые проблемы с распределением памяти в силу «The OpenCL documentation is very unclear here».
Вместе с тем, CUDA уступает OpenCL в области синхронизации потоков — данных, инструкций, памяти, чего угодно. К тому же с помощью OpenCL можно использовать внеочередное исполнение потоков (out-of-order queues) и инструкций, а CUDA до сих пор умеет только in-order. На практике это позволяет избежать простоев процессора в ожидании данных, и эффект тем заметнее, чем длиннее ковейер процессора и больше разница между скоростью работы памяти и скоростью работы вычислительных модулей. В двух словах: чем больше мощности вы выделите под OpenCL, тем больше будет отрыв в производительности. CUDA для достижения сравнимых результатов потребует написание куда более сложного кода.
Инструменты для разработки (Дебаггер, профилер, компилятор) CUDA несколько лучше, чем аналогичные у OpenCL, но CUDA реализует API через язык C, а OpenCL — через С++, упрощая работу с объектно-ориентированным программированием, при этом оба фреймворка изобилуют «локальными» хитростями, ограничениями и особенностями.
Подход NVidia в данном случае сильно напоминает то, как работает компания Apple. Закрытое решение, с большим набором ограничений и строгими правилами, но хорошо заточенное для работы на конкретном железе.
OpenCL предлагает более гибкие инструменты и возможности, но требует более высокого уровня подготовки от разработчиков. Общий код на чистом OpenCL должен запуститься на любом поддерживающем его железе, вместе с тем «оптимизированный» под конкретные решения (скажем, видеоускорители AMD или процессоры CELL) будет работать заметно быстрее.