Skip to content

Commit

Permalink
done
Browse files Browse the repository at this point in the history
  • Loading branch information
AvvALlV committed Sep 17, 2023
1 parent 493f393 commit 6bc8404
Show file tree
Hide file tree
Showing 4 changed files with 194 additions and 47 deletions.
85 changes: 85 additions & 0 deletions src/ClObjects.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,85 @@
#pragma once

#include <CL/cl.h>
#include <string>
#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<cl_device_type>(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<unsigned char, size_t>(clGetDeviceInfo, id, CL_DEVICE_NAME);
return {nameVec.begin(), nameVec.end() - 1};
}

cl_ulong globalSize() const {
return getInfo<cl_ulong>(clGetDeviceInfo, id, CL_DEVICE_GLOBAL_MEM_SIZE);
}

cl_ulong localSize() const {
return getInfo<cl_ulong>(clGetDeviceInfo, id, CL_DEVICE_LOCAL_MEM_SIZE);
}

explicit operator bool() const {
return id == nullptr;
}

};

template<typename CL_OBJECT_TYPE, typename RELEASE_FUNC>
class WrapperRAII {
CL_OBJECT_TYPE cl_object = nullptr;
RELEASE_FUNC* releaseFunc;
public:
template <typename CREATE_FUNC, typename... Args>
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<cl_kernel> kernels;
public:
KernelsInProgram(cl_program program) {
kernels = getInfoVec<cl_kernel, cl_uint>(clCreateKernelsInProgram, program);
}
std::vector<cl_kernel> getKernels() const {
return kernels;
}
~KernelsInProgram() {
for(cl_kernel kernel: kernels) {
clReleaseKernel(kernel);
}
}
};
}
45 changes: 45 additions & 0 deletions src/Util.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,45 @@
#pragma once

#include <vector>

template<typename T>
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 <typename R, typename SIZE_INFO_TYPE, typename CL_GET_INFO_F, typename... Args>
std::vector<R> getInfoVec(CL_GET_INFO_F getInfoF, Args... args) {
SIZE_INFO_TYPE infoSize = 0;
OCL_SAFE_CALL(getInfoF(args..., 0, nullptr, &infoSize));

std::vector<R> infoVec(infoSize);
OCL_SAFE_CALL(getInfoF(args..., infoSize, infoVec.data(), nullptr));
return infoVec;
}

template <typename R, typename CL_GET_INFO_F, typename... Args>
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;
}
6 changes: 5 additions & 1 deletion src/cl/aplusb.cl
Original file line number Diff line number Diff line change
Expand Up @@ -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];
}
105 changes: 59 additions & 46 deletions src/main.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -3,33 +3,12 @@
#include <libutils/fast_random.h>
#include <libutils/timer.h>

#include <cassert>
#include <fstream>
#include <iostream>
#include <sstream>
#include <stdexcept>
#include <vector>


template<typename T>
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() {
Expand All @@ -39,19 +18,38 @@ int main() {

// TODO 1 По аналогии с предыдущим заданием узнайте, какие есть устройства, и выберите из них какое-нибудь
// (если в списке устройств есть хоть одна видеокарта - выберите ее, если нету - выбирайте процессор)
auto platforms = getInfoVec<cl_platform_id, cl_uint>(clGetPlatformIDs);
cl_objects::Device selectedDevice;
for (auto platform: platforms) {
auto devices = getInfoVec<cl_device_id, cl_uint>(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<cl_context, decltype(clReleaseContext)> 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<cl_command_queue, decltype(clReleaseCommandQueue)> cmdQueue(clCreateCommandQueue, clReleaseCommandQueue, context.getObject(), selectedDevice.getId(), 0);
unsigned int n = 100 * 1000 * 1000;
// Создаем два массива псевдослучайных данных для сложения и массив для будущего хранения результата
std::vector<float> as(n, 0);
std::vector<float> bs(n, 0);
Expand All @@ -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<cl_mem, decltype(clReleaseMemObject)>
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 проверяет, что удалось считать хоть что-то)
Expand All @@ -86,30 +88,37 @@ int main() {
// TODO 7 Создайте OpenCL-подпрограмму с исходниками кернела
// см. Runtime APIs -> Program Objects -> clCreateProgramWithSource
// у string есть метод c_str(), но обратите внимание, что передать вам нужно указатель на указатель

std::vector<size_t> lengths = {kernel_sources.size()};
std::vector<const char*> sources = {kernel_sources.data()};
cl_objects::WrapperRAII<cl_program, decltype(clReleaseProgram)>
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<char> log(log_size, 0);
// if (log_size > 1) {
if (getInfo<cl_build_status>(clGetProgramBuildInfo, program.getObject(), selectedDevice.getId(), CL_PROGRAM_BUILD_STATUS) != CL_BUILD_SUCCESS)
throw std::runtime_error("build failed");

// std::vector<char> log = getInfoVec<char, size_t>(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 (чтобы дальнейшие замеры были ближе к реальности)
Expand All @@ -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%-перцентайль (как и стандартное отклонение)
Expand All @@ -141,34 +153,35 @@ 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 штук
// - Размер каждого элемента sizeof(float)=4 байта
// - Обращений к видеопамяти 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 (и рассчитайте скорость трансфера данных в гигабайтах в секунду)
{
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;
}

0 comments on commit 6bc8404

Please sign in to comment.