Первые шаги с OpenCL или сказ о том как одинаковый код на GPU и CPU запускать

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

Введение

Думаю, что пересказывать Википедию об OpenCL особого смысла нет, но если в двух словах, то OpenCL — это язык, (фреймворк и платформа), который позволяет запускать один и тот же код на разных устройствах с разными архитектурами, а в особенности на высокопараллельных процессорах, вроде видеокарт и современных центральных процессоров. Основан стандарт на C99 и поддерживается The Khronos Group, на этом ликбез будем считать завершенным.

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

Сначала я опишу достаточно тривиальный код и те, кому совсем не терпится увидеть магиюOpenCL, могут пропустить первую часть (только прочитайте последний абзац, где я описываю функцию MathCalculations, это важно. А если вы знаете об OpenCL и вам хочется увидеть результаты тестов, то идите сразу в пятый раздел, но все равно загляните в MathCalculations).

int main(int argc, char* argv[])
{
	GenerateTestData();
	PerformCalculationsOnHost();

	//Get all available platforms
	vector<cl::Platform> platforms;
	cl::Platform::get(&platforms);

	for (int iPlatform=0; iPlatform<platforms.size(); iPlatform++)
	{
		//Get all available devices on selected platform
		std::vector<cl::Device> devices;
		platforms[iPlatform].getDevices(CL_DEVICE_TYPE_ALL, &devices);

		//Perform test on each device
		for (int iDevice=0; iDevice<devices.size(); iDevice++)
		{
			try 
			{ 
				PerformTestOnDevice(devices[iDevice]);
			} 
			catch(cl::Error error) 
			{
				std::cout << error.what() << "(" << error.err() << ")" << std::endl;
			}
			CheckResults();
		}
	}

	//Clean buffers
	delete[](pInputVector1);
	delete[](pInputVector2);
	delete[](pOutputVector);
	delete[](pOutputVectorHost);

	return 0;
}

Вот так выглядит main моей небольшой программки для тестирования OpenCL, а если точнее, то для расчета некоего абстрактного математического выражения, до которого мы попозже дойдем. Итак, давайте построчно разбираться, что же тут происходит.

Часть первая — Инициализация исходных данных и традиционный способ вычислений

GenerateTestData(); не делает ничего экстраординарного, а просто выделяет память под входные и выходные массивы, а также заполняет входные массивы случайными данными.

void GenerateTestData()
{
	pInputVector1 = new float[DATA_SIZE];
	pInputVector2 = new float[DATA_SIZE];
	pOutputVector = new float[DATA_SIZE];
	pOutputVectorHost = new float[DATA_SIZE];

	srand (time(NULL));
	for (int i=0; i<DATA_SIZE; i++)
	{
		pInputVector1[i] = rand() * 1000.0 / RAND_MAX;
		pInputVector2[i] = rand() * 1000.0 / RAND_MAX;
	}
}

Дальше идет немного более интересная функция:

void PerformCalculationsOnHost()
{
	cout << "Device: Host" << endl << endl;

	//Some performance measurement
	timeValues.clear();
	__int64 start_count;
	__int64 end_count;
	__int64 freq;
	QueryPerformanceFrequency((LARGE_INTEGER*)&freq);

	for(int iTest=0; iTest<(TESTS_NUMBER/10); iTest++)
	{
		QueryPerformanceCounter((LARGE_INTEGER*)&start_count);
		for(int iJob=0; iJob<DATA_SIZE; iJob++)
		{
			//Check boundary conditions
			if (iJob >= DATA_SIZE) break; 

			//Perform calculations
			pOutputVectorHost[iJob] = MathCalculations(pInputVector1[iJob], pInputVector2[iJob]);
		}
		QueryPerformanceCounter((LARGE_INTEGER*)&end_count);
		double time = 1000 * (double)(end_count - start_count) / (double)freq;
		timeValues.push_back(time);
	}
	hostPerformanceTimeMS = std::accumulate(timeValues.begin(), timeValues.end(), 0)/timeValues.size();

	PrintTimeStatistic();
}

В ней первый цикл

