Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Task01 Николай Степанов SPbU #18

Closed
wants to merge 2 commits into from
Closed

Task01 Николай Степанов SPbU #18

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
4f5c948
ta-da!
elteammate Sep 18, 2024
f796860
redundancy
elteammate Sep 18, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 5 additions & 1 deletion src/cl/aplusb.cl
View file
Open in desktop
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 float *a, __global float *b, __global float *c, unsigned int n) {
// Узнать, какой workItem выполняется в этом потоке поможет функция get_global_id
// см. в документации https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/
// OpenCL Compiler -> Built-in Functions -> Work-Item Functions
unsigned int i = get_global_id(0);

// P.S. В общем случае количество элементов для сложения может быть некратно размеру WorkGroup, тогда размер рабочего пространства округлен вверх от числа элементов до кратности на размер WorkGroup
// и в таком случае, если сделать обращение к массиву просто по индексу=get_global_id(0), будет undefined behaviour (вплоть до повисания ОС)
// поэтому нужно либо дополнить массив данных длиной до кратности размеру рабочей группы,
// либо сделать return в кернеле до обращения к данным в тех WorkItems, где get_global_id(0) выходит за границы данных (явной проверкой)
if (i < n) {
c[i] = a[i] + b[i];
}
}
161 changes: 141 additions & 20 deletions src/main.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,29 @@ void reportError(cl_int err, const std::string &filename, int line) {

#define OCL_SAFE_CALL(expr) reportError(expr, __FILE__, __LINE__)

template<typename F>
class DeferredCall {
const char *message;
F f;

public:
explicit DeferredCall(const char *message, F f)
: message(message), f(f) {}

~DeferredCall() {
// std::cerr << message << "\n";
f();
}
};

template<typename F>
DeferredCall<F> makeDeferred(const char *message, F f) {
return DeferredCall<F>(message, f);
}

#define token_paste(a, b) a ## b
#define token_paste2(a, b) token_paste(a, b)
#define defer(code) auto token_paste2(_super_magic_deferred_, __LINE__) = makeDeferred(#code, [&](){code;})

int main() {
// Пытаемся слинковаться с символами OpenCL API в runtime (через библиотеку clew)
Expand All @@ -39,6 +62,63 @@ int main() {

// TODO 1 По аналогии с предыдущим заданием узнайте, какие есть устройства, и выберите из них какое-нибудь
// (если в списке устройств есть хоть одна видеокарта - выберите ее, если нету - выбирайте процессор)
cl_uint platform_count = 0;
OCL_SAFE_CALL(clGetPlatformIDs(0, nullptr, &platform_count));
std::vector<cl_platform_id> platforms(platform_count);
OCL_SAFE_CALL(clGetPlatformIDs(platform_count, platforms.data(), nullptr));
std::vector<cl_device_id> cpus;
std::vector<cl_device_id> gpus;
for (int platform_index = 0; platform_index < platform_count; ++platform_index) {
cl_uint device_count = 0;
OCL_SAFE_CALL(clGetDeviceIDs(platforms[platform_index], CL_DEVICE_TYPE_ALL, 0, nullptr, &device_count));

if (device_count == 0)
continue;

std::vector<cl_device_id> devices(device_count);
OCL_SAFE_CALL(clGetDeviceIDs(platforms[platform_index], CL_DEVICE_TYPE_ALL, device_count, devices.data(), nullptr));

for (int device_index = 0; device_index < device_count; ++device_index) {
cl_device_id device = devices[device_index];

cl_bool available;
OCL_SAFE_CALL(clGetDeviceInfo(device, CL_DEVICE_AVAILABLE, sizeof(available), &available, nullptr));
if (!available) {
continue;
}

cl_device_type device_type;
OCL_SAFE_CALL(clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(device_type), &device_type, nullptr));
if (device_type & CL_DEVICE_TYPE_CPU) {
cpus.push_back(device);
}
if (device_type & CL_DEVICE_TYPE_GPU) {
gpus.push_back(device);
}
}
}

std::cout << "Found " << cpus.size() << " CPUs and " << gpus.size() << " GPUs" << std::endl;

cl_device_id device;
if (!gpus.empty()) {
device = gpus[0];
std::cout << "Using GPU: ";
} else if (!cpus.empty()) {
device = cpus[0];
std::cout << "Using CPU: ";
} else {
throw std::runtime_error("No devices found!");
}

size_t device_name_size = 0;
OCL_SAFE_CALL(clGetDeviceInfo(device, CL_DEVICE_NAME, 0, nullptr, &device_name_size));
std::vector<char> device_name(device_name_size);
OCL_SAFE_CALL(clGetDeviceInfo(device, CL_DEVICE_NAME, device_name_size, device_name.data(), nullptr));
std::cout << device_name.data() << std::endl;

cl_platform_id platform;
OCL_SAFE_CALL(clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &platform, nullptr));

// TODO 2 Создайте контекст с выбранным устройством
// См. документацию https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ -> OpenCL Runtime -> Contexts -> clCreateContext
Expand All @@ -47,11 +127,22 @@ int main() {

// Контекст и все остальные ресурсы следует освобождать с помощью clReleaseContext/clReleaseQueue/clReleaseMemObject... (да, не очень RAII, но это лишь пример)

cl_int errcode_ret = 0;
cl_context ctx = clCreateContext(nullptr, 1, &device, [](const char errinfo[], const void *, size_t, void *) {
std::cerr << "OpenCL error: " << errinfo << std::endl;
}, nullptr, &errcode_ret);
OCL_SAFE_CALL(errcode_ret);
defer(clReleaseContext(ctx));
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔥


// TODO 3 Создайте очередь выполняемых команд в рамках выбранного контекста и устройства
// См. документацию https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ -> OpenCL Runtime -> Runtime APIs -> Command Queues -> clCreateCommandQueue
// Убедитесь, что в соответствии с документацией вы создали in-order очередь задач
cl_command_queue queue = clCreateCommandQueue(ctx, device, 0, &errcode_ret);
OCL_SAFE_CALL(errcode_ret);
defer(clReleaseCommandQueue(queue));

unsigned int n = 1000 * 1000;
// unsigned int n = 1000 * 1000;
unsigned int n = 1000 * 1000 * 100;
// Создаем два массива псевдослучайных данных для сложения и массив для будущего хранения результата
std::vector<float> as(n, 0);
std::vector<float> bs(n, 0);
Expand All @@ -63,11 +154,24 @@ int main() {
}
std::cout << "Data generated for n=" << n << "!" << std::endl;

std::size_t buf_size = n * sizeof(float);

// TODO 4 Создайте три буфера в памяти устройства (в случае видеокарты - в видеопамяти - VRAM) - для двух суммируемых массивов as и bs (они read-only) и для массива с результатом cs (он write-only)
// См. Buffer Objects -> clCreateBuffer
// Размер в байтах соответственно можно вычислить через sizeof(float)=4 и тот факт, что чисел в каждом массиве n штук
// Данные в as и bs можно прогрузить этим же методом, скопировав данные из host_ptr=as.data() (и не забыв про битовый флаг, на это указывающий)
// или же через метод Buffer Objects -> clEnqueueWriteBuffer
cl_mem a_buf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, buf_size, as.data(), &errcode_ret);
OCL_SAFE_CALL(errcode_ret);
defer(clReleaseMemObject(a_buf));

cl_mem b_buf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, buf_size, bs.data(), &errcode_ret);
OCL_SAFE_CALL(errcode_ret);
defer(clReleaseMemObject(b_buf));

