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 Даниил Ушков ITMO #13

Closed
wants to merge 7 commits into from
Closed

Task01 Даниил Ушков ITMO #13

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
87a883e
choose device
Sep 17, 2024
0734301
write kernel
Sep 17, 2024
1633585
compile and run kernel
Sep 17, 2024
c908862
some fixes
Sep 17, 2024
8ae5bae
free resources
Sep 17, 2024
3867f69
refactoring
Sep 17, 2024
d6588a5
bug fix
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
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,13 +11,19 @@
// - На вход дано три массива float чисел; единственное, чем они отличаются от обычных указателей - модификатором __global, т.к. это глобальная память устройства (видеопамять)
// - Четвертым и последним аргументом должно быть передано количество элементов в каждом массиве (unsigned int, главное, чтобы тип был согласован с типом в соответствующем clSetKernelArg в T0D0 10)

__kernel void aplusb(...) {
__kernel void aplusb(__global float *as, __global float *bs, __global float *cs, 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
size_t id = get_global_id(0);

// P.S. В общем случае количество элементов для сложения может быть некратно размеру WorkGroup, тогда размер рабочего пространства округлен вверх от числа элементов до кратности на размер WorkGroup
// и в таком случае, если сделать обращение к массиву просто по индексу=get_global_id(0), будет undefined behaviour (вплоть до повисания ОС)
// поэтому нужно либо дополнить массив данных длиной до кратности размеру рабочей группы,
// либо сделать return в кернеле до обращения к данным в тех WorkItems, где get_global_id(0) выходит за границы данных (явной проверкой)
if (id > n) {
return;
}

cs[id] = as[id] + bs[id];
}
128 changes: 104 additions & 24 deletions src/main.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,48 @@ void reportError(cl_int err, const std::string &filename, int line) {

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

int chooseDeviceAtPlatform(cl_platform_id platform, cl_device_type type, cl_device_id *device) {
cl_uint devicesCount;
cl_int r = clGetDeviceIDs(platform, type, 0, nullptr, &devicesCount);
if (r == CL_DEVICE_NOT_FOUND) {
return -1;
}
OCL_SAFE_CALL(r);
std::vector<cl_device_id> devices(devicesCount);
OCL_SAFE_CALL(clGetDeviceIDs(platform, type, devicesCount, devices.data(), nullptr));
*device = devices[0];
return 0;
}

cl_device_id chooseDevice() {
cl_uint platformsCount;
OCL_SAFE_CALL(clGetPlatformIDs(0, nullptr, &platformsCount));

if (platformsCount == 0) {
throw std::runtime_error("No platforms found!");
}

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

for (cl_platform_id platform : platforms) {
cl_device_id device;
int r = chooseDeviceAtPlatform(platform, CL_DEVICE_TYPE_GPU, &device);
if (r == 0) {
return device;
}
}

for (cl_platform_id platform : platforms) {
cl_device_id device;
int r = chooseDeviceAtPlatform(platform, CL_DEVICE_TYPE_CPU, &device);
if (r == 0) {
return device;
}
}

throw std::runtime_error("No appropriate devices found!");
}

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

// TODO 1 По аналогии с предыдущим заданием узнайте, какие есть устройства, и выберите из них какое-нибудь
// (если в списке устройств есть хоть одна видеокарта - выберите ее, если нету - выбирайте процессор)
cl_device_id device = chooseDevice();

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

// Контекст и все остальные ресурсы следует освобождать с помощью 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(ctx, device, 0, &rc);
OCL_SAFE_CALL(rc);

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 +116,12 @@ int main() {
// Размер в байтах соответственно можно вычислить через sizeof(float)=4 и тот факт, что чисел в каждом массиве n штук
// Данные в as и bs можно прогрузить этим же методом, скопировав данные из host_ptr=as.data() (и не забыв про битовый флаг, на это указывающий)
// или же через метод Buffer Objects -> clEnqueueWriteBuffer
cl_mem as_gpu = clCreateBuffer(ctx, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_ONLY, n * sizeof(float), as.data(), &rc);
OCL_SAFE_CALL(rc);
cl_mem bs_gpu = clCreateBuffer(ctx, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_ONLY, n * sizeof(float), bs.data(), &rc);
OCL_SAFE_CALL(rc);
cl_mem cs_gpu = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, n * sizeof(float), nullptr, &rc);
OCL_SAFE_CALL(rc);

// TODO 6 Выполните TODO 5 (реализуйте кернел в src/cl/aplusb.cl)
// затем убедитесь, что выходит загрузить его с диска (убедитесь что Working directory выставлена правильно - см. описание задания),
Expand All @@ -79,36 +133,52 @@ 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_c_str = kernel_sources.c_str();
cl_program program = clCreateProgramWithSource(ctx, 1, &kernel_sources_c_str, nullptr, &rc);
OCL_SAFE_CALL(rc);

// TODO 8 Теперь скомпилируйте программу и напечатайте в консоль лог компиляции
// см. clBuildProgram
rc = 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;
// }
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));
if (log_size > 1) {
std::cout << "Log:" << std::endl;
std::cout << log.data() << std::endl;
}
// Проверим результат clBuildProgram
OCL_SAFE_CALL(rc);

// TODO 9 Создайте OpenCL-kernel в созданной подпрограмме (в одной подпрограмме может быть несколько кернелов, но в данном случае кернел один)
// см. подходящую функцию в Runtime APIs -> Program Objects -> Kernel Objects
cl_uint kernelsCount;
OCL_SAFE_CALL(clCreateKernelsInProgram(program, 0, nullptr, &kernelsCount));
if (kernelsCount == 0) {
throw std::runtime_error("No kernels in program!");
}
std::vector<cl_kernel> kernels(kernelsCount);
OCL_SAFE_CALL(clCreateKernelsInProgram(program, kernelsCount, kernels.data(), nullptr));
cl_kernel kernel = kernels[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), &as_gpu));
OCL_SAFE_CALL(clSetKernelArg(kernel, i++, sizeof(cl_mem), &bs_gpu));
OCL_SAFE_CALL(clSetKernelArg(kernel, i++, sizeof(cl_mem), &cs_gpu));
OCL_SAFE_CALL(clSetKernelArg(kernel, i++, sizeof(unsigned int), &n));
}

// TODO 11 Выше увеличьте n с 1000*1000 до 100*1000*1000 (чтобы дальнейшие замеры были ближе к реальности)
Expand All @@ -125,8 +195,9 @@ int main() {
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, &workGroupSize, 0, nullptr, &event));
OCL_SAFE_CALL(clWaitForEvents(1, &event));
t.nextLap();// При вызове nextLap секундомер запоминает текущий замер (текущий круг) и начинает замерять время следующего круга
}
// Среднее время круга (вычисления кернела) на самом деле считается не по всем замерам, а лишь с 20%-перцентайля по 80%-перцентайль (как и стандартное отклонение)
Expand All @@ -140,34 +211,43 @@ int main() {
// - Флопс - это число операций с плавающей точкой в секунду
// - В гигафлопсе 10^9 флопсов
// - Среднее время выполнения кернела равно t.lapAvg() секунд
std::cout << "GFlops: " << 0 << std::endl;
std::cout << "GFlops: " << n / (t.lapAvg() * 1000000000) << 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) / (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...
OCL_SAFE_CALL(clEnqueueReadBuffer(queue, cs_gpu, 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: " << (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!");
}
}

// Освобождаем ресурсы
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseMemObject(cs_gpu);
clReleaseMemObject(bs_gpu);
clReleaseMemObject(as_gpu);
clReleaseCommandQueue(queue);
clReleaseContext(ctx);

return 0;
}
Loading