for(int iTest=0; iTest<(TESTS_NUMBER/10); iTest++)

нужен для того, чтобы провести тест несколько раз для получения более точного времени выполнения. Время вычисления каждого теста сохраняются в массиве timeValues из которого потом вычисляется среднее значение и сохраняется в hostPerformanceTimeMS.

Второй цикл

for(int iJob=0; iJob<DATA_SIZE; iJob++)

последовательно производит некие математические вычисления над элементами входных массивов и сохраняет их в выходном массиве.

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

Тут же стоит заглянуть в MathCalculations и увидеть, что там все совсем скучно:

float MathCalculations(float a, float b)
{
	float res = 0;
	res += a*a*0.315f + b*0.512f + 0.789f;
	res += a*a*0.15f + b*0.12f + 0.789f;
	res += a*a*0.35f + b*0.51f + 0.89f;
	res += a*a*0.31f + b*0.52f + 0.7f;
	res += a*a*0.4315f + b*0.512f + 0.4789f;
	res += a*a*0.515f + b*0.132f + 0.7859f;
	res += a*a*0.635f + b*0.521f + 0.89f;
	res += a*a*0.731f + b*0.152f + 0.7f;
	res += a*a*0.1315f + b*0.512f + 0.789f;
	res += a*a*0.115f + b*0.12f + 0.789f;
	res += a*a*0.135f + b*0.51f + 0.89f;
	res += a*a*0.131f + b*0.52f + 0.7f;
	res += a*a*0.14315f + b*0.512f + 0.4789f;
	res += a*a*0.1515f + b*0.132f + 0.7859f;
	res += a*a*0.1635f + b*0.521f + 0.89f;
	res += a*a*0.1731f + b*0.152f + 0.7f;
	return res;
}

Собственно смысла она особого не имеет (и очевидно, что может быть сильно упрощена), а служит как простая демонстрация чистых математических действий. Важно в ней то, что она находится в отельном .cpp файле и то, что в ней совершается много арифметических действий, но об этом чуть позже.

Часть вторая — Инициализация OpenCL

Итак, терпеливые дочитали до этой части и обрадовались, что начинается интересное, а нетерпеливые этого чувства испытать не смогут, они прошлый абзац пропустили:)

