summaryrefslogtreecommitdiff
path: root/src/solvers/check.cl
diff options
context:
space:
mode:
Diffstat (limited to 'src/solvers/check.cl')
-rw-r--r--src/solvers/check.cl56
1 files changed, 39 insertions, 17 deletions
diff --git a/src/solvers/check.cl b/src/solvers/check.cl
index 0ad187a..2ef6406 100644
--- a/src/solvers/check.cl
+++ b/src/solvers/check.cl
@@ -1,21 +1,43 @@
+typedef ulong u64
+typedef uint u32
+
#pragma OPENCL EXTENSION cl_intel_printf : enable
-__kernel void check(__global unsigned long* permutations, __global int* results,
- unsigned long* instructions, unsigned int n, unsigned int w, unsigned long offset) {
- int id = get_global_id(0);
- unsigned long curr_mask = mask | permutations[id + offset];
+__kernel void check(__global u64* permutations, __global u64* results,
+ u64* instructions, __local u64* i_buffer, u32 n, u32 w, u64 offset) {
+ int gid = get_global_id(0);
+ int wid = get_local_id(0);
+ int gsize = get_global_size(0);
+ int wsize = get_local_size(0);
+ int wesize = get_enqueued_local_size(0);
+
+
+ int times = wsize / wesize;
+ times += !(wsize % wesize);
+ for (int i = 0; i < times; i++) {
+ i_buffer[wid + i] = instructions[wid + i]
+ }
+ work_group_barrier(mem_fence::local);
- unsigned long tmask, sum, stones;
- stones = tmask = sum = 0;
- for (int i = 1; i <= w + 1; i++) {
- if ((curr_mask & (1 << i)) == 0) {
- stones += 1;
- tmask |= 1 << (i - sum);
- sum = i;
- }
- }
- if (tmask == (1 << (n + 1)) - 2 && stones == n) {
- printf("test");
- results[id] = id;
- }
+ u64 result = 0;
+ unsigned long own = permutations[gid + offset];
+ for (int 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)) == 0) {
+ stones += 1;
+ tmask |= 1 << (i - sum);
+ sum = i;
+ }
+ }
+ if (tmask == (1 << (n + 1)) - 2 && stones == n) {
+ result |= 1 << (j & 63);
+ }
+ if (j & !(j & 63)) {
+ results[gid * (wsize / 64) + j / 64] = result;
+ result = 0;
+ }
+ }
}