diff --git a/src/ClObjects.h b/src/ClObjects.h new file mode 100644 index 00000000..59750a45 --- /dev/null +++ b/src/ClObjects.h @@ -0,0 +1,85 @@ +#pragma once + +#include +#include +#include "Util.h" + +namespace cl_objects { + class Device { + cl_device_id id; + public: + Device() : id(nullptr) {} + + explicit Device(cl_device_id id) : id(id) {} + + Device(const Device &other) : id(other.id) {} + + const cl_device_id &getId() const { + return id; + } + + std::string type() const { + auto type = getInfo(clGetDeviceInfo, id, CL_DEVICE_TYPE); + if (type & CL_DEVICE_TYPE_GPU) + return "GPU"; + else if (type * CL_DEVICE_TYPE_CPU) + return "CPU"; + } + + std::string name() const { + auto nameVec = getInfoVec(clGetDeviceInfo, id, CL_DEVICE_NAME); + return {nameVec.begin(), nameVec.end() - 1}; + } + + cl_ulong globalSize() const { + return getInfo(clGetDeviceInfo, id, CL_DEVICE_GLOBAL_MEM_SIZE); + } + + cl_ulong localSize() const { + return getInfo(clGetDeviceInfo, id, CL_DEVICE_LOCAL_MEM_SIZE); + } + + explicit operator bool() const { + return id == nullptr; + } + + }; + + template + class WrapperRAII { + CL_OBJECT_TYPE cl_object = nullptr; + RELEASE_FUNC* releaseFunc; + public: + template + WrapperRAII(CREATE_FUNC createFunc, RELEASE_FUNC releaseFunc, Args... args): releaseFunc(releaseFunc) { + cl_int errcode_ret = 0; + cl_object = createFunc(args..., &errcode_ret); + OCL_SAFE_CALL(errcode_ret); + } + + const CL_OBJECT_TYPE& getObject() const { + return cl_object; + } + + ~WrapperRAII() { + if (cl_object) + releaseFunc(cl_object); + } + }; + + class KernelsInProgram { + std::vector kernels; + public: + KernelsInProgram(cl_program program) { + kernels = getInfoVec(clCreateKernelsInProgram, program); + } + std::vector getKernels() const { + return kernels; + } + ~KernelsInProgram() { + for(cl_kernel kernel: kernels) { + clReleaseKernel(kernel); + } + } + }; +} \ No newline at end of file diff --git a/src/Util.h b/src/Util.h new file mode 100644 index 00000000..ada4215d --- /dev/null +++ b/src/Util.h @@ -0,0 +1,45 @@ +#pragma once + +#include + +template +std::string to_string(T value) { + std::ostringstream ss; + ss << value; + return ss.str(); +} + +void reportError(cl_int err, const std::string &filename, int line) { + if (CL_SUCCESS == err) + return; + + // Таблица с кодами ошибок: + // libs/clew/CL/cl.h:103 + // P.S. Быстрый переход к файлу в CLion: Ctrl+Shift+N -> cl.h (или даже с номером строки: cl.h:103) -> Enter + std::string message = "OpenCL error code " + to_string(err) + " encountered at " + filename + ":" + to_string(line); + throw std::runtime_error(message); +} + +#define OCL_SAFE_CALL(expr) reportError(expr, __FILE__, __LINE__) + + + +template +std::vector getInfoVec(CL_GET_INFO_F getInfoF, Args... args) { + SIZE_INFO_TYPE infoSize = 0; + OCL_SAFE_CALL(getInfoF(args..., 0, nullptr, &infoSize)); + + std::vector infoVec(infoSize); + OCL_SAFE_CALL(getInfoF(args..., infoSize, infoVec.data(), nullptr)); + return infoVec; +} + +template +R getInfo(CL_GET_INFO_F getInfoF, Args... args) { + std::size_t infoSize = 0; + OCL_SAFE_CALL(getInfoF(args..., 0, nullptr, &infoSize)); + + R info = 0; + OCL_SAFE_CALL(getInfoF(args..., infoSize, &info, nullptr)); + return info; +} \ No newline at end of file diff --git a/src/cl/aplusb.cl b/src/cl/aplusb.cl index 479624ac..df4080c9 100644 --- a/src/cl/aplusb.cl +++ b/src/cl/aplusb.cl @@ -11,13 +11,17 @@ // - На вход дано три массива float чисел; единственное, чем они отличаются от обычных указателей - модификатором __global, т.к. это глобальная память устройства (видеопамять) // - Четвертым и последним аргументом должно быть передано количество элементов в каждом массиве (unsigned int, главное, чтобы тип был согласован с типом в соответствующем clSetKernelArg в T0D0 10) -__kernel void aplusb(...) { +__kernel void aplusb(__global const float* as, __global const float* bs, __global float* cs, unsigned int size) { // Узнать, какой workItem выполняется в этом потоке поможет функция get_global_id // см. в документации https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ // OpenCL Compiler -> Built-in Functions -> Work-Item Functions + const unsigned int index = get_global_id(0); + if (index >= size) + return; // P.S. В общем случае количество элементов для сложения может быть некратно размеру WorkGroup, тогда размер рабочего пространства округлен вверх от числа элементов до кратности на размер WorkGroup // и в таком случае, если сделать обращение к массиву просто по индексу=get_global_id(0), будет undefined behaviour (вплоть до повисания ОС) // поэтому нужно либо дополнить массив данных длиной до кратности размеру рабочей группы, // либо сделать return в кернеле до обращения к данным в тех WorkItems, где get_global_id(0) выходит за границы данных (явной проверкой) + cs[index] = as[index] + bs[index]; } diff --git a/src/main.cpp b/src/main.cpp index d70cde7b..f8556c4c 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -3,33 +3,12 @@ #include #include -#include #include #include #include #include #include - - -template -std::string to_string(T value) { - std::ostringstream ss; - ss << value; - return ss.str(); -} - -void reportError(cl_int err, const std::string &filename, int line) { - if (CL_SUCCESS == err) - return; - - // Таблица с кодами ошибок: - // libs/clew/CL/cl.h:103 - // P.S. Быстрый переход к файлу в CLion: Ctrl+Shift+N -> cl.h (или даже с номером строки: cl.h:103) -> Enter - std::string message = "OpenCL error code " + to_string(err) + " encountered at " + filename + ":" + to_string(line); - throw std::runtime_error(message); -} - -#define OCL_SAFE_CALL(expr) reportError(expr, __FILE__, __LINE__) +#include "ClObjects.h" int main() { @@ -39,19 +18,38 @@ int main() { // TODO 1 По аналогии с предыдущим заданием узнайте, какие есть устройства, и выберите из них какое-нибудь // (если в списке устройств есть хоть одна видеокарта - выберите ее, если нету - выбирайте процессор) + auto platforms = getInfoVec(clGetPlatformIDs); + cl_objects::Device selectedDevice; + for (auto platform: platforms) { + auto devices = getInfoVec(clGetDeviceIDs, platform, CL_DEVICE_TYPE_ALL); + for (auto device: devices) { + selectedDevice = cl_objects::Device(device); + if (selectedDevice.type() == "GPU") + break; + } + if (selectedDevice && selectedDevice.type() == "GPU") + break; + } + if (selectedDevice) + throw std::runtime_error("Device not selected "); + + std::cout << "Selected device: " << selectedDevice.name() << std::endl; + std::cout << " type: " << selectedDevice.type() << std::endl; + std::cout << " global memory size(MB): " << selectedDevice.globalSize() / (1024 * 1024) << std::endl; + std::cout << " local memory size(KB): " << selectedDevice.localSize() / 1024 << std::endl; // TODO 2 Создайте контекст с выбранным устройством // См. документацию https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ -> OpenCL Runtime -> Contexts -> clCreateContext // Не забывайте проверять все возвращаемые коды на успешность (обратите внимание, что в данном случае метод возвращает // код по переданному аргументом errcode_ret указателю) // И хорошо бы сразу добавить в конце clReleaseContext (да, не очень RAII, но это лишь пример) - + cl_objects::WrapperRAII context(clCreateContext, clReleaseContext, nullptr, 1, &(selectedDevice.getId()), nullptr, nullptr); // TODO 3 Создайте очередь выполняемых команд в рамках выбранного контекста и устройства // См. документацию https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ -> OpenCL Runtime -> Runtime APIs -> Command Queues -> clCreateCommandQueue // Убедитесь, что в соответствии с документацией вы создали in-order очередь задач // И хорошо бы сразу добавить в конце clReleaseQueue (не забывайте освобождать ресурсы) - - unsigned int n = 1000 * 1000; + cl_objects::WrapperRAII cmdQueue(clCreateCommandQueue, clReleaseCommandQueue, context.getObject(), selectedDevice.getId(), 0); + unsigned int n = 100 * 1000 * 1000; // Создаем два массива псевдослучайных данных для сложения и массив для будущего хранения результата std::vector as(n, 0); std::vector bs(n, 0); @@ -69,7 +67,11 @@ int main() { // Данные в as и bs можно прогрузить этим же методом, скопировав данные из host_ptr=as.data() (и не забыв про битовый флаг, на это указывающий) // или же через метод Buffer Objects -> clEnqueueWriteBuffer // И хорошо бы сразу добавить в конце clReleaseMemObject (аналогично, все дальнейшие ресурсы вроде OpenCL под-программы, кернела и т.п. тоже нужно освобождать) - + size_t bufferSize = sizeof(float) * n; + cl_objects::WrapperRAII + buf_as(clCreateBuffer, clReleaseMemObject, context.getObject(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, bufferSize, as.data()), + buf_bs(clCreateBuffer, clReleaseMemObject, context.getObject(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, bufferSize, bs.data()), + buf_cs(clCreateBuffer, clReleaseMemObject, context.getObject(), CL_MEM_WRITE_ONLY , bufferSize, nullptr); // TODO 6 Выполните TODO 5 (реализуйте кернел в src/cl/aplusb.cl) // затем убедитесь, что выходит загрузить его с диска (убедитесь что Working directory выставлена правильно - см. описание задания), // напечатав исходники в консоль (if проверяет, что удалось считать хоть что-то) @@ -86,30 +88,37 @@ int main() { // TODO 7 Создайте OpenCL-подпрограмму с исходниками кернела // см. Runtime APIs -> Program Objects -> clCreateProgramWithSource // у string есть метод c_str(), но обратите внимание, что передать вам нужно указатель на указатель - + std::vector lengths = {kernel_sources.size()}; + std::vector sources = {kernel_sources.data()}; + cl_objects::WrapperRAII + program(clCreateProgramWithSource, clReleaseProgram, context.getObject(), sizeof(sources.data()) / sizeof(char*), sources.data(), lengths.data()); // TODO 8 Теперь скомпилируйте программу и напечатайте в консоль лог компиляции // см. clBuildProgram - + const cl_device_id deviceList[] = {selectedDevice.getId()}; + OCL_SAFE_CALL(clBuildProgram(program.getObject(), sizeof(deviceList) / sizeof(cl_device_id), deviceList, "", nullptr, nullptr)); // А также напечатайте лог компиляции (он будет очень полезен, если в кернеле есть синтаксические ошибки - т.е. когда clBuildProgram вернет CL_BUILD_PROGRAM_FAILURE) // Обратите внимание, что при компиляции на процессоре через Intel OpenCL драйвер - в логе указывается, какой ширины векторизацию получилось выполнить для кернела // см. clGetProgramBuildInfo - // size_t log_size = 0; - // std::vector log(log_size, 0); - // if (log_size > 1) { + if (getInfo(clGetProgramBuildInfo, program.getObject(), selectedDevice.getId(), CL_PROGRAM_BUILD_STATUS) != CL_BUILD_SUCCESS) + throw std::runtime_error("build failed"); + + // std::vector log = getInfoVec(clGetProgramBuildInfo, program.getObject(), selectedDevice.getId(), CL_PROGRAM_BUILD_LOG); + // if (log.size() > 1) { // std::cout << "Log:" << std::endl; // std::cout << log.data() << std::endl; // } // TODO 9 Создайте OpenCL-kernel в созданной подпрограмме (в одной подпрограмме может быть несколько кернелов, но в данном случае кернел один) // см. подходящую функцию в Runtime APIs -> Program Objects -> Kernel Objects - + cl_objects::KernelsInProgram kernels(program.getObject()); + cl_kernel kernel = kernels.getKernels()[0]; // TODO 10 Выставите все аргументы в кернеле через clSetKernelArg (as_gpu, bs_gpu, cs_gpu и число значений, убедитесь, что тип количества элементов такой же в кернеле) { - // unsigned int i = 0; - // clSetKernelArg(kernel, i++, ..., ...); - // clSetKernelArg(kernel, i++, ..., ...); - // clSetKernelArg(kernel, i++, ..., ...); - // clSetKernelArg(kernel, i++, ..., ...); + unsigned int i = 0; + OCL_SAFE_CALL(clSetKernelArg(kernel, i++, sizeof(cl_mem), &buf_as.getObject())); + OCL_SAFE_CALL(clSetKernelArg(kernel, i++, sizeof(cl_mem), &buf_bs.getObject())); + OCL_SAFE_CALL(clSetKernelArg(kernel, i++, sizeof(cl_mem), &buf_cs.getObject())); + OCL_SAFE_CALL(clSetKernelArg(kernel, i++, sizeof(unsigned int), &n)); } // TODO 11 Выше увеличьте n с 1000*1000 до 100*1000*1000 (чтобы дальнейшие замеры были ближе к реальности) @@ -127,7 +136,10 @@ int main() { timer t;// Это вспомогательный секундомер, он замеряет время своего создания и позволяет усреднять время нескольких замеров for (unsigned int i = 0; i < 20; ++i) { // clEnqueueNDRangeKernel... + cl_event event; + OCL_SAFE_CALL(clEnqueueNDRangeKernel(cmdQueue.getObject(), kernel, 1, nullptr, &global_work_size, &workGroupSize, 0, nullptr, &event)); // clWaitForEvents... + OCL_SAFE_CALL(clWaitForEvents(1, &event)); t.nextLap();// При вызове nextLap секундомер запоминает текущий замер (текущий круг) и начинает замерять время следующего круга } // Среднее время круга (вычисления кернела) на самом деле считается не по всем замерам, а лишь с 20%-перцентайля по 80%-перцентайль (как и стандартное отклонение) @@ -141,7 +153,7 @@ int main() { // - Флопс - это число операций с плавающей точкой в секунду // - В гигафлопсе 10^9 флопсов // - Среднее время выполнения кернела равно t.lapAvg() секунд - std::cout << "GFlops: " << 0 << std::endl; + std::cout << "GFlops: " << n * 1e-9 / t.lapAvg() << std::endl; // TODO 14 Рассчитайте используемую пропускную способность обращений к видеопамяти (в гигабайтах в секунду) // - Всего элементов в массивах по n штук @@ -149,7 +161,7 @@ int main() { // - Обращений к видеопамяти 2*n*sizeof(float) байт на чтение и 1*n*sizeof(float) байт на запись, т.е. итого 3*n*sizeof(float) байт // - В гигабайте 1024*1024*1024 байт // - Среднее время выполнения кернела равно t.lapAvg() секунд - std::cout << "VRAM bandwidth: " << 0 << " GB/s" << std::endl; + std::cout << "VRAM bandwidth: " << 3 * n * sizeof(float) / ((1 << 30) * t.lapAvg()) << " GB/s" << std::endl; } // TODO 15 Скачайте результаты вычислений из видеопамяти (VRAM) в оперативную память (RAM) - из cs_gpu в cs (и рассчитайте скорость трансфера данных в гигабайтах в секунду) @@ -157,18 +169,19 @@ int main() { timer t; for (unsigned int i = 0; i < 20; ++i) { // clEnqueueReadBuffer... + OCL_SAFE_CALL(clEnqueueReadBuffer(cmdQueue.getObject(), buf_cs.getObject(), true, 0, bufferSize, cs.data(), 0, nullptr, + nullptr)); t.nextLap(); } std::cout << "Result data transfer time: " << t.lapAvg() << "+-" << t.lapStd() << " s" << std::endl; - std::cout << "VRAM -> RAM bandwidth: " << 0 << " GB/s" << std::endl; + std::cout << "VRAM -> RAM bandwidth: " << (double) bufferSize / (1<<30) / t.lapAvg() << " GB/s" << std::endl; } // TODO 16 Сверьте результаты вычислений со сложением чисел на процессоре (и убедитесь, что если в кернеле сделать намеренную ошибку, то эта проверка поймает ошибку) - // for (unsigned int i = 0; i < n; ++i) { - // if (cs[i] != as[i] + bs[i]) { - // throw std::runtime_error("CPU and GPU results differ!"); - // } - // } - + for (unsigned int i = 0; i < n; ++i) { + if (cs[i] != as[i] + bs[i]) { + throw std::runtime_error("CPU and GPU results differ!"); + } + } return 0; }