Task05 Даниил Сухоруков ITMO

.github/workflows/cmake.yml

@@ -10,6 +10,9 @@ jobs:
     runs-on: ubuntu-24.04
 
     steps:
+    - name: Delete huge unnecessary tools folder
+      run: rm -rf /opt/hostedtoolcache
+
     - uses: actions/checkout@v2
 
     - name: Install prerequisites

src/kernels/cl/radix_sort_01_local_counting.cl

@@ -1,18 +1,36 @@
-#ifdef __CLION_IDE__
+    #ifdef __CLION_IDE__
 #include <libgpu/opencl/cl/clion_defines.cl> // This file helps CLion IDE to know what additional functions exists in OpenCL's extended C99
 #endif
 
 #include "helpers/rassert.cl"
 #include "../defines.h"
 
-__attribute__((reqd_work_group_size(1, 1, 1)))
+__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
 __kernel void radix_sort_01_local_counting(
-    // это лишь шаблон! смело меняйте аргументы и используемые буфера! можете сделать даже больше кернелов, если это вызовет затруднения - смело спрашивайте в чате
-    // НЕ ПОДСТРАИВАЙТЕСЬ ПОД СИСТЕМУ! СВЕРНИТЕ С РЕЛЬС!! БУНТ!!! АНТИХАЙП!11!!1
-    __global const uint* buffer1,
-    __global       uint* buffer2,
-    unsigned int a1,
-    unsigned int a2)
+    __global const uint* values,
+    __global       uint* local_counts,
+    unsigned int byte_index,
+    unsigned int n)
 {
-    // TODO
+    const unsigned int local_id = get_local_id(0);
+    const unsigned int group_id = get_group_id(0);
+    const unsigned int global_id = get_global_id(0);
+
+    __local uint local_buckets[RADIX_BUCKET_COUNT];
+
+    if (local_id < RADIX_BUCKET_COUNT) {
+        local_buckets[local_id] = 0;
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (global_id < n) {
+        uint value = values[global_id];
+        uint byte_value = (value >> (byte_index * 8)) & 0xFF;
+        atomic_inc(&local_buckets[byte_value]);
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (local_id < RADIX_BUCKET_COUNT) {
+        local_counts[group_id * RADIX_BUCKET_COUNT + local_id] = local_buckets[local_id];
+    }
 }

src/kernels/cl/radix_sort_02_global_prefixes_scan_sum_reduction.cl

@@ -5,13 +5,19 @@
 #include "helpers/rassert.cl"
 #include "../defines.h"
 
-__attribute__((reqd_work_group_size(1, 1, 1)))
+__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
 __kernel void radix_sort_02_global_prefixes_scan_sum_reduction(
-    // это лишь шаблон! смело меняйте аргументы и используемые буфера! можете сделать даже больше кернелов, если это вызовет затруднения - смело спрашивайте в чате
-    // НЕ ПОДСТРАИВАЙТЕСЬ ПОД СИСТЕМУ! СВЕРНИТЕ С РЕЛЬС!! БУНТ!!! АНТИХАЙП!11!!1
-    __global const uint* buffer1,
-    __global       uint* buffer2,
-    unsigned int a1)
+    __global const uint* local_counts,
+    __global       uint* global_sums,
+    unsigned int num_groups)
 {
-    // TODO
+    const unsigned int bucket = get_global_id(0);
+
+    if (bucket < RADIX_BUCKET_COUNT) {
+        uint sum = 0;
+        for (unsigned int group = 0; group < num_groups; ++group) {
+            sum += local_counts[group * RADIX_BUCKET_COUNT + bucket];
+        }
+        global_sums[bucket] = sum;
+    }
 }

src/kernels/cl/radix_sort_03_global_prefixes_scan_accumulation.cl

@@ -5,14 +5,31 @@
 #include "helpers/rassert.cl"
 #include "../defines.h"
 
-__attribute__((reqd_work_group_size(1, 1, 1)))
+__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
 __kernel void radix_sort_03_global_prefixes_scan_accumulation(
-    // это лишь шаблон! смело меняйте аргументы и используемые буфера! можете сделать даже больше кернелов, если это вызовет затруднения - смело спрашивайте в чате
-    // НЕ ПОДСТРАИВАЙТЕСЬ ПОД СИСТЕМУ! СВЕРНИТЕ С РЕЛЬС!! БУНТ!!! АНТИХАЙП!11!!1
-    __global const uint* buffer1,
-    __global       uint* buffer2,
-    unsigned int a1,
-    unsigned int a2)
+    __global const uint* local_counts,
+    __global const uint* global_sums,
+    __global       uint* local_prefix_sums,
+    unsigned int num_groups)
 {
-    // TODO
+    const unsigned int bucket = get_global_id(0);
+    const unsigned int local_id = get_local_id(0);
+
+    __local uint global_prefix[RADIX_BUCKET_COUNT];
+    if (bucket == 0) {
+        global_prefix[0] = 0;
+        for (unsigned int i = 1; i < RADIX_BUCKET_COUNT; ++i) {
+            global_prefix[i] = global_prefix[i - 1] + global_sums[i - 1];
+        }
+    }
+
+    barrier(CLK_GLOBAL_MEM_FENCE);
+
+    if (local_id < RADIX_BUCKET_COUNT) {
+        uint prefix = global_prefix[local_id];
+        for (unsigned int group = 0; group < num_groups; ++group) {
+            local_prefix_sums[group * RADIX_BUCKET_COUNT + local_id] = prefix;
+            prefix += local_counts[group * RADIX_BUCKET_COUNT + local_id];
+        }
+    }
 }

