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 Igor Bereza SPbU #21

Closed
wants to merge 2 commits into from
Closed

Task01 Igor Bereza SPbU #21

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
73ecda7
Done task 1.
TheRealLL4 Sep 19, 2024
3eb7770
n is 100 millions, not 1 billion
TheRealLL4 Sep 20, 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
8 changes: 7 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,7 +11,7 @@
// - На вход дано три массива 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
Expand All @@ -20,4 +20,10 @@ __kernel void aplusb(...) {
// и в таком случае, если сделать обращение к массиву просто по индексу=get_global_id(0), будет undefined behaviour (вплоть до повисания ОС)
// поэтому нужно либо дополнить массив данных длиной до кратности размеру рабочей группы,
// либо сделать return в кернеле до обращения к данным в тех WorkItems, где get_global_id(0) выходит за границы данных (явной проверкой)
size_t index = get_global_id(0);
if (index >= n) {
return;
}

c[index] = a[index] + b[index];
}
118 changes: 91 additions & 27 deletions src/main.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,27 +31,57 @@ void reportError(cl_int err, const std::string &filename, int line) {

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


int main() {
// Пытаемся слинковаться с символами OpenCL API в runtime (через библиотеку clew)
if (!ocl_init())
throw std::runtime_error("Can't init OpenCL driver!");

// TODO 1 По аналогии с предыдущим заданием узнайте, какие есть устройства, и выберите из них какое-нибудь
// (если в списке устройств есть хоть одна видеокарта - выберите ее, если нету - выбирайте процессор)
cl_uint platformsCount = 0;
OCL_SAFE_CALL(clGetPlatformIDs(0, nullptr, &platformsCount));

std::vector<cl_platform_id> platforms(platformsCount);
OCL_SAFE_CALL(clGetPlatformIDs(platformsCount, platforms.data(), nullptr));

cl_device_id device_id;
for (int i = 0; i < platforms.size(); i++) {
cl_uint devicesCount = 0;
OCL_SAFE_CALL(clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, nullptr, &devicesCount));

std::vector <cl_device_id> devices(devicesCount, 0);
OCL_SAFE_CALL(clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, devices.size(), devices.data(), nullptr));

for (int j = 0; j < devices.size(); j++) {
cl_device_type device_type;
OCL_SAFE_CALL(clGetDeviceInfo(devices[j], CL_DEVICE_TYPE, sizeof(device_type), &device_type, nullptr));

if (device_type == CL_DEVICE_TYPE_GPU) {
device_id = devices[j];
break;
} else if (device_type == CL_DEVICE_TYPE_CPU) {
device_id = devices[j];
}
}
}

// TODO 2 Создайте контекст с выбранным устройством
// См. документацию https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ -> OpenCL Runtime -> Contexts -> clCreateContext
// Не забывайте проверять все возвращаемые коды на успешность (обратите внимание, что в данном случае метод возвращает
// код по переданному аргументом errcode_ret указателю)
cl_int error_code;
cl_context context = clCreateContext(nullptr, 1, &device_id, nullptr, nullptr, &error_code);
OCL_SAFE_CALL(error_code);

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

// 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(context, device_id, 0, &error_code);
OCL_SAFE_CALL(error_code);

unsigned int n = 1000 * 1000;
unsigned int n = 100 * 1000 * 1000;
// Создаем два массива псевдослучайных данных для сложения и массив для будущего хранения результата
std::vector<float> as(n, 0);
std::vector<float> bs(n, 0);
Expand All @@ -68,6 +98,14 @@ int main() {
// Размер в байтах соответственно можно вычислить через sizeof(float)=4 и тот факт, что чисел в каждом массиве n штук
// Данные в as и bs можно прогрузить этим же методом, скопировав данные из host_ptr=as.data() (и не забыв про битовый флаг, на это указывающий)
// или же через метод Buffer Objects -> clEnqueueWriteBuffer
cl_mem device_as = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * n, as.data(), &error_code);
OCL_SAFE_CALL(error_code);

cl_mem device_bs = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * n, bs.data(), &error_code);
OCL_SAFE_CALL(error_code);

cl_mem device_cs = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * n, nullptr, &error_code);
OCL_SAFE_CALL(error_code);

