From 12aee4b494a91dd80fdd00ad58b173b6f3e0acb7 Mon Sep 17 00:00:00 2001 From: vatican1 Date: Thu, 14 Sep 2023 10:17:39 +0300 Subject: [PATCH] task01 done --- src/cl/aplusb.cl | 10 +++- src/main.cpp | 126 +++++++++++++++++++++++++++++++++++++---------- 2 files changed, 108 insertions(+), 28 deletions(-) diff --git a/src/cl/aplusb.cl b/src/cl/aplusb.cl index 479624ac..1d1ba7fb 100644 --- a/src/cl/aplusb.cl +++ b/src/cl/aplusb.cl @@ -11,7 +11,15 @@ // - На вход дано три массива float чисел; единственное, чем они отличаются от обычных указателей - модификатором __global, т.к. это глобальная память устройства (видеопамять) // - Четвертым и последним аргументом должно быть передано количество элементов в каждом массиве (unsigned int, главное, чтобы тип был согласован с типом в соответствующем clSetKernelArg в T0D0 10) -__kernel void aplusb(...) { +__kernel void aplusb(__global float *a, + __global float *b, + __global float *c, + unsigned int n) +{ + int i = get_global_id(0); + if (i >= n) + return; + c[i] = a[i] + b[i]; // Узнать, какой workItem выполняется в этом потоке поможет функция get_global_id // см. в документации https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ // OpenCL Compiler -> Built-in Functions -> Work-Item Functions diff --git a/src/main.cpp b/src/main.cpp index d70cde7b..eac008f8 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -31,6 +31,41 @@ void reportError(cl_int err, const std::string &filename, int line) { #define OCL_SAFE_CALL(expr) reportError(expr, __FILE__, __LINE__) +std::pair getBetterDevice() +{ + cl_uint platformsCount = 0; + OCL_SAFE_CALL(clGetPlatformIDs(0, nullptr, &platformsCount)); + std::vector platforms(platformsCount); + OCL_SAFE_CALL(clGetPlatformIDs(platformsCount, platforms.data(), nullptr)); + + for (int platformIndex = 0; platformIndex < platformsCount; ++platformIndex) + { + cl_platform_id platform = platforms[platformIndex]; + + cl_uint devicesCount = 0; + OCL_SAFE_CALL(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &devicesCount)); + std::vector devices(devicesCount); + OCL_SAFE_CALL(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, devicesCount, devices.data(), nullptr)); + + for (int deviceIndex = 0; deviceIndex < devicesCount; ++deviceIndex) + { + cl_device_id device = devices[deviceIndex]; + cl_device_type device_type; + OCL_SAFE_CALL(clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(cl_device_type), &device_type, nullptr)); + if (device_type & CL_DEVICE_TYPE_GPU) + { + std::cout << "find GPU!" << std::endl; + return std::make_pair(platform, device); + } + } + } + cl_uint devicesCount = 0; + OCL_SAFE_CALL(clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, 0, nullptr, &devicesCount)); + std::vector devices(devicesCount); + OCL_SAFE_CALL(clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, devicesCount, devices.data(), nullptr)); + std::cout << "don't find GPU!" << std::endl; + return std::make_pair(platforms[0], devices[0]); +} int main() { // Пытаемся слинковаться с символами OpenCL API в runtime (через библиотеку clew) @@ -40,18 +75,34 @@ int main() { // TODO 1 По аналогии с предыдущим заданием узнайте, какие есть устройства, и выберите из них какое-нибудь // (если в списке устройств есть хоть одна видеокарта - выберите ее, если нету - выбирайте процессор) + std::pair device = getBetterDevice(); + + cl_ulong deviceNameSize = 0; + OCL_SAFE_CALL(clGetDeviceInfo(device.second, CL_DEVICE_NAME, 0, nullptr, &deviceNameSize)); + std::vector deviceName(deviceNameSize, 0); + OCL_SAFE_CALL(clGetDeviceInfo(device.second, CL_DEVICE_NAME, deviceNameSize, deviceName.data(), nullptr)); + std::cout << "DeviceName: " << deviceName.data() << 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_int errcode_res = CL_SUCCESS; + + cl_context context = clCreateContext(nullptr, 1, &device.second, nullptr, nullptr, &errcode_res);; + OCL_SAFE_CALL(errcode_res); + // 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_command_queue queue = clCreateCommandQueue(context, device.second, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &errcode_res); + OCL_SAFE_CALL(errcode_res); + + unsigned int n = 1000 * 1000 * 100; // Создаем два массива псевдослучайных данных для сложения и массив для будущего хранения результата std::vector as(n, 0); std::vector bs(n, 0); @@ -70,6 +121,14 @@ int main() { // или же через метод Buffer Objects -> clEnqueueWriteBuffer // И хорошо бы сразу добавить в конце clReleaseMemObject (аналогично, все дальнейшие ресурсы вроде OpenCL под-программы, кернела и т.п. тоже нужно освобождать) + cl_mem as_mem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * n, as.data(), &errcode_res); + OCL_SAFE_CALL(errcode_res); + cl_mem bs_mem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * n, bs.data(), &errcode_res); + OCL_SAFE_CALL(errcode_res); + cl_mem cs_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * n, cs.data(), &errcode_res); + OCL_SAFE_CALL(errcode_res); + + // TODO 6 Выполните TODO 5 (реализуйте кернел в src/cl/aplusb.cl) // затем убедитесь, что выходит загрузить его с диска (убедитесь что Working directory выставлена правильно - см. описание задания), // напечатав исходники в консоль (if проверяет, что удалось считать хоть что-то) @@ -80,36 +139,45 @@ int main() { if (kernel_sources.size() == 0) { throw std::runtime_error("Empty source file! May be you forgot to configure working directory properly?"); } - // std::cout << kernel_sources << std::endl; +// std::cout << kernel_sources << std::endl; } // TODO 7 Создайте OpenCL-подпрограмму с исходниками кернела // см. Runtime APIs -> Program Objects -> clCreateProgramWithSource // у string есть метод c_str(), но обратите внимание, что передать вам нужно указатель на указатель + const char * kernel_sources_pointer = kernel_sources.c_str(); + cl_program program = clCreateProgramWithSource(context, 1, &kernel_sources_pointer, nullptr, &errcode_res); + OCL_SAFE_CALL(errcode_res); // TODO 8 Теперь скомпилируйте программу и напечатайте в консоль лог компиляции // см. clBuildProgram + OCL_SAFE_CALL(clBuildProgram(program, 1, &device.second, nullptr, 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) { - // std::cout << "Log:" << std::endl; - // std::cout << log.data() << std::endl; - // } + size_t log_size = 0; + OCL_SAFE_CALL(clGetProgramBuildInfo(program, device.second, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size)); + std::vector log(log_size, 0); + OCL_SAFE_CALL(clGetProgramBuildInfo(program, device.second, CL_PROGRAM_BUILD_LOG, log_size, log.data(), nullptr)); + 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 + const char * kernel_name= "aplusb"; + cl_kernel kernel = clCreateKernel(program, kernel_name, &errcode_res); + OCL_SAFE_CALL(errcode_res); // 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; + clSetKernelArg(kernel, i++, sizeof(cl_mem), &as_mem); + clSetKernelArg(kernel, i++, sizeof(cl_mem), &bs_mem); + clSetKernelArg(kernel, i++, sizeof(cl_mem), &cs_mem); + clSetKernelArg(kernel, i++, sizeof(unsigned int), &n); } // TODO 11 Выше увеличьте n с 1000*1000 до 100*1000*1000 (чтобы дальнейшие замеры были ближе к реальности) @@ -125,15 +193,17 @@ int main() { size_t workGroupSize = 128; size_t global_work_size = (n + workGroupSize - 1) / workGroupSize * workGroupSize; timer t;// Это вспомогательный секундомер, он замеряет время своего создания и позволяет усреднять время нескольких замеров - for (unsigned int i = 0; i < 20; ++i) { - // clEnqueueNDRangeKernel... - // clWaitForEvents... + for (unsigned int i = 0; i < 20; ++i) + { + cl_event start_kernel; + OCL_SAFE_CALL(clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global_work_size, nullptr, 0, nullptr, &start_kernel)); + OCL_SAFE_CALL(clWaitForEvents(1, &start_kernel)); t.nextLap();// При вызове nextLap секундомер запоминает текущий замер (текущий круг) и начинает замерять время следующего круга } // Среднее время круга (вычисления кернела) на самом деле считается не по всем замерам, а лишь с 20%-перцентайля по 80%-перцентайль (как и стандартное отклонение) // подробнее об этом - см. timer.lapsFiltered // P.S. чтобы в CLion быстро перейти к символу (функции/классу/много чему еще), достаточно нажать Ctrl+Shift+Alt+N -> lapsFiltered -> Enter - std::cout << "Kernel average time: " << t.lapAvg() << "+-" << t.lapStd() << " s" << std::endl; + std::cout << "Kernel average time: " << t.lapAvg() << " +- " << t.lapStd() << " s" << std::endl; // TODO 13 Рассчитайте достигнутые гигафлопcы: // - Всего элементов в массивах по n штук @@ -141,7 +211,7 @@ int main() { // - Флопс - это число операций с плавающей точкой в секунду // - В гигафлопсе 10^9 флопсов // - Среднее время выполнения кернела равно t.lapAvg() секунд - std::cout << "GFlops: " << 0 << std::endl; + std::cout << "GFlops: " << n / t.lapAvg() / 1e9<< std::endl; // TODO 14 Рассчитайте используемую пропускную способность обращений к видеопамяти (в гигабайтах в секунду) // - Всего элементов в массивах по n штук @@ -149,26 +219,28 @@ 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: " << (double) (3 * n * sizeof(float)) / t.lapAvg() / (1024*1024*1024) << " GB/s" << std::endl; } // TODO 15 Скачайте результаты вычислений из видеопамяти (VRAM) в оперативную память (RAM) - из cs_gpu в cs (и рассчитайте скорость трансфера данных в гигабайтах в секунду) { timer t; - for (unsigned int i = 0; i < 20; ++i) { - // clEnqueueReadBuffer... + for (unsigned int i = 0; i < 20; ++i) + { + cl_event event; + OCL_SAFE_CALL(clEnqueueReadBuffer(queue, cs_mem, CL_TRUE, 0, n * sizeof(float), cs.data(), 0, nullptr, &event)); 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)n * sizeof(float) / t.lapAvg() / (1024*1024*1024) << " 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; }