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..c5fd66eb 100644
--- a/src/cl/bitonic.cl
+++ b/src/cl/bitonic.cl
@@ -1,4 +1,23 @@
-__kernel void bitonic()
-{
+__kernel void bitonic(__global int *const as,
+                      uint const log_top_block_size,
+                      uint const log_cur_block_size) {
+    uint const gid = get_global_id(0);
+    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 << 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) {
+        as[idx1] = right_value;
+        as[idx2] = left_value;
+    }
 }
diff --git a/src/main_bitonic.cpp b/src/main_bitonic.cpp
index 9b508a6c..4753c90a 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<typename T>
@@ -59,7 +60,7 @@ int main(int argc, char **argv) {
     const std::vector<int> cpu_sorted = computeCPU(as);
 
     // remove me
-    return 0;
+
 
     gpu::gpu_mem_32i as_gpu;
     as_gpu.resizeN(n);
@@ -74,6 +75,13 @@ int main(int argc, char **argv) {
             t.restart();// Запускаем секундомер после прогрузки данных, чтобы замерять время работы кернела, а не трансфер данных
 
             /*TODO*/
+            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);
+                }
+            }
 
             t.nextLap();
         }
@@ -90,4 +98,4 @@ int main(int argc, char **argv) {
     }
 
     return 0;
-}
\ No newline at end of file
+}