что новенького: согласно кроносу в версии 2.0 и 2.1 можно юзать фичи C++но уже в 1.2 на AMD можно делать так (на всякий случай дублирую пикчей)[code lang="python"] import numpy as npimport pyopencl as clkernsrc = """class Test{ public: void setX (int value); private: int x; private: int y;};void Test::setX (int value){ x = value;}__kernel void foo (__global Test* InClass){ if (get_global_id(0) == 0) InClass->setX(5);}"""mf = cl.mem_flagsctx = cl.create_some_context()queue = cl.CommandQueue(ctx)arr = np.array([1,2], dtype=np.int32)res = np.zeros(2, dtype=np.int32)clarr = cl.Buffer(ctx, mf.READ_WRITE| mf.COPY_HOST_PTR, hostbuf=arr)prg = cl.Program(ctx, kernsrc).build(options=" -x clc++ ")prg.foo(queue, (1,), None, clarr)cl.enqueue_copy(queue, res, clarr)print("Passed:", arr)print("Got modified:", res)[/code]и оно даже работает. Может пригодиться при портировании плюсового хлама, например.Список открытого софта с поддержкой OpenCL, которым пользуюсь я:Нейросетевой фреймворк caffe - пока отдельный бранч, но уже вполне рабочий и ускоренныйПакет 3d моделирования blender - с версии 2.76 (или даже раньше) рендер cycles работает поверх OpenCLРедактор raw фото darktable - многие фильтры идут с поддержкой OpenCL, есть возможность гибко настраивать через конфиг, например разные устройства для превью и полной отрисовки
Да, приветствуется обмен опытом, просьбы о помосчи у опытных товарищей, и так далее.
Лучше бы сделал туториал для ньюфагов с хелловорлдом.
OpenCL де факто мертв на данный момент. Нвидиа подмяла рынок ГПУ вычислений полностью.
>>796811Чем кудахтахтах лучше?
>>796812Спонсорством за спиной.
>>796818И все? А для обычного рядового погромана?
>>796793Можно начать с классики писалось ещё при Стиве:https://developer.apple.com/library/prerelease/content/samplecode/OpenCL_Hello_World_Example/Listings/hello_c.htmlА можно с pyopencl - легче не придумать. Простенький примерчик, как создать на хосте массив структур и обработать.[code lang="python"] import numpy as npimport pyopencl as clkernsrc = """ // Тут комментарии в цэ-стиле// Декларируем простенькую структуруstruct Test{ int x; int y;};__kernel void foo (__global struct Test* InStruct){ size_t gid = get_global_id(0); //Глобальный индекс int dx = gid; // Локальный индекс (внутри фронта) треды внутри одного фронта // могут работать с общей памятью с префиксом __local, с __global такой фокус не пройдёт // - записать можно что угодно, но прочитать записанное, в отличие от __local, не получится. int dy = get_local_id(0); // Изменяем переданную структуру InStruct[gid].x += dx; InStruct[gid].y += dy;}"""mf = cl.mem_flagsctx = cl.create_some_context() # Создаём дефольный контекстqueue = cl.CommandQueue(ctx) # Очередьarr = np.zeros(16, dtype=np.int32).reshape(-1, 2) # Сперва готовим данные на хостеres = np.zeros(arr.shape, dtype=np.int32) # Пустой массив для записи результатаclarr = cl.Buffer(ctx, mf.READ_WRITE| mf.COPY_HOST_PTR, hostbuf=arr) # Буфер на устройстве, заполненный данными хостаprg = cl.Program(ctx, kernsrc).build() # Программаprg.foo(queue, (arr.shape[0],), (4,), clarr) # Вызываем наше ядро (оно может быть и не одно)cl.enqueue_copy(queue, res, clarr) # Получаем взад отправленные данные# Сравниваемprint("Passed:\n", arr)print("Got modified:\n", res)[/code]у кого-нибудь работает эта ебучая разметка?
>>796811Да, есть такое. И раньше начала и учёных спонсировала и сё такое прочее. Но:>>796820OpenCL таки лучше. Поддерживается на: интелах начиная с хасфейла. голимом CPU на амуде и интел и, как минимум посредством pocl, на остальных - не-OpenCL версию кода можно вообще не писать, заработает везде. на армах - на всех новых SOC оно уже есть на FPGA и прочих dsp* а скоро и в вебе - WebCL маячит где-то в конце тоннеля.Короче - куда имеет куда меньший спектр возможных применений и добавляет анальную зависимость от вендора.
>>796834ебучая разметка. Забыл уточнить - поддержка на интелах именно на GPU, на CPU она у них давно есть.
>>796834>голимом CPU на амуде и интелНахуя он вообще нужен на ЦПУ, если и так можно вычисление распараллелить на потоки?Интересуют то видюшечки.
>>796826Это типо инлайн си в пайтоне? Компилится на лету?
>>796841Ты уверен, что оно компайлится, а не отправляется текстом на видюху, и компайлится уже там?
>>796839Затем, что видюшечки есть не везде и не у всех, а коду как-то работать надо. В случае с кудой придётся соснуть или писать CPU-версию алгоритма, в случае с OpenCL заработает и так.
>>796841Компилится драйвером OpenCL и скомпиленное летит на видяху
>>796844Но ГПУ же от ЦПУ пиздец дохуя отличается, написать одинаковый алгоритм чтобы работало похуй где - один хуй же не получится, не?
>>796826А что этот пример у тебя делает? Ты создал пустой массив, обработал его ядром? Че за ядро? Градиент штоле делает?
>>796778 (OP)ОП, можешь пояснить за SyCL? Его уже можно юзать?Не люблю си-подобную лапшу, люблю лаконичное дженерик программирование, чтобы ехали шаблоны через шаблоны, и ими лямбды погоняли.
>>796847> Че за ядро?Функция с пометкой __kernel> Градиент штоле делает?Просто пихает значения глобального и локального индексов в пустой массив.>>796846OpenCL код очень легко векторизуется, CPU реализации размазывают его по процессорам с использованием всяких AVX/SSE . Например, в 512-битном регистре AVX 16 флоатов * 4 ядра получится = 64 потока вычислений внутри 4х софтовых тредов.
>>796852>SyCLПонятия не имею, я работаю через питон, в-основним, на сишечке только примеры для багтрекеорв хуячу. Посмотри в сторону clBlas/ViennaCL - там ускоренная работа с матрицами. Если твои потребности ограничиваются простыми операциями над большими массивами данных, должно хватить.
>>796846Ещё поясню: OpenCL юзает параллелизм данных, условно говоря, у тебя есть массив на 1000000 значений и тебе нужно получить их синусы. Далее всё просто: задаёшь global_size в 1000000, внутри __kernel получаешь индекс текущего global_id и по нему работаешь с нужным элементом массива. Сколько у юзера потоков исполнения тебя не волнует - всё отработает само.
Помню, раньше в плане GPGPU карты от AMD били карты от NVIDIA по всем фронтам. Это правда? Даже если в AMD-карте реализовано на OpenCL, а в NVIDIA - на CUDA? Как сейчас с этим?
CUDA просто работает. У нее есть компилятор nvcc, который в один проход даст тебе и CPU и GPU код, подсветит тебе ошибки и выдаст бинарник, и все это будет работе на куче железа Nvidia. При этом уже лет 5 как поддерживается С++98 код (шаблоны! перегрузка операторов!) и божественная библиотека thrust, автоматизирующая с помощью RAII и функторов рутинные операции типа map/reduce. Вот к примеру код генерации лучей для моего мини-трассировщика, который я написал позавчера:struct Ray { float3 origin; float3 dir; __host__ __device__ Ray(){} Ray(float3 origin_, float3 dir_) : origin(origin_), dir(dir_) {}};struct RndToRays { float3 p1; float3 v1; float3 v2; float3 u; float3 v; float half_size; float3 dir; RndToRays(float3 p1_, float3 v1_, float3 v2_, float angle_size_, float3 dir_) : p1(p1_), v1(v1_), v2(v2_), half_size(cosf(angle_size_/2)), dir(dir_) { u = norm(v1); v = norm(v2); } Ray operator()(float4 rnd) { float3 point; point.x = p1.x + rnd.x v1.x + rnd.y v2.x; point.y = p1.y + rnd.x v1.y + rnd.y v2.y; point.z = p1.z + rnd.x v1.z + rnd.y v2.z; float phi = rnd.z 2.0f M_PI; float theta = acosf(half_size + rnd.w (1.0f - half_size)); float3 dir2; dir2.x = sinf(theta) (cosf(phi) u.x + sinf(phi) v.x) + cosf(theta) dir.x; dir2.y = sinf(theta) (cosf(phi) u.y + sinf(phi) v.y) + cosf(theta) dir.y; dir2.z = sinf(theta) (cosf(phi) u.z + sinf(phi) v.z) + cosf(theta) dir.z; return Ray(point - dir2, dir2); }};//и вызов thrust::device_vector<float4> d_points(N); curandGenerateUniform(qrng, (float)thrust::raw_pointer_cast(&d_points[0]), 4 N); thrust::device_vector<Ray> d_rays(N); thrust::transform(d_points.begin(), d_points.end(), d_rays.begin(), RndToRays({-10, -10, 0}, {20, 0, 0}, {0, 20, 0}, 0.5fM_PI/180, vec));Здесь ВООБЩЕ нет хуйни типа __device__, __kernel__ или подобного, тем не менее, это полноценный CUDA-код. Случайные числа я генерю с помощью curand, которая встроена в CUDA и затюнена Nvidia. Есть ли аналог curand под OpenCL? Есть, от васяна на гитхабе. Короче, по уровню поддержки это примерно как коммерческий юникс и линукс в 90-е. Наверное, у открытого OpenCL большое будущее, но сейчас и в ближайшие годы это говнецо.OpenCL - это несовместимая между собой жопа. Nvidia откровенно на него подзабила, что конечно делает их пидорасами типа Билли Гейтса, но нужно жить в реальном мире.К примеру, библиотека clFFT медленнее на Nvidia в 6 раз, потому что AMD спонсор clFFT. Поэтому CUDA-код все равно приходится писать - и писать его первым, потому что разработка на CUDA намного удобнее.Компилятор у AMD, например, хуевый, и любит вместо вывода сообщений об ошибках тупо падать, причем вовсе не в том месте, где должен. К примеру, у меня в коде был такой баг:context = clCreateContext (contextProperties, 1/deviceIdCount/, deviceIds.data (), nullptr, nullptr, &error);Закоменченный код вызывал падение внутри библиотеки clFFT, что привело сначала к дебагу этой либы, потом дебагу кода вокруг (методом "закомментируй код и посмотри упадет или нет", так как падает все все равно в случайном месте). При том, что на картах Nvidia все работает нормально даже с OpenCL. Не говоря уже о cuFFT, которая, как водится, круче clFFT.Далее нормально распространяются только сырцы, в которых может быть твое ноу-хау. Из-за этого приходится генерить кучу бинарников на разном железе. Кронос придумали для этой цели SPIR, который уже есть в виде совершенно разных SPIR и SPIR-V. Но эта хуйня не везде поддерживается. В итоге приходится иметь CUDA для nvidia, и зоопарк OpenCL-бинарников для AMD.Поэтому в общем-то CUDA - промышленный стандарт для суровых перцев, занимающихся симуляциями и машинлернингом. А OpenCL - это хуитка для consumer-рынка, покрыть долю людей, у которых AMD вместо видеокарты. Я кончил.
>>797077AMD били в гигафлопсах на ватт и на бакс. Но только за счет архитектуры. Грубо говоря, биткоины майнить на амд удобнее, но вот какой-то нетривиальный код, требущий умной работы с глобальной памятью и т. п., быстрее уже на нвидии.
>>797098Это всё оно так, но для меня лично нивелируется 1.5 словом: vendor-lockвот где невидии нет и не будет, там не взлетит. Поэтому писать код, который будет работать у юзера на утюге/ноуте/десктопе - лучше на OpenCL. >Поэтому в общем-то CUDA - промышленный стандарт для суровых перцев, занимающихся симуляциями и машинлернингом.ну, у caffe уже есть OpenCL, мне норм. Хотя и сильно позже, да. Опять же, у интела есть Xeon Phi, на котором OpenCL есть, а куды нет. То есть уже имеем минимум 2х крупных игроков только с OpenCL и одного с кудой. Если твой мини трассировщик готовится вырасти во что-то эдакое макси года через 3 - откуда тебе знать, который из 3х вендоров будет тогда лидировать? С OpenCL в любом случае не пропадёшь, а куду возможно придётся переписывать.> context = clCreateContext (contextProperties, 1/deviceIdCount/,> deviceIds.data (), nullptr, nullptr, &error);ну ёпты, пользуйся pycuda/pyopencl на прототипе, чтобы не ебаться с host-кодом. Как с архитектурой определишься, переноси на си/плюсы/расты
>>797386>писать код, который будет рандомно не компилиться, крешить юзеру утюг/ноут/десктоп, работать некорректно и медленно - лучше на OpenCLftfyOpenCL - это, в основном, AMD-шный высер, на который AMD же и забил болт. На какой ещё платформе, скажем, нужен mul24 (ещё и без mul24_hi)?А для phi конпелятор здорового человека - это няшная сишечка с родными интринсиками.
>>797658> OpenCL - это, в основном, AMD-шный высерВообще это задумка Apple, которой не понравилась вендорлочность куды.> на который AMD же и забил болтэто стакой же стандарт, как и OpenGL. Развивается бюрократической организацией согласно интересам всех участников.> рандомно не компилитьсятак ты и пиши так, чтобы компилился. Выведи часть логики за уровень ядра например, всё равно при большом числе условных переходов у тебя исполнители будут простаивать. Заодно и код упростится == станет более портабельным.> А для phi конпелятор здорового человека - это няшная сишечка с родными интринсиками.итого: можно иметь 3 реализации алгоритма под каждый конпелятор, а можно одну. Какого покемона выберешь ты?
>>797386>Это всё оно так, но для меня лично нивелируется 1.5 словом: vendor-lock1000 строк кода переписываются с CUDA на OpenCL за вечер. При этом код на CUDA уже будет отлажен и ты будешь уверен, что в нем нет багов, темлейты придется перепердолить в дефайны, а вот если ты начнешь с OpenCL, тебя может ждать что угодно.Ну и да, OpenCL - это тоже lock-in, потому что вот конкретно под AMD нужно профайлить и тюнить алгоритмы. clFFT например залочен под AMD, я писал выше.
>ну ёпты, пользуйся pycuda/pyopencl на прототипе, чтобы не ебаться с host-кодомТак у CUDA я и не ебусь с host-кодом, это у AMD спецификация OpenCL реализована криво. Это в общем-то основная проблема OpenCL, при заявляемой совместимости код не совместим и реализации содержат баги в разных местах. В итоге проще для nvidia писать CUDA-код и не ебать мозг.>Если твой мини трассировщик готовится вырасти во что-то эдакое макси года через 3 - откуда тебе знать, который из 3х вендоров будет тогда лидировать?Я GPGPU занимаюсь с 2012, знаешь, что изменилось с тех пор? Нихуя.
>>797741Для линукса у AMD есть открытый драйвер с реализацией OpenCL (в частности, беспокоящий тебя clCreateContext в mesa -> clover)Бери, правь. Ах да, ты не отрепортил в clFFT? Собери всё с дебагом и запости трейс и дамп в https://github.com/clMathLibraries/clFFT/issues - на то он и опенсорс.> Я GPGPU занимаюсь с 2012, знаешь, что изменилось с тех пор? Нихуя.интел только сейчас начал подтягивать свои GPU до приличного уровня и реализовывать в них OpenCL. Глядишь, скоро упрётся в невозможность нарастить мощь проца, кроме как наращиванием числа ядер/GPU/AVX (GPU там уже больше 50% кристалла) и тоже начнёт ещё активнее агитировать за OpenCL, так как иначе всю мощь будет не выжать.
>>797716>пиши так, чтобы компилилсяТеоретик вообще не палится.
>>797098Что характерно, у нвидии OpenCL умеет в asm-блоки (пусть и в PTX), а у AMD - нет. Там даже сраный carry flag не вытащить, хотя в железе он есть.
>>797808Всё оч просто: пишешь так, чтобы конпелялось на самой убогой платформе, например pocl или beignet. На остальных заработает автоматическтки.>>797814> у нвидии OpenCL умеет в asm-блоки (пусть и в PTX), а у AMD - нетВо-первых, эксперементальные реализации уже есть, причём нативного асма. Во-вторых, это не нужно в большинстве случаев: если очень хочется писать непереносимый код - бери невидию и ебись с кудой.
>>798197>не нужноТипичный амдебил типичен.
>>798202> Типичный амдебилда я больше кода отладил на beignet на рабочем ноуте, чем дома на амд, так что мимо.
Сколько мс на реальных проектах занимает переброска с девайса на хост и/или обратно 1 Гб?
>>7993171G в мою видяху не лезет, бо всего 2. Вот тебе тест на 0.5http://ideone.com/Vo4AnPНа моей 7870:Copy to device: 0.409sRun: 0.222sCopy from device: 0.219sНа A8-5500, CPU вариант (GPU не будет, ибо у месы пока кривой OpenCL):Copy to device: 0.411sRun: 0.917sCopy from device: 0.399s
>>799317Меньше секунды, а так хуй знает. Оно ограничено шиной, сам считай, не твоя армия.
>>800070>>800084Спасибо за тред и ответ! Как появится свободное время, то покопаюсь детальней в куде. Вроде не особо много чего придется учить, чтобы начать использовать.
Для тех, у кого нет видеокарты амд/невидии, но кто тоже хочет приобщиться святых таинств:Опенсорс варианты:http://cgit.freedesktop.org/beignet/ - реализация для intel gpu начиная с HD4000 (не очень бажная, если репортить - быстро фиксят)http://portablecl.org/ - CPU-only реализация (весьма бажная, зато тоже можно слать багрепорты)Для ценителей проприетарных продуктов у интел и амд также есть CPU реализации:https://software.intel.com/en-us/articles/opencl-drivershttp://developer.amd.com/tools-and-sdks/opencl-zone/amd-accelerated-parallel-processing-app-sdk/
>>800113Асло сейчас и зеленые, и красные движутся в сторону unified memory - когда трансферить явно ничего не надо, типа DMA и тормозит оно при каждом обращении.
>>801106Оно и сейчас как-бы есть через CL_MEM_USE_HOST_PTR только нужно, чтобы код брал предсказуемые данные, а не случайно определял индекс во время исполнения. Тогда можно даже выиграть по времени за счёт того, что подгрузка данных будет идти одновременно с исполнением.