src/kernels/cl/radix_sort_04_scatter.cl

@@ -5,15 +5,29 @@
 #include "helpers/rassert.cl"
 #include "../defines.h"
 
-__attribute__((reqd_work_group_size(1, 1, 1)))
+__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1)))
 __kernel void radix_sort_04_scatter(
-    // это лишь шаблон! смело меняйте аргументы и используемые буфера! можете сделать даже больше кернелов, если это вызовет затруднения - смело спрашивайте в чате
-    // НЕ ПОДСТРАИВАЙТЕСЬ ПОД СИСТЕМУ! СВЕРНИТЕ С РЕЛЬС!! БУНТ!!! АНТИХАЙП!11!!1
-    __global const uint* buffer1,
-    __global const uint* buffer2,
-                   uint* buffer3,
-    unsigned int a1,
-    unsigned int a2)
+    __global const uint* values,
+    __global const uint* local_prefix_sums,
+    __global       uint* output,
+    unsigned int byte_index,
+    unsigned int n)
 {
-    // TODO
-}
+    const unsigned int global_id = get_global_id(0);
+    const unsigned int group_id = get_group_id(0);
+    const unsigned int local_id = get_local_id(0);
+
+    __local uint local_prefix[RADIX_BUCKET_COUNT];
+
+    if (local_id < RADIX_BUCKET_COUNT) {
+        local_prefix[local_id] = local_prefix_sums[group_id * RADIX_BUCKET_COUNT + local_id];
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (global_id < n) {
+        uint value = values[global_id];
+        uint byte_value = (value >> (byte_index * 8)) & 0xFF;
+        uint position = atomic_inc(&local_prefix[byte_value]);
+        output[position] = value;
+    }
+}

src/kernels/defines.h

@@ -5,6 +5,8 @@
 #define GROUP_SIZE_X 16
 #define GROUP_SIZE_Y 16
 
+#define RADIX_BUCKET_COUNT 256
+
 #define RASSERT_ENABLED 0 // disabled by default, enable for debug by changing 0 to 1, disable before performance evaluation/profiling/commiting
 
 #endif // pragma once

src/main_radix_sort.cpp

@@ -85,37 +85,61 @@ void run(int argc, char** argv)
         std::cout << "CPU std::sort effective RAM bandwidth: " << memory_size_gb / t.elapsed() << " GB/s (" << n / 1000 / 1000 / t.elapsed() << " uint millions/s)" << std::endl;
     }
 
-    // Аллоцируем буферы в VRAM
+    unsigned int num_groups = (n + GROUP_SIZE - 1) / GROUP_SIZE;
+
+    unsigned int local_counts_size = num_groups * RADIX_BUCKET_COUNT;
+    unsigned int global_sums_size = RADIX_BUCKET_COUNT;
+    unsigned int local_prefix_sums_size = num_groups * RADIX_BUCKET_COUNT;
+
     gpu::gpu_mem_32u input_gpu(n);
