From fe22472f3cb9298c91e0bbeb1808c82384c2d1b1 Mon Sep 17 00:00:00 2001 From: Evgenii Date: Sat, 14 Dec 2024 00:50:30 +0300 Subject: [PATCH 1/2] add task06 solution --- .gitignore | 1 + src/cl/bitonic.cl | 25 +++++++++++++++++++++++-- src/main_bitonic.cpp | 11 +++++++++-- 3 files changed, 33 insertions(+), 4 deletions(-) diff --git a/.gitignore b/.gitignore index 55308b36..7fb808b9 100644 --- a/.gitignore +++ b/.gitignore @@ -3,3 +3,4 @@ build cmake-build* .vs +CMakeLists.txt.user diff --git a/src/cl/bitonic.cl b/src/cl/bitonic.cl index ecf336b9..80b6a765 100644 --- a/src/cl/bitonic.cl +++ b/src/cl/bitonic.cl @@ -1,4 +1,25 @@ -__kernel void bitonic() +__kernel void bitonic(__global int *const as, + uint const top_block_size, + uint const cur_block_size) { - + uint const gid = get_global_id(0); + + uint const top_block_id = gid / (top_block_size / 2); + uint const cur_step = cur_block_size / 2; + uint const cur_block_id = gid / cur_step; + + bool const is_order_increase = top_block_id % 2 == 0; + uint const idx1 = gid + (cur_block_id * cur_step); + uint const idx2 = idx1 + cur_step; + + int const left_value = as[idx1]; + int const right_value = as[idx2]; + bool const is_lesser = left_value < right_value; + + if ((is_order_increase && !is_lesser) || (!is_order_increase && is_lesser)) + { + as[idx1] = right_value; + as[idx2] = left_value; + } } + diff --git a/src/main_bitonic.cpp b/src/main_bitonic.cpp index 9b508a6c..7750956a 100644 --- a/src/main_bitonic.cpp +++ b/src/main_bitonic.cpp @@ -13,6 +13,7 @@ const int benchmarkingIters = 10; const int benchmarkingItersCPU = 1; +// обязательно степень двойки, или докидывать нулей в массив тогда const unsigned int n = 32 * 1024 * 1024; template @@ -59,7 +60,7 @@ int main(int argc, char **argv) { const std::vector cpu_sorted = computeCPU(as); // remove me - return 0; + gpu::gpu_mem_32i as_gpu; as_gpu.resizeN(n); @@ -74,6 +75,12 @@ int main(int argc, char **argv) { t.restart();// Запускаем секундомер после прогрузки данных, чтобы замерять время работы кернела, а не трансфер данных /*TODO*/ + for (uint top_block_size = 1; top_block_size <= n; top_block_size *= 2) { + for (uint cur_block_size = top_block_size; 1 < cur_block_size; cur_block_size /= 2) { + bitonic.exec(gpu::WorkSize(64, n / 2), as_gpu, top_block_size, cur_block_size); + as_gpu.readN(as.data(), n); + } + } t.nextLap(); } @@ -90,4 +97,4 @@ int main(int argc, char **argv) { } return 0; -} \ No newline at end of file +} From b0b74e842548f8e64e0b8b6a46345bf854fbe7bd Mon Sep 17 00:00:00 2001 From: Evgenii Date: Sat, 14 Dec 2024 02:04:30 +0300 Subject: [PATCH 2/2] add version with binary shifts --- src/cl/bitonic.cl | 22 ++++++++++------------ src/main_bitonic.cpp | 7 ++++--- 2 files changed, 14 insertions(+), 15 deletions(-) diff --git a/src/cl/bitonic.cl b/src/cl/bitonic.cl index 80b6a765..c5fd66eb 100644 --- a/src/cl/bitonic.cl +++ b/src/cl/bitonic.cl @@ -1,25 +1,23 @@ __kernel void bitonic(__global int *const as, - uint const top_block_size, - uint const cur_block_size) -{ + uint const log_top_block_size, + uint const log_cur_block_size) { uint const gid = get_global_id(0); - - uint const top_block_id = gid / (top_block_size / 2); - uint const cur_step = cur_block_size / 2; - uint const cur_block_id = gid / cur_step; + uint const clam_log_cur_block_size = 0 == log_cur_block_size ? 0 : (log_cur_block_size - 1); + + uint const top_block_id = gid >> (log_top_block_size - 1); + uint const cur_step = 1 << clam_log_cur_block_size; + uint const cur_block_id = gid >> clam_log_cur_block_size; bool const is_order_increase = top_block_id % 2 == 0; - uint const idx1 = gid + (cur_block_id * cur_step); + uint const idx1 = gid + (cur_block_id << clam_log_cur_block_size); uint const idx2 = idx1 + cur_step; - + int const left_value = as[idx1]; int const right_value = as[idx2]; bool const is_lesser = left_value < right_value; - if ((is_order_increase && !is_lesser) || (!is_order_increase && is_lesser)) - { + if (is_order_increase != is_lesser) { as[idx1] = right_value; as[idx2] = left_value; } } - diff --git a/src/main_bitonic.cpp b/src/main_bitonic.cpp index 7750956a..4753c90a 100644 --- a/src/main_bitonic.cpp +++ b/src/main_bitonic.cpp @@ -75,9 +75,10 @@ int main(int argc, char **argv) { t.restart();// Запускаем секундомер после прогрузки данных, чтобы замерять время работы кернела, а не трансфер данных /*TODO*/ - for (uint top_block_size = 1; top_block_size <= n; top_block_size *= 2) { - for (uint cur_block_size = top_block_size; 1 < cur_block_size; cur_block_size /= 2) { - bitonic.exec(gpu::WorkSize(64, n / 2), as_gpu, top_block_size, cur_block_size); + for (uint log_top_block_size = 1; (1 << log_top_block_size) <= n; ++log_top_block_size) { + for (uint log_cur_block_size = log_top_block_size; + /* костыль для ансайнда */ (1 << log_cur_block_size) <= n; --log_cur_block_size) { + bitonic.exec(gpu::WorkSize(64, n >> 1), as_gpu, log_top_block_size, log_cur_block_size); as_gpu.readN(as.data(), n); } }