Сначала я скажу о том, что OpenCL Runtime API представляет из себя именно API для C, а не для C++. В целом, в этом нет ничего плохого кроме того, что для проверки ошибок надо проверять код, возвращаемый каждой функцией и это не очень удобно. А также надо вручную следить за освобождением выделенных ресурсов.
Но есть также и официальная C++ обертка (ее можно найти на сайте Khronos), которая представляет из себя набор классов, соответствующих объектам OpenCL и поддерживающим подсчеты ссылок (reference counting который) и бросание исключений в случае ошибок (исключения надо включать при помощи #define __CL_ENABLE_EXCEPTIONS). Вот эту самую обертку я и буду использовать в нашем тесте.

Итак первым делом мы получаем список доступных платформ:

vector<cl::Platform> platforms;
cl::Platform::get(&platforms);

Платформа в OpenCL соответствует вендору, т.е. у NVidia будет одна платформа с ее устройствами, у Intel другая итд итп. В моем случае мне доступны как раз две платформы NVidia и Intel.

Сразу еще один маленький трюк, C++ wrapper может пользоваться своими собственными векторами (если ему об этом сказать) или векторами из STD, так что если где-то в примерах попадется что-то вроде cl::vector, не пугайтесь, он знает оба формата.

После того как мы получили список платформ, для каждой платформы мы получаем список доступных устройств:

std::vector<cl::Device> devices;
platforms[iPlatform].getDevices(CL_DEVICE_TYPE_ALL, &devices);

Собственно устройства — это то, что будет выполнять наши вычисления. Это может быть и GPU, и CPU и какой-то специальный ускоритель, который подключен к хосту, т.е. той системе, на которой запускается OpenCL. Вместо CL_DEVICE_TYPE_ALL можно передать CL_DEVICE_TYPE_GPU, тогда он будет выдавать только видеокарты или CL_DEVICE_TYPE_CPU для центральных процессоров.

Для каждого найденного устройства я запускаю тест, о котором расскажу чуть ниже, и пытаюсь отловить исключения, которые бросит OpenCL в случае проблем, а если все прошло хорошо, то CheckResults сравнивает результаты с теми, которые мы насчитали в первой части на хосте и рассчитывает статистику ошибок.

Часть третья — Создание и запуск ядра

Тут мы подходим к самому интересному — вычислениям.

void PerformTestOnDevice(cl::Device device)
{
	cout << endl << "-------------------------------------------------" << endl;
	cout << "Device: " << device.getInfo<CL_DEVICE_NAME>() << endl << endl;

	//For the selected device create a context
	vector<cl::Device> contextDevices;
	contextDevices.push_back(device);
	cl::Context context(contextDevices);

	//For the selected device create a context and command queue
	cl::CommandQueue queue(context, device);

	//Clean output buffers
	fill_n(pOutputVector, DATA_SIZE, 0);

	//Create memory buffers
	cl::Buffer clmInputVector1 = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pInputVector1);
	cl::Buffer clmInputVector2 = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pInputVector2);
	cl::Buffer clmOutputVector = cl::Buffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pOutputVector);

	//Load OpenCL source code
	std::ifstream sourceFile("OpenCLFile1.cl");
	std::string sourceCode(std::istreambuf_iterator<char>(sourceFile),(std::istreambuf_iterator<char>()));

	//Build OpenCL program and make the kernel
	cl::Program::Sources source(1, std::make_pair(sourceCode.c_str(), sourceCode.length()+1));
	cl::Program program = cl::Program(context, source);
	program.build(contextDevices);
	cl::Kernel kernel(program, "TestKernel");

	//Set arguments to kernel
	int iArg = 0;
	kernel.setArg(iArg++, clmInputVector1);
	kernel.setArg(iArg++, clmInputVector2);
	kernel.setArg(iArg++, clmOutputVector);
	kernel.setArg(iArg++, DATA_SIZE);

	//Some performance measurement
	timeValues.clear();
	__int64 start_count;
	__int64 end_count;
	__int64 freq;
	QueryPerformanceFrequency((LARGE_INTEGER*)&freq);

	//Run the kernel on specific ND range
	for(int iTest=0; iTest<TESTS_NUMBER; iTest++)
	{
		QueryPerformanceCounter((LARGE_INTEGER*)&start_count);

		queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(DATA_SIZE), cl::NDRange(128));
		queue.finish();

		QueryPerformanceCounter((LARGE_INTEGER*)&end_count);
		double time = 1000 * (double)(end_count - start_count) / (double)freq;
		timeValues.push_back(time);
	}

	PrintTimeStatistic();

	// Read buffer C into a local list
	queue.enqueueReadBuffer(clmOutputVector, CL_TRUE, 0, DATA_SIZE * sizeof(float), pOutputVector);
}

Первым делом мы выводим имя устройства, полученное таким путем:

device.getInfo<CL_DEVICE_NAME>()

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

Затем мы создаем контекст:

vector<cl::Device> contextDevices;
contextDevices.push_back(device);
cl::Context context(contextDevices);

С контекстами все не так просто… При создании контекста, мы передаем список устройств, которые мы хотим в него включить, но тут есть ограничение: только устройства на одной платформе могут быть в одном контексте, т.е. сделать контекст с GPU и CPU (в случае Intel/NVidia) не получится. В случае нескольких устройств в одном контексте, все буферы будут синхронизироваться автоматически на разных устройствах. С одной стороны, это упрощает поддержку multi-GPU, а с другой стороны никто не знает как, что и когда драйвер будет синхронизировать, а эффективность передачи данных является критичным для получения высокой производительности ради которой все и затевается. Поэтому я обычно создаю отдельный контекст для каждого устройства и вручную распределяю данные. Таким образом всегда известно, что, где, когда происходит.

Следующий шаг — это создание очереди команд для устройства:

cl::CommandQueue queue(context, device);