-    gpu::gpu_mem_32u buffer1_gpu(n), buffer2_gpu(n), buffer3_gpu(n), buffer4_gpu(n); // TODO это просто шаблонка, можете переименовать эти буферы, сделать другого размера/типа, удалить часть, добавить новые
+    gpu::gpu_mem_32u buffer1_gpu(local_counts_size);
+    gpu::gpu_mem_32u buffer2_gpu(global_sums_size);
+    gpu::gpu_mem_32u buffer3_gpu(local_prefix_sums_size);
     gpu::gpu_mem_32u buffer_output_gpu(n);
 
-    // Прогружаем входные данные по PCI-E шине: CPU RAM -> GPU VRAM
     input_gpu.writeN(as.data(), n);
-    // Советую занулить (или еще лучше - заполнить какой-то уникальной константой, например 255) все буферы
-    // В некоторых случаях это ускоряет отладку, но обратите внимание, что fill реализован через копию множества нулей по PCI-E, то есть он очень медленный
-    // Если вам нужно занулять буферы в процессе вычислений - используйте кернел который это сделает (см. кернел fill_buffer_with_zeros)
     buffer1_gpu.fill(255);
     buffer2_gpu.fill(255);
     buffer3_gpu.fill(255);
-    buffer4_gpu.fill(255);
     buffer_output_gpu.fill(255);
 
-    // Запускаем кернел (несколько раз и с замером времени выполнения)
+    gpu::gpu_mem_32u* values_in = &input_gpu;
+    gpu::gpu_mem_32u* values_out = &buffer_output_gpu;
+    gpu::gpu_mem_32u* local_counts = &buffer1_gpu;
+    gpu::gpu_mem_32u* global_sums = &buffer2_gpu;
+    gpu::gpu_mem_32u* local_prefix_sums = &buffer3_gpu;
+
     std::vector<double> times;
     for (int iter = 0; iter < 10; ++iter) { // TODO при отладке запускайте одну итерацию
         timer t;
 
-        // Запускаем кернел, с указанием размера рабочего пространства и передачей всех аргументов
-        // Если хотите - можете удалить ветвление здесь и оставить только тот код который соответствует вашему выбору API
         if (context.type() == gpu::Context::TypeOpenCL) {
-            // TODO
-            throw std::runtime_error(CODE_IS_NOT_IMPLEMENTED);
-            // ocl_fillBufferWithZeros.exec();
-            // ocl_radixSort01LocalCounting.exec();
-            // ocl_radixSort02GlobalPrefixesScanSumReduction.exec();
-            // ocl_radixSort03GlobalPrefixesScanAccumulation.exec();
-            // ocl_radixSort04Scatter.exec();
+            input_gpu.writeN(as.data(), n);
+            values_in = &input_gpu;
+            values_out = &buffer_output_gpu;
+
+            for (unsigned int byte_index = 0; byte_index < 4; ++byte_index) {
+                gpu::WorkSize workSize1(GROUP_SIZE, n);
+                ocl_radixSort01LocalCounting.exec(workSize1, *values_in, *local_counts, byte_index, n);
+
+                gpu::WorkSize workSize2(GROUP_SIZE, RADIX_BUCKET_COUNT);
+                ocl_radixSort02GlobalPrefixesScanSumReduction.exec(workSize2, *local_counts, *global_sums, num_groups);
+
+                gpu::WorkSize workSize3(GROUP_SIZE, RADIX_BUCKET_COUNT);
+                ocl_radixSort03GlobalPrefixesScanAccumulation.exec(workSize3, *local_counts, *global_sums, *local_prefix_sums, num_groups);
+
+                gpu::WorkSize workSize4(GROUP_SIZE, n);
+                ocl_radixSort04Scatter.exec(workSize4, *values_in, *local_prefix_sums, *values_out, byte_index, n);
+
+                std::swap(values_in, values_out);
+            }
+
+            if (values_in != &buffer_output_gpu) {
+                std::vector<unsigned int> temp = values_in->readVector();
+                buffer_output_gpu.writeN(temp.data(), n);
+            }
+
+            input_gpu.writeN(as.data(), n);
         } else if (context.type() == gpu::Context::TypeCUDA) {
             // TODO
             throw std::runtime_error(CODE_IS_NOT_IMPLEMENTED);