Принимайте участие в зарплатном опросе! Уже собрано почти 8 000 анкет.
×Закрыть

OpenCL, GPU и виртуальные машины

Имеется проблема в том чтобы запускать софт, использующих технологию OpenCL на виртуальных машинах. Вижу два с половиной варианта:
1. VMware ESX, которая умеет пробрасывать видеокарту с хоста на виртуалку.
2. Использовать проект pocl и собирать OpenCL под CPU (что лучше чем ничего).
2.5. Использовать OpenGL для вычислений (но это секс на лыжах и в гамаке).

Кто как решает проблему тяжелых вычислений на виртуальных машинах?

LinkedIn
Допустимые теги: blockquote, a, pre, code, ul, ol, li, b, i, del.
Ctrl + Enter
Допустимые теги: blockquote, a, pre, code, ul, ol, li, b, i, del.
Ctrl + Enter

Чет у меня появилось подозрение, что я с OpenCL именно так работаю
www.anekdot.ru/...​a-i-shurupovert_72491.gif

Ага. Но каждому нужно попробовать закатывать солнышко ручками.

Можно ли как-то повысить производительность этого кода, так как разница в производительности на GTX1070 и i7-4770K незначительна?

Имеем изображение 2.5 Мпикс, фильтр свертки 19×19.

Максимум производительности дало помещение фильтра свертки в константную память (насколько я понял система автоматически мапит константную память в локальную), остальные шаманства или не дают никакого результата или ухудшают результат на CPU (на GPU изменений нет).

naive + constant memory code

kernel void ImageKernelProceed(global read_only uint *bitmapData, __constant read_only float *kernelData,
global write_only uint *resData, global read_only int *params)
{
	int lImgWidth = params[0];
	int lImgHeight = params[1];
	int lKernelWidth = params[2];
	int lKernelHeight = params[3];
	int ly = get_global_id(0);
	int lx = get_global_id(1);

	int lKernelHeightRadius = lKernelHeight / 2;
	int lKernelWidthRadius = lKernelWidth / 2;
	float lIrR = 0;
	float lIrG = 0;
	float lIrB = 0;
	if ((ly < lKernelHeightRadius) | ((lImgHeight - lKernelHeightRadius) <= ly)
		| (lx < lKernelWidthRadius) | ((lImgWidth - lKernelWidthRadius) <= lx))
	{ 
		for (int lkY = 0; lkY < lKernelHeight; ++lkY)
		{ 
			int lSrcYPos = ly + lkY - lKernelHeightRadius;
			if ((lSrcYPos < 0) | (lImgHeight <= lSrcYPos))
				lSrcYPos = ly;
				for (int lkX = 0; lkX < lKernelWidth; ++lkX)
				{
					int lXPos = lx + lkX - lKernelWidthRadius;
					if ((lXPos < 0) | (lImgHeight <= lXPos))
						lXPos = lx;
					float lMult = kernelData[lkY * lKernelWidth + lkX];
					int lColorVal = bitmapData[lSrcYPos * lImgWidth + lXPos];
					lIrR += ((lColorVal >> 16) & 0xff) * lMult;
					lIrG += ((lColorVal >> 8) & 0xff) * lMult;
					lIrB += (lColorVal & 0xff) * lMult;
				}				
		}	
	}
	else
	{ 
		for (int lkY = 0; lkY < lKernelHeight; ++lkY)
		{ 
			int lSrcYPos = ly + lkY - lKernelHeightRadius;
			for (int lkX = 0; lkX < lKernelWidth; ++lkX)
			{
				int lXPos = lx + lkX - lKernelWidthRadius;
				float lMult = kernelData[lkY * lKernelWidth + lkX];
				int lColorVal = bitmapData[lSrcYPos * lImgWidth + lXPos];
				lIrR += ((lColorVal >> 16) & 0xff) * lMult;
				lIrG += ((lColorVal >> 8) & 0xff) * lMult;
				lIrB += (lColorVal & 0xff) * lMult;
			}
		}	
	}

	//fill result
	int lrB = (int)(lIrB + 0.5f);
	int lrR = (int)(lIrR + 0.5f);
	int lrG = (int)(lIrG + 0.5f);
	if ((255 < lrB) | (255 < lrG) | (255 < lrR) | (lrB < 0) | (lrG < 0) | (lrR < 0))
	{
		if (lrB < 0) lrB = 0;
		if (255 < lrB) lrB = 255;
		if (lrG < 0) lrG = 0;
		if (255 < lrG) lrG = 255;
		if (lrR < 0) lrR = 0;
		if (255 < lrR) lrR = 255;
	}
	resData[ly * lImgWidth + lx] = (0xff << 24) | ((lrR & 0xff) << 16) | ((lrG & 0xff) << 8) | (lrB & 0xff);
}

