From 1a9d0d91e08d0dee9b48c05ef0ac72cc732a42ab Mon Sep 17 00:00:00 2001 From: Dennis Kobert Date: Mon, 6 Jan 2020 08:33:57 +0100 Subject: Backport to opencl 1.2 --- src/solvers/check.cl | 23 ++++++++++++----------- 1 file changed, 12 insertions(+), 11 deletions(-) (limited to 'src/solvers/check.cl') diff --git a/src/solvers/check.cl b/src/solvers/check.cl index 2ef6406..7ff39df 100644 --- a/src/solvers/check.cl +++ b/src/solvers/check.cl @@ -1,22 +1,23 @@ -typedef ulong u64 -typedef uint u32 +//#pragma OPENCL EXTENSION cl_intel_printf : enable +typedef ulong u64; +typedef uint u32; -#pragma OPENCL EXTENSION cl_intel_printf : enable __kernel void check(__global u64* permutations, __global u64* results, - u64* instructions, __local u64* i_buffer, u32 n, u32 w, u64 offset) { + __local 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 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); + //int times = wsize / wesize; + //times += !(wsize % wesize); + //for (int i = 0; i < times; i++) { + int i = 0; + i_buffer[wid + i] = instructions[wid + i]; + //} + barrier(CLK_LOCAL_MEM_FENCE); u64 result = 0; unsigned long own = permutations[gid + offset]; -- cgit v1.2.3-54-g00ecf