cl_mem c_buf = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, buf_size, nullptr, &errcode_ret);
OCL_SAFE_CALL(errcode_ret);
defer(clReleaseMemObject(c_buf));

// TODO 6 Выполните TODO 5 (реализуйте кернел в src/cl/aplusb.cl)
// затем убедитесь, что выходит загрузить его с диска (убедитесь что Working directory выставлена правильно - см. описание задания),
Expand All @@ -85,30 +189,43 @@ int main() {
// TODO 7 Создайте OpenCL-подпрограмму с исходниками кернела
// см. Runtime APIs -> Program Objects -> clCreateProgramWithSource
// у string есть метод c_str(), но обратите внимание, что передать вам нужно указатель на указатель
const char *kernel_sources_ptr = kernel_sources.c_str();
cl_program program = clCreateProgramWithSource(ctx, 1, &kernel_sources_ptr, nullptr, &errcode_ret);
OCL_SAFE_CALL(errcode_ret);
defer(clReleaseProgram(program));

// TODO 8 Теперь скомпилируйте программу и напечатайте в консоль лог компиляции
// см. clBuildProgram
cl_int build_error = clBuildProgram(program, 1, &device, nullptr, 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) {
// std::cout << "Log:" << std::endl;
// std::cout << log.data() << std::endl;
// }
if (build_error == CL_BUILD_PROGRAM_FAILURE) {
size_t log_size = 0;
OCL_SAFE_CALL(clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size));
std::vector<char> log(log_size, 0);
OCL_SAFE_CALL(clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, log.data(), nullptr));
std::cout << "Log:" << std::endl;
std::cout << log.data() << std::endl;
return 1;
} else {
OCL_SAFE_CALL(build_error);
}

// TODO 9 Создайте OpenCL-kernel в созданной подпрограмме (в одной подпрограмме может быть несколько кернелов, но в данном случае кернел один)
// см. подходящую функцию в Runtime APIs -> Program Objects -> Kernel Objects
cl_kernel kernel = clCreateKernel(program, "aplusb", &errcode_ret);
OCL_SAFE_CALL(errcode_ret);
defer(clReleaseKernel(kernel));

// 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 *), &a_buf));
OCL_SAFE_CALL(clSetKernelArg(kernel, i++, sizeof(cl_mem *), &b_buf));
OCL_SAFE_CALL(clSetKernelArg(kernel, i++, sizeof(cl_mem *), &c_buf));
OCL_SAFE_CALL(clSetKernelArg(kernel, i++, sizeof(cl_uint), &n));
}

// TODO 11 Выше увеличьте n с 1000*1000 до 100*1000*1000 (чтобы дальнейшие замеры были ближе к реальности)
Expand All @@ -127,6 +244,9 @@ int main() {
for (unsigned int i = 0; i < 20; ++i) {
// clEnqueueNDRangeKernel...
// clWaitForEvents...
cl_event event;
OCL_SAFE_CALL(clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global_work_size, &workGroupSize, 0, nullptr, &event));
OCL_SAFE_CALL(clWaitForEvents(1, &event));
t.nextLap();// При вызове nextLap секундомер запоминает текущий замер (текущий круг) и начинает замерять время следующего круга
}
// Среднее время круга (вычисления кернела) на самом деле считается не по всем замерам, а лишь с 20%-перцентайля по 80%-перцентайль (как и стандартное отклонение)
Expand All @@ -140,34 +260,35 @@ int main() {
// - Флопс - это число операций с плавающей точкой в секунду
// - В гигафлопсе 10^9 флопсов
// - Среднее время выполнения кернела равно t.lapAvg() секунд
std::cout << "GFlops: " << 0 << std::endl;
std::cout << "GFlops: " << (double)n / t.lapAvg() / 1e9 << 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: " << (double)(3 * buf_size) / (1024 * 1024 * 1024) / 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(queue, c_buf, CL_TRUE, 0, buf_size, 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)buf_size / (1024 * 1024 * 1024) / 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;
}
Loading