vector optimized + constant memory code

kernel void ImageKernelProceed(global read_only uint *bitmapData, __constant read_only float *kernelData,
global write_only uint *resData, global read_only int *params)
{
	int lImgWidth = params[0];
	int lImgHeight = params[1];
	int lKernelWidth = params[2];
	int lKernelHeight = params[3];
	int ly = get_global_id(0);
	int lx = get_global_id(1);

	int lKernelHeightRadius = lKernelHeight / 2;
	int lKernelWidthRadius = lKernelWidth / 2;
	float4 lIr = 0;

	//edge calculation
	if ((ly < lKernelHeightRadius) | ((lImgHeight - lKernelHeightRadius) <= ly)
		| (lx < lKernelWidthRadius) | ((lImgWidth - lKernelWidthRadius) <= lx))
	{ 
		for (int lkY = 0; lkY < lKernelHeight; ++lkY)
		{ 
			int lSrcYPos = ly + lkY - lKernelHeightRadius;
			if ((lSrcYPos < 0) | (lImgHeight <= lSrcYPos)) lSrcYPos = ly;
			for (int lkX = 0; lkX < lKernelWidth; ++lkX)
			{
					int lXPos = lx + lkX - lKernelWidthRadius;
					if ((lXPos < 0) | (lImgHeight <= lXPos)) lXPos = lx;
					float lMult = kernelData[lkY * lKernelWidth + lkX];
					int lColorVal = bitmapData[lSrcYPos * lImgWidth + lXPos];
					lIr += (float4)((lColorVal & 0xff), ((lColorVal >> 8) & 0xff), ((lColorVal >> 16) & 0xff), 0) * (float4)lMult;
			}				
		}
		//fill result
		lIr.s3 = 0xFF;
		resData[ly * lImgWidth + lx] = (uint)convert_uchar4_sat(lIr);
		return;

	}
	//inside image calculations
	for (int lkY = 0; lkY < lKernelHeight; ++lkY)
	{ 
		int lSrcYPos = ly + lkY - lKernelHeightRadius;
		int lXInd = lSrcYPos * lImgWidth + lx - lKernelWidthRadius; 
		//4 pixels aligned
		for (int lkX = 0; lkX < lKernelWidth / 4; ++lkX)
		{ 
			//load
			int lLInd = lkY * lKernelWidth + lkX * 4;
			float4 lMult4 = (float4)(kernelData[lLInd], kernelData[lLInd + 1], kernelData[lLInd + 2], kernelData[lLInd + 3]);
			lLInd = lXInd + lkX * 4;
			int4 lColorVal4 = (int4)(bitmapData[lLInd], bitmapData[lLInd + 1], bitmapData[lLInd + 2], bitmapData[lLInd + 3]);
			lIr += (float4)((lColorVal4.s0 & 0xff), ((lColorVal4.s0 >> 8) & 0xff), ((lColorVal4.s0 >> 16) & 0xff), 0) * (float4)lMult4.s0;
			lIr += (float4)((lColorVal4.s1 & 0xff), ((lColorVal4.s1 >> 8) & 0xff), ((lColorVal4.s1 >> 16) & 0xff), 0) *	(float4)lMult4.s1;
			lIr += (float4)((lColorVal4.s2 & 0xff), ((lColorVal4.s2 >> 8) & 0xff), ((lColorVal4.s2 >> 16) & 0xff), 0) * (float4)lMult4.s2;
			lIr += (float4)((lColorVal4.s3 & 0xff), ((lColorVal4.s3 >> 8) & 0xff), ((lColorVal4.s3 >> 16) & 0xff), 0) * (float4)lMult4.s3;				
		}
		//rest
		for (int lkX = (lKernelWidth / 4) * 4; lkX < lKernelWidth; ++lkX)
		{
			float lMult = kernelData[lkY * lKernelWidth + lkX];
			int lColorVal = bitmapData[lXInd + lkX];
			lIr += (float4)((lColorVal & 0xff), ((lColorVal >> 8) & 0xff), ((lColorVal >> 16) & 0xff), 0) * (float4)lMult;				
		}
	}	

	//fill result
	lIr.s3 = 0xFF;
	resData[ly * lImgWidth + lx] = (uint)convert_uchar4_sat(lIr);
}

