diff options
author | Dennis Kobert <dennis@kobert.dev> | 2020-01-12 03:47:14 +0100 |
---|---|---|
committer | Dennis Kobert <dennis@kobert.dev> | 2020-01-12 03:47:14 +0100 |
commit | 1650906f010574e8810c8b0b98334e22fac5894d (patch) | |
tree | fe27a9d727e143353c1fcf0286890d549c443303 /src/solvers/gpu/check.cl | |
parent | 6b6f830f8e6d4c0b0d1328b7b22f810ad039d038 (diff) |
Restructuring
Diffstat (limited to 'src/solvers/gpu/check.cl')
-rw-r--r-- | src/solvers/gpu/check.cl | 43 |
1 files changed, 43 insertions, 0 deletions
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; +} + |