From 1650906f010574e8810c8b0b98334e22fac5894d Mon Sep 17 00:00:00 2001
From: Dennis Kobert <dennis@kobert.dev>
Date: Sun, 12 Jan 2020 03:47:14 +0100
Subject: Restructuring

---
 src/solvers/gpu/check.cl | 43 +++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 43 insertions(+)
 create mode 100644 src/solvers/gpu/check.cl

(limited to 'src/solvers/gpu/check.cl')

diff --git a/src/solvers/gpu/check.cl b/src/solvers/gpu/check.cl
new file mode 100644
index 0000000..af9bc51
--- /dev/null
+++ b/src/solvers/gpu/check.cl
@@ -0,0 +1,43 @@
+//#pragma OPENCL EXTENSION cl_intel_printf : enable
+typedef ulong u64;
+typedef uint u32;
+
+__kernel void check(__global u64* permutations, __global u64* results, 
+    __global u64* instructions, __local u64* i_buffer, u32 n, u32 w, u64 offset) {
+    int wid = get_local_id(0);
+    i_buffer[wid] = instructions[wid];
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    int gid = get_global_id(0);
+    int gsize = get_global_size(0);
+    int wsize = get_local_size(0);
+
+    u32 w64 = wsize / 64 + !!(wsize & 63);
+    u64 result = 0;
+    u64 own =  permutations[gid + offset];
+    u64 correct_mask = (1 << (n + 1)) - 2;
+
+
+    int j;
+    for (j = 0; j < wsize; j++) {
+        unsigned long curr_mask = i_buffer[j] | own;
+        unsigned long tmask, sum, stones;
+        stones = tmask = sum = 0;
+        for (int i = 1; i <= w + 1; i++) {
+            if (!(curr_mask & (1 << i))) {
+                stones++;
+                tmask |= 1 << (i - sum);
+                sum = i;
+            }
+        }
+        if (tmask == correct_mask &&  stones == n) {
+            result |= 1 << (j & 63);   
+        }
+        if (j % 64 == 0 && j) {
+            results[gid * w64 + j / 64] = result;
+            result = 0;
+        }
+    }
+    results[gid * w64 + j / 64] =  result;
+}
+
-- 
cgit v1.2.3-70-g09d2