Я бы сходу не смог подсказать. Профилировать пробовал? Разными путями, от профилировщика NVIDIA и Интела до метода «дихотомии» — вырезаешь куски кода и смотришь на скорость работы.
Ну и CPU и GPU полностью по разном у с памятью работают — она в корне разная по доступу.
Hy и CPU на 5ГГц может, а GPU только на 2ГГц.

Ну и я бы попробовал бы еще IPP и аналогичную функцию свертки из CUDA.

Причин же может быть безумно много, от условий не в нужных местах до тормозов на PCI.

Вообще задачка низкоуровневой оптимизации кода всегда сильно нетривиальная. Тут и кеши повсплывают и память и конвейеры, причем то что на одном проце будет оптимально, на другом таким не будет.

Ну и при фильтре длины 19 я бы уже и через FFT его сделал для сравнения.

Причин же может быть безумно много, от условий не в нужных местах до тормозов на PCI.

Замедление PCI я учел, просто запускаешь пустой кернел и смотришь. Так вот если убрать задержку PCI то на GPU скорость всего в 1.5 раза выше чем на CPU. Ну и фичи, если загружать в вектор как uchar4 (вместо получения данных через сдвиг), то на GPU скорость остается та же самая, а на CPU падает в два раза.

Так вот если убрать задержку PCI то на GPU скорость всего в 1.5 раза выше чем на CPU.

Тут тоже многое зависит от количества ядер GPU, размера обрабатываемых данных.
Разница основная с CPU, что на GPU все ядра в параллель работают.
Но в детали я глубоко не лазил. Нет смысла настолько глубоко лезть в GPU. Спроса на знание GPU глубоко в постсовке нет. А когда нужно что-то оптимизнуть, то достигаешь приемлемого и останавливаешься, большей частью стараюсь на уровне алгоритма это делать, а не глубокой оптимизации в железе.
На том же CPU есть векторные команды — SSE, AVX, AVX512 — часто и этого достаточно.

Ну так я и думаю что за счет AVX CPU показывает в 60 раз большую производительность чем однопоточный .NET код.

Хотя да, GPU я пока готовить не умею. Думаю буду использовать готовую интеловскую библиотеку сеток.

Если без CUDA, то я бы на твоем месте сосредоточился на интеловских либах, MKL, IPP, OpenVINO. Их обогнать сложно и они сами определяют, что у тебя за проц и юзают его максимуму.
OpenCL — это больше поиграться за счет работодателя.

А как интеловские либы перенести на ARM?

Никак. И там SSE, AVX нет, там Neon и свои странные GPU.

Ну так OpenVino это и есть OpenCL, правда сам код оптимизирован под Intel, но хорошо оптимизированный нормально пойдет и на странном Neon

techdecoded.intel.io/...​-in-the-openvino-toolkit

Ну оптимизация и здесь неплоха на том же CPU в 12 раз быстрее чем многопоточный код на TPL(.NET) и где-то в 60 раз быстрее чем однопоточный код на .NET. Хотя, хотелось бы чтобы эта цифра превысила 100.

Хосподи, но это же не С, чтобы так писать.

	//fill result
	int lrB = (int)(lIrB + 0.5f);
	int lrR = (int)(lIrR + 0.5f);
	int lrG = (int)(lIrG + 0.5f);
	if ((255 < lrB) | (255 < lrG) | (255 < lrR) | (lrB < 0) | (lrG < 0) | (lrR < 0))
	{
		if (lrB < 0) lrB = 0;
		if (255 < lrB) lrB = 255;
		if (lrG < 0) lrG = 0;
		if (255 < lrG) lrG = 255;
		if (lrR < 0) lrR = 0;
		if (255 < lrR) lrR = 255;
	}