Эта самая очередь привязывается к конкретному устройству и, в теории, может быть Out of Order, но по факту, я такого поведения не замечал. Очередей для одного устройства может быть несколько, причем можно синхронизировать команды из разных очередей, но в пределах одного контекста.

Далее мы создаем буферы для выходных и выходного векторов:

//Create memory buffers
cl::Buffer clmInputVector1 = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pInputVector1);
cl::Buffer clmInputVector2 = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pInputVector2);
cl::Buffer clmOutputVector = cl::Buffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pOutputVector);

При создании буфера указывается контекст (а не конкретное устройство), его объем и, при желании и использовании флага CL_MEM_COPY_HOST_PTR, указатель на данные, которые будут в него скопированы при создании. Как я говорил ранее, C++ wrapper использует подсчет ссылок, поэтому удалять буфер вручную не надо, в отличие от чистого C API.

Далее нам необходимо создать ядро, код которого хранится в файле «OpenCLFile1.cl». Для этого мы читаем текст из файла, создаем OpenCL программу, компилируем ее и получаем из нее ядро с именем «TestKernel», которое вы увидите в следующей части.

cl::Program::Sources source(1, std::make_pair(sourceCode.c_str(), sourceCode.length()+1));
cl::Program program = cl::Program(context, source);
program.build(contextDevices);
cl::Kernel kernel(program, "TestKernel");

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

Далее нам нужно установить аргументы, которые будут передаваться ядру. В отличие от CUDA, нужно вызывать специальные функции (в случае C++ wrapper'а, методы) для каждого аргумента и при необходимости указывать размер аргумента.

int iArg = 0;
kernel.setArg(iArg++, clmInputVector1);
kernel.setArg(iArg++, clmInputVector2);
kernel.setArg(iArg++, clmOutputVector);
kernel.setArg(iArg++, DATA_SIZE);

Теперь мы подошли к самому главному — запуску ядра:

queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(DATA_SIZE), cl::NDRange(128));

Собственно queue.enqueueNDRangeKernel добавляет команду запуска ядра в очередь команд и устанавливает количество элементов, которые будут обработаны, а также размер группы. О группах я расскажу отдельно (в другой статье), но сейчас упомяну лишь тот факт, что все элементы всегда разбиваются на группы и от размера группы может сильно зависеть производительность. В нашем случае количество элементов равно DATA_SIZE, а размер группы 128. Во время выполнения ядра, оно будет запущено DATA_SIZE раз (в неизвестной последовательности и возможно одновременно) и при каждом запуске ему будет передана информация о том, какой именно элемент обрабатывается.
enqueueNDRangeKernel является не блокирующей, поэтому после запуска ядра, мы должны дождаться его завершения, для чего и служит:

queue.finish();

Фактически finish выполняет две задачи:
1) Пересылает все команды в устройство (выполнение enqueueNDRangeKernel гарантирует, что драйвер получил команду и поставил ее в очередь, но не гарантируют ее запуск на устройстве, причем довольно часто может проходить достаточно длительное время перед реальным запуском ядра).
2) Ждет завершения всех команд в очереди.
Если нужно выполнить только первую часть, существует команда push (clFlush), которая является не блокирующей, но заставляет драйвер начать выполнение команд из очереди.

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

queue.enqueueReadBuffer(clmOutputVector, CL_TRUE, 0, DATA_SIZE * sizeof(float), pOutputVector);

В зависимости от второго аргумента, enqueueReadBuffer может быть блокирующей или не блокирующей. В нашем случае, она блокирующая, поэтому нет необходимости вызывать finish отдельно. Синтаксис простой: первый аргумент — откуда читать, четвертый аргумент — сколько читать и последний аргумент — куда читать. Есть еще параметр, который задает смещение от начала входного буфера, которое надо использовать в случае, если нужно считать данные не сначала, так как мы не можем использовать адресную арифметику для буферов OpenCL на хосте.

Часть четвертая — Код OpenCL kernel

А вот тут мы и дошли до того места, где нам надо начинать писать код (хотя это и кодом назвать сложно, так… баловство:)) на OpenCL. Вот так выглядит OpenCLFile1.cl:

