Сучасна диджитал-освіта для дітей — безоплатне заняття в GoITeens ×
Mazda CX 5
×

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

Підписуйтеся на Telegram-канал «DOU #tech», щоб не пропустити нові технічні статті

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

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

👍ПодобаєтьсяСподобалось0
До обраногоВ обраному1
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);
}
Причин же может быть безумно много, от условий не в нужных местах до тормозов на PCI.

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

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

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

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

Ну так 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

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

вот с этой штукой ничего не ясно, если в 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 было.

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

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

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

кстати а как насчет Vulkan?

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

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

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

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

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

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

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

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

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

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

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

photos.app.goo.gl/vKYUSxAJmkH91W5K7

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

и аналога на CUDA

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

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

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

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

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

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

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

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

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

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

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