int lrB = (int)clamp(lIrB + 0.5f, 0.0f, 255.0f);
int lrR = (int)clamp(lIrR + 0.5f, 0.0f, 255.0f);
int lrG = (int)clamp(lIrG + 0.5f, 0.0f, 255.0f);

					float lMult = kernelData[lkY * lKernelWidth + lkX];
					int lColorVal = bitmapData[lSrcYPos * lImgWidth + lXPos];
					lIrR += ((lColorVal >> 16) & 0xff) * lMult;
					lIrG += ((lColorVal >> 8) & 0xff) * lMult;
					lIrB += (lColorVal & 0xff) * lMult;

Используй сэмплеры! См. функцию read_imagef(). Весь код вверху — это одно умножение на вектор и один вызов функции. Сэмплеры эффективно используют кеш и могут читать данные наперёд, делают автоматическую конверсию входных типов во float.

                                for (int lkX = 0; lkX < lKernelWidth; ++lkX)
				{
					int lXPos = lx + lkX - lKernelWidthRadius;
					if ((lXPos < 0) | (lImgHeight <= lXPos))
						lXPos = lx;

Проверки на внутреннем цикле — это же полный писец, ты пускаешь нити последовательно! Что мешает посчитать это выше цикла и цикл крутить не от 0 до lKernelWidth, а только в нужных пределах?

Проверки на внутреннем цикле — это же полный писец, ты пускаешь нити последовательно! Что мешает посчитать это выше цикла и цикл крутить не от 0 до lKernelWidth, а только в нужных пределах?

Ну это цикл работает только в пределах границ и не особо замедляет. Я его даже убирал влияние на производительность около 5%.

int lrB = (int)clamp(lIrB + 0.5f, 0.0f, 255.0f);

Ну я это сделал в

vector optimized + constant memory code

 Но все равно это скорости сильно не прибавило.

Используй сэмплеры! См. функцию read_imagef().

Может помочь, попробую потом. Семплеры я еще не пробовал.

Ну это цикл работает только в пределах границ и не особо замедляет. Я его даже убирал влияние на производительность около 5%.

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

Я просто физически удалял весь двойной цикл после //edge calculation. Это дало 5% увеличения производительности.

Но у тебя же там как раз вся логика, т.е. ты просто убил весь кернел и быстрее не стало?

Нет я оставил //inside image calculations

Ну и оставил

  if ((ly < lKernelHeightRadius) | ((lImgHeight - lKernelHeightRadius) <= ly)
  | (lx < lKernelWidthRadius) | ((lImgWidth - lKernelWidthRadius) <= lx))
  {
    return;
  }

Самое печальное, что между оптимизированной версией через вектора и самой первой версий, которая взята с C# разница в производительности 10%.

Можна ще закинути вашу задачу на AWS, там є типи інстансів з підтримкою OpenCL. Буде швидко, а якщо налаштуєте запуск/зупинку за потреби — може бути не так і дорого. docs.aws.amazon.com/...​-computing-instances.html

Нашел самый простой способ, на виртуалку ставится AMD SDK и так как мы думает поддерживать с 1.2 то для тестов прокатит.
community.amd.com/thread/227948
vm6502q.readthedocs.io/...​latest/opencl.html#vmware

Но ты же выше спрашивал про:

Кто как решает проблему тяжелых вычислений на виртуальных машинах?

А так просто заюзать OpenCL на CPU просто и даже очень просто. Но это все одно будет тормознее, чем напрямую проц юзать, через те же MKL, IPP (даже для AMD).

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

Но правильно заданный и сформулированный вопрос в себе содержит обычно 60% ответа.

вот с этой штукой ничего не ясно, если в Xen и KVM четко указано что это SR-IOV то тут пёс его знает как графика прокидывается

Использовать OpenGL для вычислений

интересно как предполагаете получить OpenGL в виртуалке?

VMware Workstation Pro
Support has been added for the following standards:
DirectX 10
OpenGL 3.3
docs.vmware.com/...​ion-12-release-notes.html

OpenGL 3.3

только тут compute shader нету, нужно 4.3 это точно что в гамаке

DirectX 10

тут есть только очень порезаные

только тут compute shader нету, нужно 4.3 это точно что в гамаке

Люди умудрялись это делать на пиксельных шейдерах, это я про «секс, стоя на лыжах и в гамаке»

судя по тому что в DirectX 10
The Z dimension of numthreads is limited to 1
это и есть пиксельные шейдера

сто лет в обед — как openGL поддерживается vmWare player.
Если в игрушки играю- то только в виртуалке.

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

Ну ради справедливости заметим, что сравнение CUDA и OpenCL было.

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

Основная разница между ними в поддержке нвидией интелом и амд этих технологий.

Ну OpenCL поддерживается еще и ARM-ом, Qualcomm-ом. Так что в отличие от CUDA, переносимость больше как минимум на порядок.

Уже лет 15 как и состояние поддержки в общем-то не изменилось за 15 лет.

Ну да, CUDA — это решение одного производителя, что плохо, с OpenCL есть выбор и он довольно внушителен. Реально OpenCL начал распространение с 2015-го, когда был полностью принят стандарт 1.2 (OpenCL 1.2 extensions).

Согласен с тобой в этом. Но пока с кудой все продвинутее, чем с ОpenCL.
Может сейчас этот консорциум таки опомнится и начнет плотную работу с OpenCl.

Но я пока на куде посижу, с ней гемора сильно меньше, чем с OpenCL.

Без понятия. Слышал что такой есть и даже либы и даже кто-то его юзает, но не более. Может лет через дохреннадцать не будет CUDA и OpenCL, а будет Вулкан или еще какая хрень.

Vulkan — это OpenGL в ручном режиме. В плане числодробилок те же компьют шейдеры.

я вкурсе, и это же хорошо что все руками

У него поддержка только на топовом оборудовании, причем, к примеру, Apple не поддерживает Vulkan вообще.

ну тоесть Apple это не топовое оборудование.
вообще Vulkan была заявлено работа на железе которое суппортит OpenGL ES 3.1 для примера nvidia суппортит вулкан 1.1 на серии GeForce 600 которая вышла семь лет назад

ну тоесть Apple это не топовое оборудование.

с кривой клавиатурой, что царапается и прилично дороже топового среди ноутов. Какая прелесть.
Мне это напоминает анекдоты из 90-х про клоунов в малиновых пиджаках.

Официальной поддержки нет, есть опенсоурсная поддержка Vulkan.
vulkan.lunarg.com

типичные макопроблемы, в то время как вулкан можно запустить везде начиная от моб телефона заканчивая клаудом, для эпл нада пилить на ихней никому не нужной поделке

Так техника Apple позволяет исполнить самые сокровенные мечты рубиста. Ну прям слоган.

Apple и OpenGL не поддерживает %)