#include "MathCode.cpp"
__kernel void TestKernel(
	__global const float* pInputVector1, 
	__global const float* pInputVector2, 
	__global float* pOutputVectorHost, 
	int elementsNumber)
{
    //Get index into global data array
    int iJob = get_global_id(0);

    //Check boundary conditions
    if (iJob >= elementsNumber) return; 

    //Perform calculations
    pOutputVectorHost[iJob] = MathCalculations(pInputVector1[iJob], pInputVector2[iJob]);
}

Итак по порядку:
Первым делом мы включаем в наш код файл MathCode.cpp, который содержит математическую функцию, ту самую на которую я просил обратить внимание ранее и ту самую, которая используется для традиционных вычислений на хосте. Как вы видите, мы даже не копируем код, мы используем один и тот же файл с математическим кодом.
Дальше мы создаем ядро, которое помечаем ключевым словом __kernel. Некоторые аргументы ядра также помечены ключевым словом __global, которое указывает на то, что это буфер в глобальной памяти устройства, созданный нами в коде хоста.
В коде ядра мы получаем номер элемента, который необходимо обработать:

int iJob = get_global_id(0);

Параметр get_global_id указывает на измерение, так как обрабатываемые элементы могут представлять из себя 1, 2 или 3мерный массив.
Затем проверяем граничные условия:

if (iJob >= elementsNumber) return; 

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

pOutputVectorHost[iJob] = MathCalculations(pInputVector1[iJob], pInputVector2[iJob]);

Часть пятая — Тестирование и замеры производительности

Вот и пришло время запустить приложение, оценить производительность и сделать некоторые выводы.

Я запускал тест на двух машинах и получил интересные результаты:
Ноутбук (CPU: Intel® Core™ i7-820QM, GPU: NVidia Quadro FX 2800M):

Host: 959.256 ms
CPU: 82.4163 ms (13.106X faster then host)
GPU: 9.90836 ms (109.014X faster then host)

Десктоп (CPU: Intel® Core™ i7-2600, GPU: NVidia GeForce GTX 580):

Host: 699.031 ms
CPU: 27.7833 ms (25.159X faster then host)
GPU: 2.06257 ms (338.897X faster then host)

Device: Host

Calculation time statistic: (20 runs)
Med: 959.256 ms (1.12602X faster then host)
Avg: 1080.15 ms
Min: 933.554 ms
Max: 1319.19 ms


-------------------------------------------------
Device: Quadro FX 2800M

Calculation time statistic: (200 runs)
Med: 9.90836 ms (109.014X faster then host)
Avg: 10.7231 ms
Min: 9.82841 ms
Max: 135.924 ms

Errors:
avgRelAbsDiff = 5.25777e-008
maxRelAbsDiff = 5.83678e-007

-------------------------------------------------
Device: Intel(R) Core(TM) i7 CPU       Q 820  @ 1.73GHz

Calculation time statistic: (200 runs)
Med: 82.4163 ms (13.106X faster then host)
Avg: 85.2226 ms
Min: 79.4138 ms
Max: 113.03 ms

Errors:
avgRelAbsDiff = 3.64332e-008
maxRelAbsDiff = 4.84797e-007

Device: Host

Calculation time statistic: (20 runs)
Med: 699.031 ms (0.999956X faster then host)
Avg: 699.1 ms
Min: 691.544 ms
Max: 715.233 ms


-------------------------------------------------
Device: GeForce GTX 580

Calculation time statistic: (200 runs)
Med: 2.06257 ms (338.897X faster then host)
Avg: 2.4 ms
Min: 2.03873 ms
Max: 82.0514 ms

Errors:
avgRelAbsDiff = 3.50006e-008
maxRelAbsDiff = 4.92271e-007

-------------------------------------------------
Device:         Intel(R) Core(TM) i7-2600 CPU @ 3.40GHz

Calculation time statistic: (200 runs)
Med: 27.7833 ms (25.159X faster then host)
Avg: 27.49 ms
Min: 27.0154 ms
Max: 35.8386 ms

Errors:
avgRelAbsDiff = 3.64377e-008
maxRelAbsDiff = 4.89584e-007

