added task06

.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

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)

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```

src/cl/bitonic.cl

@@ -0,0 +1,3 @@
+__kernel void bitonic(__global float *as) {
+    // TODO
+}

src/cl/prefix_sum.cl

@@ -0,0 +1 @@
+// TODO

src/main_merge.cpp → src/main_bitonic.cpp

@@ -5,7 +5,7 @@
 #include <libutils/timer.h>
 
 // Этот файл будет сгенерирован автоматически в момент сборки - см. convertIntoHeader в CMakeLists.txt:18
-#include "cl/merge_cl.h"
+#include "cl/bitonic_cl.h"
 
 #include <iostream>
 #include <stdexcept>
@@ -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!");

src/main_prefix_sum.cpp

@@ -0,0 +1,83 @@
+#include <libgpu/context.h>
+#include <libgpu/shared_device_buffer.h>
+#include <libutils/misc.h>
+#include <libutils/timer.h>
+#include <libutils/fast_random.h>
+
+// Этот файл будет сгенерирован автоматически в момент сборки - см. convertIntoHeader в CMakeLists.txt:18
+#include "cl/prefix_sum_cl.h"
+
+
+template<typename T>
+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<unsigned int>(1023, std::numeric_limits<int>::max() / n);
+		std::cout << "n=" << n << " values in range: [" << 0 << "; " << values_range << "]" << std::endl;
+
+		std::vector<unsigned int> as(n, 0);
+		FastRandom r(n);
+		for (int i = 0; i < n; ++i) {
+			as[i] = r.next(0, values_range);
+		}
+
+		std::vector<unsigned int> bs(n, 0);
+		{
+			for (int i = 0; i < n; ++i) {
+				bs[i] = as[i];
+				if (i) {
+					bs[i] += bs[i-1];
+				}
+			}
+		}
+		const std::vector<unsigned int> reference_result = bs;
+
+		{
+			{
+				std::vector<unsigned int> 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<unsigned int> 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
+		}
+	}
+}