А можно вопрос?
Оно хоть что-то поддерживает?

Оно хоть что-то поддерживает?

Кошку поддерживает, когда она на макбуке спит.

Сын прислал фото, сказал, что кошка сменила мак на делл %)

photos.app.goo.gl/vKYUSxAJmkH91W5K7

Даже кошки что-то подозревают %)

Ну и еще дам совет. Не забывай, что через PCI пропихать большие объемы долго. Так что часто оказывается выгоднее что на CPU считать (интеловскими либами), чем запихать в GPU и назад выпихать.

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

Да знаю я это, но на легких алгоритмах и managed кода хватает. Паритет по вычислительной мощности, как показал эксперимент на GTX1070 и i7-4770k идет когда мы имеем 9 итераций на точку, если выше, то GPU быстрее. Хотя если запускать на процессоре с нулевым копированием через маппинг буфферов, то почему-то через OpenCL в 3.5 раза быстрее чем через TPL на managed коде.

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

Хотя если запускать на процессоре с нулевым копированием через маппинг буфферов

Это на встроенных картах в мамку? Или на чем-то типа Джетсонов или Коралов?

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

При обработке картинок, видео, начинает влиять и какая память и какой процессор и какие кеши у него и какая PCI и какая карточка. Причем простых зависимостей нет.