// TODO 6 Выполните TODO 5 (реализуйте кернел в src/cl/aplusb.cl)
// затем убедитесь, что выходит загрузить его с диска (убедитесь что Working directory выставлена правильно - см. описание задания),
Expand All @@ -77,38 +115,49 @@ int main() {
std::ifstream file("src/cl/aplusb.cl");
kernel_sources = std::string(std::istreambuf_iterator<char>(file), std::istreambuf_iterator<char>());
if (kernel_sources.size() == 0) {
throw std::runtime_error("Empty source file! May be you forgot to configure working directory properly?");
throw std::runtime_error("Empty source file! Maybe 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 *source_text = kernel_sources.data();
size_t source_length = kernel_sources.length();
cl_program shader = clCreateProgramWithSource(context, 1, &source_text, &source_length, &error_code);
OCL_SAFE_CALL(error_code);

// TODO 8 Теперь скомпилируйте программу и напечатайте в консоль лог компиляции
// см. clBuildProgram
clBuildProgram(shader, 1, &device_id, 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;
// }
size_t log_size = 0;
OCL_SAFE_CALL(clGetProgramBuildInfo(shader, device_id, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size));

std::vector<char> log(log_size, 0);
OCL_SAFE_CALL(clGetProgramBuildInfo(shader, device_id, 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
cl_kernel kernel = clCreateKernel(shader, "aplusb", &error_code);
OCL_SAFE_CALL(error_code);

// 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(device_as), &device_as);
clSetKernelArg(kernel, i++, sizeof(device_bs), &device_bs);
clSetKernelArg(kernel, i++, sizeof(device_cs), &device_cs);
clSetKernelArg(kernel, i++, sizeof(n), &n);
}

// TODO 11 Выше увеличьте n с 1000*1000 до 100*1000*1000 (чтобы дальнейшие замеры были ближе к реальности)
Expand All @@ -122,11 +171,13 @@ int main() {
// - Дождаться завершения полунного события - см. в документации подходящий метод среди Event Objects
{
size_t workGroupSize = 128;
size_t global_work_size = (n + workGroupSize - 1) / workGroupSize * workGroupSize;
size_t global_work_size = ((n + workGroupSize - 1) / workGroupSize) * workGroupSize;
timer t;// Это вспомогательный секундомер, он замеряет время своего создания и позволяет усреднять время нескольких замеров
for (unsigned int i = 0; i < 20; ++i) {
// clEnqueueNDRangeKernel...
// clWaitForEvents...
cl_event event;
OCL_SAFE_CALL(clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global_work_size, nullptr, 0, nullptr, &event));
OCL_SAFE_CALL(clWaitForEvents(1, &event));
OCL_SAFE_CALL(clReleaseEvent(event));
t.nextLap();// При вызове nextLap секундомер запоминает текущий замер (текущий круг) и начинает замерять время следующего круга
}
// Среднее время круга (вычисления кернела) на самом деле считается не по всем замерам, а лишь с 20%-перцентайля по 80%-перцентайль (как и стандартное отклонение)
Expand All @@ -140,34 +191,47 @@ int main() {
// - Флопс - это число операций с плавающей точкой в секунду
// - В гигафлопсе 10^9 флопсов
// - Среднее время выполнения кернела равно t.lapAvg() секунд
std::cout << "GFlops: " << 0 << std::endl;
std::cout << "GFlops: " << 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: " << 3 * n * sizeof(float) / (1024 * 1024 * 1024.0 * 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, device_cs, CL_TRUE, 0, n * sizeof(float), 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: " << 3 * n * sizeof(float) / (1024 * 1024 * 1024.0 * 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!");
}
}

OCL_SAFE_CALL(clReleaseKernel(kernel));

OCL_SAFE_CALL(clReleaseProgram(shader));

OCL_SAFE_CALL(clReleaseMemObject(device_as));
OCL_SAFE_CALL(clReleaseMemObject(device_bs));
OCL_SAFE_CALL(clReleaseMemObject(device_cs));

OCL_SAFE_CALL(clReleaseCommandQueue(queue));

OCL_SAFE_CALL(clReleaseContext(context));

return 0;
}

Loading