Task05 Кудрявцев Федор HSE

src/cl/merge.cl

@@ -5,9 +5,62 @@
 #line 5
 
 
+unsigned int lower_bound(__global const int *array, unsigned int len, int value)
+{
+    unsigned int r = len;
+    unsigned int l = 0;
+
+    while (l < r) {
+        int m = (l + r) / 2;
+        if (array[m] < value) {
+            l = m + 1;
+        } else {
+            r = m;
+        }
+    }
+    return l;
+}
+
+unsigned int upper_bound(__global const int *array, unsigned int len, int value)
+{
+    unsigned int r = len;
+    unsigned int l = 0;
+
+    while (l < r) {
+        int m = (l + r) / 2;
+        if (array[m] <= value) {
+            l = m + 1;
+        } else {
+            r = m;
+        }
+    }
+    return l;
+}
+
 __kernel void merge_global(__global const int *as, __global int *bs, unsigned int block_size)
 {
+    const int i = get_global_id(0);
+    const int value = as[i];
+
+    const int double_block = block_size * 2;
+    const unsigned int block_index = i % block_size;
+    const unsigned int block_start = i - block_index;
+    const unsigned int double_block_index = i % double_block ;
+    const unsigned int double_block_start = i - double_block_index;
+    unsigned int other_block_index = 0;
+
+    __global const int *double_block_start_pointer = as + double_block_start;
+
+    if (double_block_index < block_size ) {
+        __global const int *other = double_block_start_pointer + block_size;
+        other_block_index = lower_bound(other, block_size, value);
+    } else {
+        __global const int *other = double_block_start_pointer;
+        other_block_index = upper_bound(other, block_size, value);
+    }
 
+    bs[double_block_start + block_index + other_block_index] = value;
+    return;
 }
 
 __kernel void calculate_indices(__global const int *as, __global unsigned int *inds, unsigned int block_size)

src/main_merge.cpp

@@ -11,6 +11,7 @@
 #include <stdexcept>
 #include <vector>
 
+const int WORK_GROUP_SIZE = 128;
 const int benchmarkingIters = 10;
 const int benchmarkingItersCPU = 1;
 const unsigned int n = 32 * 1024 * 1024;
@@ -58,8 +59,6 @@ int main(int argc, char **argv) {
 
     const std::vector<int> cpu_sorted = computeCPU(as);
 
-    // remove me for task 5.1
-    return 0;
 
     gpu::gpu_mem_32i as_gpu;
     gpu::gpu_mem_32i bs_gpu;
@@ -75,7 +74,11 @@ int main(int argc, char **argv) {
         for (int iter = 0; iter < benchmarkingIters; ++iter) {
             as_gpu.writeN(as.data(), n);
             t.restart();
-            // TODO
+            for (uint32_t size = 1; size < n; size *= 2) {
+                gpu::WorkSize ws(WORK_GROUP_SIZE, n);
+                merge_global.exec(ws, as_gpu, bs_gpu, size);
+                std::swap(as_gpu, bs_gpu);
+            }
             t.nextLap();
         }
         std::cout << "GPU global: " << t.lapAvg() << "+-" << t.lapStd() << " s" << std::endl;