А с виртуалками и пробросами GPU добавится еще пачка мест тормозов.

Это на встроенных картах в мамку?

Нет это именно на target CPU_DEVICE.

А с виртуалками и пробросами GPU добавится еще пачка мест тормозов.

Это не страшно, тесты будут проходить чуть подольше, у клиента все равно будет нейтив.

Причем тут нейтив некий. Бутылочное горлышко — PCI.

Тем что код можно запускать на разных тарджетах, легкий код где в основном работа с памятью, на CPU без использования PCI, тяжелый (где пересылки ничто) на GPU. Ну или делать множество операций на GPU без извлечения буфера, извлекая только конечный результат.

Ну, пробуйте. Если потом сделаешь тут статью, как вы справились с OpenCL — будет круто. Смотришь можно будет переходить на более дешевые видеокарточки от AMD.

Если потом сделаешь тут статью, как вы справились с OpenCL — будет круто.

Для технических статей времени нет, я обычно техно-политические статьи пишу (это мое хобби).

я обычно техно-политические статьи пишу

рекламируй. где смотреть?

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

Но, если именно за это тебе платят, то почему не поиграться. Но всё одно не взлетит оное.

Просто нужна заглушка OpenCL для тестов библиотеки на сборочной машине. Билд и тесты проходит 2(сокращенный вариант)-6(полный) часа. У клиента оно будет работать на нативном железе с нативными библиотеками OpenCL.

У клиента оно будет работать на нативном железе с нативными библиотеками OpenCL

Чаще всего не будет. Вам придется еще заморочиться, чтобы у клиента таки OpenCL завелся. А включить его — это почти всегда еще тот танец с бубном.
С кудой тоже не всегда всё хорошо, а вот с OpenCL почти всегда гемор.

1. VMware ESX, которая умеет пробрасывать видеокарту с хоста на виртуалку.

Так это только один раз, вторую виртуалку в pass-through не пустишь.

2. Использовать проект pocl и собирать OpenCL под CPU (что лучше чем ничего).

Проще на CPU и посчитать — будет эффективнее.

2.5. Использовать OpenGL для вычислений (но это секс на лыжах и в гамаке).

OpenGL/OpenGL ES — Compute Shaders, не OpenCL, но и не обычные шейдеры, как по мне неплохой вариант (VMWare)

QEMU поддерживает GVT-g режим виртуализации интелловских GPU под линуксовым хостом. гостевая ОС видит GPU как буд-то в режиме pass-through. Потери от 5% до 70% в зависимости от видов загруки, например, glmark2 показывает отвратительные результаты (из-за того, что загружают GPU миллионами мелких тасков, вызывая аццкий driver overhead), а вот glBenchmark/glBench на высоте, но в среднем порядка 10-15% потерь по сравнению с хостом.

OpenCL стек под линукс для Intel GPU открытый: 01.org/compute-runtime

Кто как решает проблему тяжелых вычислений на виртуальных машинах?

Но вообще сама постановка задачи немного глуповатая — виртуалка и тяжелые вычисления как бы несовместимы из-за того, что кеш процессора очень мал и несколько виртуалок порвут его на куски как тузик грелку. Ты просто даже не сможешь полноценно загрузить GPU без простоев.

Ты просто даже не сможешь полноценно загрузить GPU без простоев.

Основная цель в том чтобы тяжелый CV(computer vision), проект мог использовать нечто быстрое для операций свертки и прочих матрично-векторных вычислений. Скорости на Managed коде уже не хватает даже если выполнять код параллельно на TPL или ThreadManager (своя технология, аналог TPL, работающая на .NET, Java, C++).

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

Похоже, пора думать, как туда подключить нативные куски.

Так уже придумано. OpenCL 1.2 (>90% устройств поддерживает) позволяет компилировать код в промежуточный, это уменьшает время компиляции OpenCL кода у хостовых runtimе библиотек. А компиляция OpenCL идет всегда под нативную платформу, нативным компилятором. Лучшее решение для managed кода, с максимальной совместимостью (у нас универсальный код от мейнфреймов до смартфонов).

