diff --git a/.github/workflows/cmake.yml b/.github/workflows/cmake.yml index 553bce0a..d8394a38 100644 --- a/.github/workflows/cmake.yml +++ b/.github/workflows/cmake.yml @@ -28,6 +28,10 @@ jobs: - name: Build run: cmake --build ${{github.workspace}}/build --config ${{env.BUILD_TYPE}} - - name: merge + - name: bitonic working-directory: ${{github.workspace}}/build - run: ./merge + run: ./bitonic + + - name: prefix_sum + working-directory: ${{github.workspace}}/build + run: ./prefix_sum diff --git a/CMakeLists.txt b/CMakeLists.txt index 26b389ef..9564aa1e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -20,6 +20,10 @@ endif() # Она считывает все байты из файла src/cl/aplusb.cl (т.е. весь исходный код кернела) и преобразует их в массив байтов в файле src/cl/aplusb_cl.h aplusb_kernel # Обратите внимание что это происходит на этапе компиляции, кроме того необходимо чтобы файл src/cl/aplusb_cl.h был перечислен среди исходников для компиляции при вызове add_executable -convertIntoHeader(src/cl/merge.cl src/cl/merge_cl.h merge_kernel) -add_executable(merge src/main_merge.cpp src/cl/merge_cl.h) -target_link_libraries(merge libclew libgpu libutils) +convertIntoHeader(src/cl/bitonic.cl src/cl/bitonic_cl.h bitonic_kernel) +add_executable(bitonic src/main_bitonic.cpp src/cl/bitonic_cl.h) +target_link_libraries(bitonic libclew libgpu libutils) + +convertIntoHeader(src/cl/prefix_sum.cl src/cl/prefix_sum_cl.h prefix_sum_kernel) +add_executable(prefix_sum src/main_prefix_sum.cpp src/cl/prefix_sum_cl.h) +target_link_libraries(prefix_sum libclew libgpu libutils) diff --git a/README.md b/README.md index acde8fae..c3d2aff7 100644 --- a/README.md +++ b/README.md @@ -2,26 +2,26 @@ [Остальные задания](https://github.com/GPGPUCourse/GPGPUTasks2023/). -# Задание 5. Merge sort +# Задание 6. Bitonic sort, prefix sum -[![Build Status](https://github.com/GPGPUCourse/GPGPUTasks2023/actions/workflows/cmake.yml/badge.svg?branch=task05&event=push)](https://github.com/GPGPUCourse/GPGPUTasks2023/actions/workflows/cmake.yml) +[![Build Status](https://github.com/GPGPUCourse/GPGPUTasks2023/actions/workflows/cmake.yml/badge.svg?branch=task06&event=push)](https://github.com/GPGPUCourse/GPGPUTasks2023/actions/workflows/cmake.yml) 0. Сделать fork проекта -1. Выполнить задание -2. Отправить **Pull-request** с названием ```Task05 <Имя> <Фамилия> <Аффиляция>``` (указав вывод программы при исполнении на вашем компьютере - в тройных кавычках для сохранения форматирования) +1. Выполнить задания 6.1, 6.2 +2. Отправить **Pull-request** с названием ```Task06 <Имя> <Фамилия> <Аффиляция>``` (указав вывод каждой программы при исполнении на вашем компьютере - в тройных кавычках для сохранения форматирования) -**Дедлайн**: 23:59 8 октября. +**Дедлайн**: 23:59 22 октября. -Задание +Задание 6.1. Bitonic sort ========= -**Задание 5.1.** Базовое, за него можно получить 10/10 баллов. +Реализуйте bitonic sort для вещественных чисел -Реализуйте merge sort для вещественных чисел. Более простым вариантом будет реализовать бинпоиск по строчкам и столбцам. +Файлы: ```src/main_bitonic.cpp``` и ```src/cl/bitonic.cl``` -**Задание 5.2**. Альтернативное, на +3 бонусных балла, то есть можно получить 13/10. +Задание 6.2. Prefix sum +========= -Реализуйте merge sort в версии с использованием локальной памяти на всех уровнях, с разбиением на рабочие группы путем запуска поиска по диагональкам. Здесь имеет смысл сразу писать поиск по диагональке, тк он все равно пригодится для разбиения на рабочие группы. -Но приступайте, если имеете хорошую уверенность в своей готовности разобраться с индексацией, она может сожрать:) +Реализуйте prefix sum в модели массового параллелизма -Файлы: ```src/main_merge.cpp``` и ```src/cl/merge.cl``` +Файлы: ```src/main_prefix_sum.cpp``` и ```src/cl/prefix_sum.cl``` diff --git a/src/cl/bitonic.cl b/src/cl/bitonic.cl new file mode 100644 index 00000000..6050ae62 --- /dev/null +++ b/src/cl/bitonic.cl @@ -0,0 +1,3 @@ +__kernel void bitonic(__global float *as) { + // TODO +} diff --git a/src/cl/merge.cl b/src/cl/merge.cl deleted file mode 100644 index 8b137891..00000000 --- a/src/cl/merge.cl +++ /dev/null @@ -1 +0,0 @@ - diff --git a/src/cl/prefix_sum.cl b/src/cl/prefix_sum.cl new file mode 100644 index 00000000..0ffdd02f --- /dev/null +++ b/src/cl/prefix_sum.cl @@ -0,0 +1 @@ +// TODO \ No newline at end of file diff --git a/src/main_merge.cpp b/src/main_bitonic.cpp similarity index 86% rename from src/main_merge.cpp rename to src/main_bitonic.cpp index 207672bc..aef94f94 100644 --- a/src/main_merge.cpp +++ b/src/main_bitonic.cpp @@ -5,7 +5,7 @@ #include // Этот файл будет сгенерирован автоматически в момент сборки - см. convertIntoHeader в CMakeLists.txt:18 -#include "cl/merge_cl.h" +#include "cl/bitonic_cl.h" #include #include @@ -53,22 +53,25 @@ int main(int argc, char **argv) { /* gpu::gpu_mem_32f as_gpu; as_gpu.resizeN(n); + { - ocl::Kernel merge(merge_kernel, merge_kernel_length, "merge"); - merge.compile(); + ocl::Kernel bitonic(bitonic_kernel, bitonic_kernel_length, "bitonic"); + bitonic.compile(); + timer t; for (int iter = 0; iter < benchmarkingIters; ++iter) { as_gpu.writeN(as.data(), n); - t.restart();// Запускаем секундомер после прогрузки данных, чтобы замерять время работы кернела, а не трансфера данных - unsigned int workGroupSize = 128; - unsigned int global_work_size = (n + workGroupSize - 1) / workGroupSize * workGroupSize; - merge.exec(gpu::WorkSize(workGroupSize, global_work_size), as_gpu, n); - t.nextLap(); + + t.restart();// Запускаем секундомер после прогрузки данных, чтобы замерять время работы кернела, а не трансфер данных + + // TODO } std::cout << "GPU: " << t.lapAvg() << "+-" << t.lapStd() << " s" << std::endl; std::cout << "GPU: " << (n / 1000 / 1000) / t.lapAvg() << " millions/s" << std::endl; + as_gpu.readN(as.data(), n); } + // Проверяем корректность результатов for (int i = 0; i < n; ++i) { EXPECT_THE_SAME(as[i], cpu_sorted[i], "GPU results should be equal to CPU results!"); diff --git a/src/main_prefix_sum.cpp b/src/main_prefix_sum.cpp new file mode 100644 index 00000000..2b6b9857 --- /dev/null +++ b/src/main_prefix_sum.cpp @@ -0,0 +1,83 @@ +#include +#include +#include +#include +#include + +// Этот файл будет сгенерирован автоматически в момент сборки - см. convertIntoHeader в CMakeLists.txt:18 +#include "cl/prefix_sum_cl.h" + + +template +void raiseFail(const T &a, const T &b, std::string message, std::string filename, int line) +{ + if (a != b) { + std::cerr << message << " But " << a << " != " << b << ", " << filename << ":" << line << std::endl; + throw std::runtime_error(message); + } +} + +#define EXPECT_THE_SAME(a, b, message) raiseFail(a, b, message, __FILE__, __LINE__) + + +int main(int argc, char **argv) +{ + int benchmarkingIters = 10; + unsigned int max_n = (1 << 24); + + for (unsigned int n = 4096; n <= max_n; n *= 4) { + std::cout << "______________________________________________" << std::endl; + unsigned int values_range = std::min(1023, std::numeric_limits::max() / n); + std::cout << "n=" << n << " values in range: [" << 0 << "; " << values_range << "]" << std::endl; + + std::vector as(n, 0); + FastRandom r(n); + for (int i = 0; i < n; ++i) { + as[i] = r.next(0, values_range); + } + + std::vector bs(n, 0); + { + for (int i = 0; i < n; ++i) { + bs[i] = as[i]; + if (i) { + bs[i] += bs[i-1]; + } + } + } + const std::vector reference_result = bs; + + { + { + std::vector result(n); + for (int i = 0; i < n; ++i) { + result[i] = as[i]; + if (i) { + result[i] += result[i-1]; + } + } + for (int i = 0; i < n; ++i) { + EXPECT_THE_SAME(reference_result[i], result[i], "CPU result should be consistent!"); + } + } + + std::vector result(n); + timer t; + for (int iter = 0; iter < benchmarkingIters; ++iter) { + for (int i = 0; i < n; ++i) { + result[i] = as[i]; + if (i) { + result[i] += result[i-1]; + } + } + t.nextLap(); + } + std::cout << "CPU: " << t.lapAvg() << "+-" << t.lapStd() << " s" << std::endl; + std::cout << "CPU: " << (n / 1000.0 / 1000.0) / t.lapAvg() << " millions/s" << std::endl; + } + + { + // TODO: implement on OpenCL + } + } +}