Skip to content

Task01 Вадим Козлов ИТМО #1042

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

Open
six-nine wants to merge 4 commits into
base: task01
Choose a base branch
from
Open

Task01 Вадим Козлов ИТМО #1042

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
f750cd9
init version
six-nine Feb 22, 2026
7ae95e2
Fix
six-nine Feb 22, 2026
c45a3a0
Another fix
six-nine Feb 22, 2026
205c7ab
Fix perf
six-nine Feb 22, 2026
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
12 changes: 11 additions & 1 deletion src/kernels/cl/aplusb_matrix_bad.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@

#include "../defines.h"

__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
__kernel void aplusb_matrix_bad(__global const uint* a,
__global const uint* b,
__global uint* c,
Expand All @@ -16,5 +17,14 @@ __kernel void aplusb_matrix_bad(__global const uint* a,
// т.е. если в матрице сделать шаг вправо или влево на одну ячейку - то в памяти мы шагнем на 4 байта
// т.е. если в матрице сделать шаг вверх или вниз на одну ячейку - то в памяти мы шагнем на так называемый stride=width*4 байта

// TODO реализуйте этот кернел - просуммируйте две матрицы так чтобы получить максимально ПЛОХУЮ производительность с точки зрения memory coalesced паттерна доступа
const unsigned int id = get_global_id(0);
const unsigned int j = id / height;
const unsigned int i = id % height;

if (i >= height || j >= width)
return;

const unsigned int idx = i * width + j;

c[idx] = a[idx] + b[idx];
}
10 changes: 9 additions & 1 deletion src/kernels/cl/aplusb_matrix_good.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@

#include "../defines.h"

__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
__kernel void aplusb_matrix_good(__global const uint* a,
__global const uint* b,
__global uint* c,
Expand All @@ -16,5 +17,12 @@ __kernel void aplusb_matrix_good(__global const uint* a,
// т.е. если в матрице сделать шаг вправо или влево на одну ячейку - то в памяти мы шагнем на 4 байта
// т.е. если в матрице сделать шаг вверх или вниз на одну ячейку - то в памяти мы шагнем на так называемый stride=width*4 байта

// TODO реализуйте этот кернел - просуммируйте две матрицы так чтобы получить максимально ХОРОШУЮ производительность с точки зрения memory coalesced паттерна доступа
const unsigned int id = get_global_id(0);
const unsigned int i = id / width;
const unsigned int j = id % width;

if (i >= height || j >= width)
return;

c[id] = a[id] + b[id];
}
53 changes: 28 additions & 25 deletions src/main_aplusb_matrix.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,9 +42,6 @@ void run(int argc, char** argv)
unsigned int height = task_size * 128;
std::cout << "matrices size: " << width << "x" << height << " = 3 * " << (sizeof(unsigned int) * width * height / 1024 / 1024) << " MB" << std::endl;

// TODO Удалите эту строку, она для того чтобы моя заготовка (не работающий код) не пыталась запуститься на CI
throw std::runtime_error(CODE_IS_NOT_IMPLEMENTED);

std::vector<unsigned int> as(width * height, 0);
std::vector<unsigned int> bs(width * height, 0);
for (size_t i = 0; i < width * height; ++i) {
Expand All @@ -55,8 +52,9 @@ void run(int argc, char** argv)
// Аллоцируем буферы в VRAM
gpu::gpu_mem_32u a_gpu(width * height), b_gpu(width * height), c_gpu(width * height);

// TODO Удалите этот rassert - прогрузите входные данные по PCI-E шине: CPU RAM -> GPU VRAM
rassert(false, 5462345134123);
// rassert(false, 5462345134123);
a_gpu.writeN(as.data(), width * height);
b_gpu.writeN(bs.data(), width * height);

{
std::cout << "Running BAD matrix kernel..." << std::endl;
Expand All @@ -69,34 +67,21 @@ void run(int argc, char** argv)
// Настраиваем размер рабочего пространства (n) и размер рабочих групп в этом рабочем пространстве (GROUP_SIZE=256)
// Обратите внимание что сейчас указана рабочая группа размера 1х1 в рабочем пространстве width x height, это не то что вы хотите
// TODO И в плохом и в хорошем кернеле рабочая группа обязана состоять из 256 work-items
gpu::WorkSize workSize(1, 1, width, height);
gpu::WorkSize workSize(GROUP_SIZE, width * height);

// Запускаем кернел, с указанием размера рабочего пространства и передачей всех аргументов
// Если хотите - можете удалить ветвление здесь и оставить только тот код который соответствует вашему выбору API
// TODO раскомментируйте вызов вашего API и поправьте его
if (context.type() == gpu::Context::TypeOpenCL) {
// ocl_aplusb_matrix_bad.exec(workSize, a_gpu, ...);
} else if (context.type() == gpu::Context::TypeCUDA) {
// cuda::aplusb_matrix_bad(workSize, a_gpu, ...);
} else if (context.type() == gpu::Context::TypeVulkan) {
struct {
unsigned int width;
unsigned int height;
} params = { width, height };
// vk_aplusb_matrix_bad.exec(params, workSize, a_gpu, ...);
} else {
rassert(false, 4531412341, context.type());
}
ocl_aplusb_matrix_bad.exec(workSize, a_gpu, b_gpu, c_gpu, width, height);

times.push_back(t.elapsed());
}
std::cout << "a + b matrix kernel times (in seconds) - " << stats::valuesStatsLine(times) << std::endl;

// TODO Удалите этот rassert - вычислите достигнутую эффективную пропускную способность видеопамяти
rassert(false, 54623414231);
double memory_size_gb = sizeof(unsigned int) * 3 * width * height / 1024.0 / 1024.0 / 1024.0;
std::cout << "a + b kernel median VRAM bandwidth: " << memory_size_gb / stats::median(times) << " GB/s" << std::endl;

// TODO Считываем результат по PCI-E шине: GPU VRAM -> CPU RAM
std::vector<unsigned int> cs(width * height, 0);
c_gpu.readN(cs.data(), width * height);

// Сверяем результат
for (size_t i = 0; i < width * height; ++i) {
Expand All @@ -107,10 +92,28 @@ void run(int argc, char** argv)
{
std::cout << "Running GOOD matrix kernel..." << std::endl;

// TODO Почти тот же код что с плохим кернелом, но теперь с хорошим, рекомендуется копи-паста
std::vector<double> times;
for (int iter = 0; iter < 10; ++iter) {
timer t;

// Настраиваем размер рабочего пространства (n) и размер рабочих групп в этом рабочем пространстве (GROUP_SIZE=256)
// Обратите внимание что сейчас указана рабочая группа размера 1х1 в рабочем пространстве width x height, это не то что вы хотите
// TODO И в плохом и в хорошем кернеле рабочая группа обязана состоять из 256 work-items
gpu::WorkSize workSize(GROUP_SIZE, width * height);

// Запускаем кернел, с указанием размера рабочего пространства и передачей всех аргументов
// Если хотите - можете удалить ветвление здесь и оставить только тот код который соответствует вашему выбору API
ocl_aplusb_matrix_good.exec(workSize, a_gpu, b_gpu, c_gpu, width, height);

times.push_back(t.elapsed());
}
std::cout << "a + b matrix kernel times (in seconds) - " << stats::valuesStatsLine(times) << std::endl;

double memory_size_gb = sizeof(unsigned int) * 3 * width * height / 1024.0 / 1024.0 / 1024.0;
std::cout << "a + b kernel median VRAM bandwidth: " << memory_size_gb / stats::median(times) << " GB/s" << std::endl;

// TODO Считываем результат по PCI-E шине: GPU VRAM -> CPU RAM
std::vector<unsigned int> cs(width * height, 0);
c_gpu.readN(cs.data(), width * height);

// Сверяем результат
for (size_t i = 0; i < width * height; ++i) {
Expand Down