Для нативных кусков пришлось бы увеличить команду раз так в 10. Намного проще сделать легкую деградацию вычислительных алгоритмов (выбрать менее точные, но легкие), чтобы поднять производительность (я так и делал 7 лет, подход вполне оправданый).

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

А придется, если в OpenCL полезете.

Там как раз все просто, есть стандарт, есть набор функций. Куча кода легко при небольших правках портируется с CUDA. Написал раз — будет работать везде. При этом OpenCL 1.2 (да чуть сложнее разработка чем на 2.0) поддерживается почти всем кроме виртуалок.

Так что с OpenCL сложность как при установке JVM, а дальше оно будет работать без какого-то изменения кода.

Вы сами выбираете свой путь. Я всего-лишь предупредил о подводных скалах.

нечто быстрое для операций свертки

Тебе просто нужно на OpenCL запрограммировать операцию свертки. Кстати у АМД есть либы для этого, но по качеству они сильно хуже аналогичных интеловских для CPU и нвидиевских для CUDA.

Фактически, если свяжешься с OpenCL, тебе придется самому написать на OpenCL аналоги BLAS, LAPACK, IPP (интеловского) и аналога на CUDA (от нвидия). Что-то уже есть в инете и от АМД, но почти все сильно не доделанное.

Ну и 3 твоих пути в самом верху — все это еще сильнее усложнят.

Ну и прямой путь для вас, это числодробление переносить в С, С++ и Куду (благо у нвидии море приличных либ для оного уже есть) и не забывайте про MKL, IPP — написать самому на таком же уровне качества и скорости очень и очень не просто.

нечто быстрое для операций свертки

Под сверткой я понимаю весь класс функций. Нейросеть — это тоже свертка и SVM — свертка.

и аналога на CUDA

Код CUDA а автоматическом режиме можно портировать под OpenCL, да скорость упадет в 3 раза, но всегда в процессе можно оптимизировать.

Ну и прямой путь для вас, это числодробление переносить в С, С++ и Куду

Клиенту при покупке библиотеки выдавать нивидевскую теслу в качестве подарка? В таком случае путь деградации алгоритмов(как сейчас) более удобен. Вместо размытия по по гаусу использовать размытия по среднему через интегральные окна, для предварительной обработки доунскейлить изображения в 4-8 раз по стороне(16-64 раза ускорение).

Под сверткой я понимаю весь класс функций. Нейросеть — это тоже свертка и SVM — свертка.

В твоем случае это матричная арифметика уже получается. Умножение, сложение матриц — это blas, как стандарт. У АМД какой-то blas есть.

А вот для нейросеток у нвидии есть cuDNN специализированный. Его во всех основных движках нейронок и юзают. TFLite гугл оптимизировала немного и под свои TPU.

для нейросеток у нвидии есть cuDNN специализированный

Опять же код нейросеток есть и под OpenCL, может он хуже оптимизирован, но даже падение производительности в 3 раза, компенсируется переносимостью.
github.com/intel/clDNN

Не думаю, что хуже. Но 99% ML сейчас сидит на CUDA и соответсвенно разные баги находятся быстрее и фиксятся нвидией.
И не забывай, что GPU с OpenCL от AMD, а не от Интела и здесь вы будете ловить много различных багов по несовместимости. А вот Nvidiа к OpenCL относится сильно не позитивно.

А вот Nvidiа к OpenCL относится сильно не позитивно.

Ну да они до сих пор указывают совместимость с OpenCL 1.2, хоть часть функций 2.0 уже работает.

поддерживает GVT-g

есть способ запустить это на windows host?

Сомневаюсь, GVT-g — это хостовая технология, которая трапается, когда VM отправляет batch buffer в GPU для обработки, этот пакет внаглую используется хостом для посылки в GPU (почти zero-copy, но не всегда). Т.е. клиенты/VM используют немодифицированные графические драйвера под Intel — как они есть, а вот на хосте драйвер интегрируется с хостовым виртуалайзером, чтобы иметь возможность залезать в память клиента. Поэтому врядли Интел сделала специальные драйвера для винды и множества виртуализаторов. Может разве что только под Xen, нужно проверять, информации нет.

Подписаться на комментарии