Итак, приступим к разбору результатов, а результаты, надо сказать, очень даже впечатляющие. GPU на ноутбуке в ~110X быстрее хоста, а на десктопе и вовсе в ~340X быстрее, впечатляющий результат, однако. Перед тем, как в меня начнут бросать тапки и говорить, что такое сравнение не правильное, я скажу, что в нем действительно есть несколько лукавств, но не более того.

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

Во-вторых, помните как выглядел математический код? Для тех, кто не смотрел на него, скажу, что это много много математических операций над одними и теми же данными, причем получился он путем простого копипаста и замены цифр в коэффициентах, а изначально он был проще и занимал всего одну строку, только вот когда я начал тестировать, результаты были не такие радостные, GPU было всего в 4-5 раз быстрее. Как думаете, почему? (вопрос риторический, можно не думать:)). А все просто, мы уперлись в производительность памяти. Я надеюсь, что попозже у меня дойдут руки и я напишу статью о взаимосвязи производительности памяти и процессора, но это отдельная история, в этой статье нам интересен лишь тот факт, что с данным ядром у нас получился чистый тест арифметической производительности процессора.

Учитывая эти два момента, можно сказать, что GPU действительно в сотни раз быстрее не-параллельного кода на CPU для чистой арифметики, что в целом, соответствует разнице в теоретической производительности. (Еще одна надежда на то, что дойдут руки замерить реальные цифры и их соответствие теории для другой статьи).

Но о том, что GPU быстро считает мы знаем, а в результате нашего теста получилось, что и CPU выполняет OpenCL код довольно быстро, если быть точным, то в 13X и 25Х раз быстрее, чем обычный код скомпилированный MSVC10 с дефолтными настройками. Давайте разбираться, как так получается и откуда взялись эти цифры.

Оба процессора содержат 4 реальных и 8 виртуальных ядер, а OpenCL как раз и сделан для того, чтобы все ядра использовать, но улучшение у нас гораздо больше, чем 4Х. А тут надо сказать спасибо Intel, которая в своей реализации OpenCL, добавила поддержку автоматической векторизации, т.е. без каких-либо изменений в коде, OpenCL использует SSE или AVX, в зависимости от того, что доступно. Учитывая, что SSE у нас 128битное, а AVX работает с 256битами, получается, что производительность должна подняться в 16X и 32X соответственно. Это уже ближе к истине, но все еще не совсем точное совпадение. А дальше нам надо вспомнить о такой радостной штуке, как TurboBoost. Процессоры эти работают на частотах 1,73GHz/3,06GHz (ноутбук) и 3,4GHz/3,8GHz (десктоп), но по факту могу сказать, что частота ноутбучного процессора скачет от 1,73 до 2,8 непрерывно, да и греется он весьма сильно (тут следует бросить большой каметь в Dell за кривую систему охлаждения), поэтому реально во время теста частоты 3,06GHz сколь нибудь значимое время мы не увидим. Плюс не надо забывать, что практический результат всегда меньше теоретически возможного (десктоп по идее должен работать быстрее), но как мы видим, 25Х улучшение производительности можно получить практически бесплатно на одном и том же железе.

Заключение

Задачей этой статьи не была попытка объяснить все детали работы с OpenCL, скорее это была попытка показать, что все не так уж сложно (вот тут я уже писал, что не все так просто) и в идеальных условиях можно получить очень впечатляющую производительность, причем даже на одном и том же железе, да к тому же можно использовать один и тот же код для всех устройств. Но помните, что это почти идеальные условия, которые бывают далеко не всегда.

PS: Для тех, кто хочет побаловаться с кодом и посмотреть тесты на другом железе, проект (и даже собранный экзешник) лежит на гитхабе. Для запуска может понадобиться OpenCL SDK от производителей вашего железа.

PS2: Если у кого-нибудь есть Ivy Bridge, было бы интересно посмотреть на тест встроенного видеоядра. Дело в том, что в последней версии OpenCL SDK, Intel открыла доступ к IGP, но только для последнего поколения процессоров, а таких у меня под рукой нет. Да и на результаты AMD интересно взглянуть.

Автор: Akson87