hw05 - attempt1

src/cl/merge.cl

@@ -1,21 +1,60 @@
 #ifdef __CLION_IDE__
-#include "clion_defines.cl"
+    #include "clion_defines.cl"
 #endif
 
 #line 5
-
-
-__kernel void merge_global(__global const int *as, __global int *bs, unsigned int block_size)
-{
-
+int lower_bound(__global const int* as, int size, int value) {
+    int l = -1;
+    int r = size;
+    while (r - l > 1) {
+        int m = (l + r) / 2;
+        if (as[m] >= value) {
+            r = m;
+        } else {
+            l = m;
+        }
+    }
+    return r;
 }
 
-__kernel void calculate_indices(__global const int *as, __global unsigned int *inds, unsigned int block_size)
-{
+int upper_bound(__global const int* as, int size, int value) {
+    int l = -1;
+    int r = size;
+    while (r - l > 1) {
+        int m = (l + r) / 2;
+        if (as[m] > value) {
+            r = m;
+        } else {
+            l = m;
+        }
+    }
+    return r;
+}
 
+__kernel void merge_global(__global const int *as, __global int *bs, unsigned int block_size) {
+
+    const unsigned  int full_block = 2 * block_size;
+    const unsigned int block = get_global_id(0) / full_block;
+    unsigned int idx = get_global_id(0) % full_block;
+    __global const int *begin = as + block * full_block;
+    __global const int *mid = begin + block_size;
+    __global int *out = bs + block * full_block;
+    int tmp;
+
+    if (idx >= block_size) {
+         tmp = begin[idx];
+        __global int *pos = out + upper_bound(begin, block_size, tmp) + idx - block_size;
+        *pos = tmp;
+    } else {
+        tmp = begin[idx];
+        __global int *pos = out + lower_bound(mid, block_size, tmp) + idx;
+        *pos = tmp;
+    }
 }
 
-__kernel void merge_local(__global const int *as, __global const unsigned int *inds, __global int *bs, unsigned int block_size)
-{
+__kernel void calculate_indices(__global const int *as, __global unsigned int *inds, unsigned int block_size) {
+}
 
+__kernel void merge_local(__global const int *as, __global const unsigned int *inds, __global int *bs,
+                          unsigned int block_size) {
 }

src/main_merge.cpp

@@ -24,6 +24,7 @@ void raiseFail(const T &a, const T &b, std::string message, std::string filename
 }
 
 #define EXPECT_THE_SAME(a, b, message) raiseFail(a, b, message, __FILE__, __LINE__)
+const unsigned int WORK_SIZE = 128;
 
 std::vector<int> computeCPU(const std::vector<int> &as)
 {
@@ -59,7 +60,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 +75,10 @@ int main(int argc, char **argv) {
         for (int iter = 0; iter < benchmarkingIters; ++iter) {
             as_gpu.writeN(as.data(), n);
             t.restart();
-            // TODO
+            for (int block_size = 1; block_size < n; block_size *= 2) {
+                merge_global.exec(gpu::WorkSize(std::min(WORK_SIZE, n), n), as_gpu, bs_gpu, block_size);
+                std::swap(as_gpu, bs_gpu);
+            }
             t.nextLap();
         }
         std::cout << "GPU global: " << t.lapAvg() << "+-" << t.lapStd() << " s" << std::endl;