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

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

    Введение


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

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

    Сначала я опишу достаточно тривиальный код и те, кому совсем не терпится увидеть магиюOpenCL, могут пропустить первую часть (только прочитайте последний абзац, где я описываю функцию MathCalculations, это важно. А если вы знаете об OpenCL и вам хочется увидеть результаты тестов, то идите сразу в пятый раздел, но все равно загляните в MathCalculations).
    int main(int argc, char* argv[])
    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()
    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()
    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 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)
    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 than host)
    GPU: 9.90836 ms (109.014X faster than host)
    

    Десктоп (CPU: Intel® Core™ i7-2600, GPU: NVidia GeForce GTX 580):
    Host: 699.031 ms
    CPU: 27.7833 ms (25.159X faster than host)
    GPU: 2.06257 ms (338.897X faster than host)
    

    Полные результаты
    Device: Host
    
    Calculation time statistic: (20 runs)
    Med: 959.256 ms (1.12602X faster than 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 than 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 than 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 than 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 than 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 than 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 интересно взглянуть.
    Поделиться публикацией
    Ой, у вас баннер убежал!

    Ну. И что?
    Реклама
    Комментарии 21
      0
      Есть нетбук на AMD C60 (Fusion) — в кристалле встроенный Radeon. Если интересно — погоняю и отпишусь. На майнинге выдает 8 мегахэш.
        +3
        Как то просто выглядит функция float MathCalculations(float a, float b).

        Вы уверены, что компилятор ее сам не привел к виду res = a*a*A + b*B + C? Компиляторы нынче умные… да и случай простой. Может asm покажите?
          –1
          Я достаточно уверен, что не приводит. Изначально она такой упрощенной и была, тогда цифры сходились с разницей в производительности памяти, а тут сходятся с разницей в производительности арифметики. Возможно если сказать компилятору, чтобы он сильнее оптимизировал, то он заменит.
            0
            Все верно, компилятор(VS2010) не сообразит(на счет других не знаю). Пробовал включать sse2, оптимизацию по скорости, float point model(FAST). Код изменился незначительно(функция MathCalculations):
            БЫЛО:
            ...
            fld     dword ptr [ebp+8H] 
            fmul    st(0), st(0)           
            fld     dword ptr [ebp+0CH]    
            fld     qword ptr [?_0812]     
            fmul    st(0), st(1)           
            fld     st(1)                  
            fmul    qword ptr [?_0811]     
            fst     qword ptr [ebp-20H]    
            fld     st(2)                  
            fmul    qword ptr [?_0810]     
            fst     qword ptr [ebp-28H]
            ...
            


            СТАЛО:
            ...
            movaps  xmm6, xmm0               
            mulss   xmm6, dword ptr [?_0801] 
            movss   dword ptr [ebp-4H], xmm6 
            movaps  xmm2, xmm0               
            mulss   xmm2, dword ptr [?_0800] 
            movaps  xmm3, xmm0               
            mulss   xmm3, dword ptr [?_0799] 
            movaps  xmm4, xmm0               
            mulss   xmm4, dword ptr [?_0798] 
            movaps  xmm5, xmm0               
            mulss   xmm5, dword ptr [?_0797] 
            ...
            


            Полагаю, что оптимизация не получилось из-за того, что компилятор боится переставить операции местами, так как это повлияет на результат.
              0
              Почему повлияет на результат?
                0
                Из за формата представления чисел с точкой. Пример на python:
                Python 2.7.3 (default, Apr 10 2012, 23:24:47) [MSC v.1500 64 bit (AMD64)] on win32
                Type "help", "copyright", "credits" or "license" for more information.
                >>> a = 0.1 * 1 + 0.1 * 2 + 0.1 * 3 + 0.1 * 4 + 0.1 * 5 + 0.1 * 6 + 0.1 * 7 + 0.1 * 8 + 0.1 * 9
                >>> b = 0.1 * (1 + 2 + 3 + 4 + 5 + 6 + 7 + 8 + 9)
                >>> print (a, " ", b)
                (4.500000000000001, ' ', 4.5)
                >>> print (a == b)
                False
                >>>
                
                0
                Ну так не O3 единым, ещё надо поставить опцию /fp:fast. Попробуйте, пожалуйста :)
                  0
                  Уже. float point model(FAST) = /fp:fast
            +2
            Host реализация — это просто всё счетать в одном потоке. Можно написать тест на pthreads с очередями или какой-то другой библиотекой для работы с потоками и посмотреть как это соотносится с реализацией OpenCL на CPU.

            Не понятно почему MSVC10 не смог использовать SSE оптимизации. Интересно было бы попробовать современный GCC, ну должен же он найти вектора и начать с ними работать нормально.
              –1
              Где-то видел на просторах Интернета, что MSVC в принципе забил на SSE. Что 6.0 что 2010.
                0
                Ага, поэтому даже в настройках компилятора можно выставить поддержку SSE?
                Кстати, в настройках компилятора разрешено использовать SSE? (С/С++ -> Code Generation -> Enable Enhanced Instruction Set).
                  0
                  Насколько я понимаю, поддержка != автоматическая векторизация. Хотя можно попробовать. Интеловский компилятор вроде может автоматом векторизовывать.
                    0
                    попробовал, ничего не изменилось.
                      0
                      В ICC есть флаг –vec-report для получения отчета о выполненной векторизации. И, если не выполнилась, то выводится причина. Может и в MSVC есть что-то подобное?
                0
                Хост = один поток, все по дефолту. Многопоточную версию посмотреть было бы интересно, но больше 4х она при всем желании не сделает:)
                  0
                  Так вроде GCC и MSVC10 требуют флагов для SSE
                  0
                  Akson87
                  есть дома ivy bridge, но на мониторе сгорел hdmi. как смогу, так проведу тест.
                    0
                    Спасибо, было бы интересно взглянуть
                    0
                    А как различается точность вычислений с плавающей запятой на CPU и использованием OpenCL и без?
                    Мне для повышения точности моделирования приходится собирать приложение в Debug (чтобы получить 80-разрядные вычисления на double).
                      0
                      Больше 32/64 не будет в OpenCL, насколько я понимаю. На процессоре оно использует сейчас SSE/AVX, а там 80битной арифметики нет. На видеокартах тоже нет.
                      0
                      #include "MathCode.cpp"
                      


                      А можно также включать в opencl код файлы кода написанные на C#?

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

                      Самое читаемое