infmachine 0.1.1

//! This module provides cycle stage executor and its interface.

use crate::*;
use gatenative::cpu_build_exec::*;
use gatenative::opencl_build_exec::*;
use gatenative::*;
use opencl3::command_queue::CommandQueue;
use opencl3::context::Context;
use opencl3::device::Device;
use opencl3::error_codes::ClError;
use opencl3::kernel::{ExecuteKernel, Kernel};
use opencl3::program::Program;
use opencl3::types::cl_ulong;
use rayon::prelude::*;

use std::convert::Infallible;
use std::sync::atomic::{self, AtomicU32};
use std::sync::Arc;

/// Cycle stage executor trait.
///
/// It provides methods to execute stages: 1 (read), 2 (clear before write), 3 (write).
pub trait CycleStage<'a, DR, DW, D, E>
where
    DR: DataReader,
    DW: DataWriter,
    D: DataHolder<'a, DR, DW>,
    E: Executor<'a, DR, DW, D>,
{
    type ErrorType;

    /// Create new cycle state executor. `exec` is circuit executor,
    /// `config` is machine configuration, `env_config` is environment configuration.
    fn new_from_executor(
        exec: &E,
        config: InfParMachineConfig,
        env_config: InfParEnvConfig,
    ) -> Self;
    /// Eexecutes read stage - main memory read.
    /// `buffer` holds internal processor states (`proc_int_data`) for all processors,
    /// `memory` is main memory.
    fn execute_read(&self, buffer: &mut D, memory: &D) -> Result<(), Self::ErrorType>;
    /// Executes clear stage - clear main memory before write.
    /// `buffer` holds internal processor states (`proc_int_data`) for all processors,
    /// `memory` is main memory.
    fn execute_clear(&self, buffer: &mut D, memory: &mut D) -> Result<(), Self::ErrorType>;
    /// Executes write stage - write main memory.
    /// `buffer` holds internal processor states (`proc_int_data`) for all processors,
    /// `memory` is main memory.
    fn execute_write(&self, buffer: &mut D, memory: &mut D) -> Result<(), Self::ErrorType>;
}

// CPU utils for atomics

fn get_atomic_slice(t: &mut [u32]) -> &[AtomicU32] {
    unsafe { &*std::ptr::slice_from_raw_parts(t.as_mut_ptr().cast::<AtomicU32>(), t.len()) }
}

/// CPU implementation for cycle stage executor.
pub struct CPUCycleStage {
    data_config: ProcIntDataConfig,
    max_mem_size_in_cells: u64,
}

impl CPUCycleStage {
    fn new(config: InfParMachineConfig, env_config: InfParEnvConfig) -> Self {
        let max_mem_size = env_config.max_mem_size.unwrap_or(u64::MAX);
        let max_mem_size_in_cells =
            u64::try_from(((max_mem_size as u128) << 3) >> config.cell_len_bits).unwrap();
        let data_config = ProcIntDataConfig::new(config, env_config);
        Self {
            data_config,
            max_mem_size_in_cells,
        }
    }

    #[inline]
    fn get_mem_address_and_set_state(
        &self,
        memory_len_in_cells: u64,
        entry: &[u32],
        mem_addr: &mut [u32],
        global_state: &Arc<AtomicU32>,
    ) -> (u64, bool) {
        self.data_config.mem_address_slice(entry, mem_addr);
        let mem_address_len = self.data_config.mem_address_len() >> 5;
        // handle memory address
        let mut illegal = false;
        let mem_addr_64 = if mem_address_len >= 2 {
            (mem_addr[0] as u64) | ((mem_addr[1] as u64) << 32)
        } else {
            mem_addr[0] as u64
        };
        if mem_addr_64 >= memory_len_in_cells {
            global_state.fetch_or(GLOBAL_STATE_ILLEGAL, atomic::Ordering::SeqCst);
            illegal = true;
        }
        if mem_address_len > 2 {
            if mem_addr[2..].iter().any(|x| *x != 0) {
                global_state.fetch_or(GLOBAL_STATE_ILLEGAL, atomic::Ordering::SeqCst);
                illegal = true;
            }
        }
        (mem_addr_64, illegal)
    }
}

const CYCLE_STAGE_CHUNK_LEN: usize = 8 * 1024;

impl<'a> CycleStage<'a, CPUDataReader<'a>, CPUDataWriter<'a>, CPUDataHolder, CPUExecutor>
    for CPUCycleStage
{
    type ErrorType = Infallible;

    fn new_from_executor(
        _: &CPUExecutor,
        config: InfParMachineConfig,
        env_config: InfParEnvConfig,
    ) -> Self {
        Self::new(config, env_config)
    }

    fn execute_read(
        &self,
        buffer: &mut CPUDataHolder,
        memory: &CPUDataHolder,
    ) -> Result<(), Infallible> {
        let mut buffer = buffer.get_mut();
        let buffer = buffer.get_mut();
        let memory = memory.get();
        let memory = memory.get();
        let dc = &self.data_config;
        let entry_len = dc.len();
        let cell_len_bits = dc.config().cell_len_bits;
        let cell_len = 1 << cell_len_bits;
        let mem_address_len = dc.mem_address_len() >> 5;
        let mem_addr_cell_mask = if cell_len_bits < 5 {
            (1usize << (5 - cell_len_bits)) - 1
        } else {
            0
        };
        let cell_mask = if cell_len_bits < 5 {
            (1u32 << cell_len) - 1
        } else {
            0
        };
        assert_eq!((buffer.len() - 64) % entry_len, 0);
        let global_state = Arc::new(AtomicU32::new(buffer[0]));
        let memory_len_in_cells = std::cmp::min(
            u64::try_from(((memory.len() as u128) << 5) >> cell_len_bits).unwrap(),
            self.max_mem_size_in_cells,
        );
        buffer[64..]
            .chunks_mut(entry_len * CYCLE_STAGE_CHUNK_LEN)
            .par_bridge()
            .for_each(|chunk| {
                let mut mem_addr = vec![0; mem_address_len];
                let empty_cell = vec![0; (cell_len + 31) >> 5];
                for entry in chunk.chunks_mut(entry_len) {
                    if dc.mem_read(entry) {
                        let (mem_addr_64, illegal) = self.get_mem_address_and_set_state(
                            memory_len_in_cells,
                            entry,
                            &mut mem_addr,
                            &global_state,
                        );
                        if !illegal {
                            let mem_addr = usize::try_from(mem_addr_64).unwrap();
                            // read memory
                            if cell_len < 32 {
                                let cell = (memory[mem_addr >> (5 - cell_len_bits)]
                                    >> ((mem_addr & mem_addr_cell_mask) << cell_len_bits))
                                    & cell_mask;
                                dc.set_read_mem_cell_all(entry, &[cell]);
                            } else if cell_len == 32 {
                                dc.set_read_mem_cell_all(entry, &memory[mem_addr..mem_addr + 1]);
                            } else {
                                let mem_addr = mem_addr << (cell_len_bits - 5);
                                let mem_addr_p1 = mem_addr + (1 << (cell_len_bits - 5));
                                dc.set_read_mem_cell_all(entry, &memory[mem_addr..mem_addr_p1]);
                            }
                        }
                    } else {
                        // zeroing if no memory read
                        dc.set_read_mem_cell_all(entry, &empty_cell);
                    }
                }
            });
        buffer[0] = global_state.load(atomic::Ordering::SeqCst);
        Ok(())
    }

    fn execute_clear(
        &self,
        buffer: &mut CPUDataHolder,
        memory: &mut CPUDataHolder,
    ) -> Result<(), Infallible> {
        let mut buffer = buffer.get_mut();
        let buffer = buffer.get_mut();
        let mut memory = memory.get_mut();
        let memory = memory.get_mut();
        let memory = get_atomic_slice(memory);
        let dc = &self.data_config;
        let entry_len = dc.len();
        let cell_len_bits = dc.config().cell_len_bits;
        let cell_len = 1 << cell_len_bits;
        let mem_address_len = dc.mem_address_len() >> 5;
        let mem_addr_cell_mask = if cell_len_bits < 5 {
            (1usize << (5 - cell_len_bits)) - 1
        } else {
            0
        };
        let cell_mask = if cell_len_bits < 5 {
            (1u32 << cell_len) - 1
        } else {
            0
        };
        assert_eq!((buffer.len() - 64) % entry_len, 0);
        let global_state = Arc::new(AtomicU32::new(buffer[0]));
        let memory_len_in_cells = std::cmp::min(
            u64::try_from(((memory.len() as u128) << 5) >> cell_len_bits).unwrap(),
            self.max_mem_size_in_cells,
        );
        buffer[64..]
            .chunks(entry_len * CYCLE_STAGE_CHUNK_LEN)
            .par_bridge()
            .for_each(|chunk| {
                let mut mem_addr = vec![0; mem_address_len];
                for entry in chunk.chunks(entry_len) {
                    if dc.mem_write(entry) {
                        let (mem_addr_64, illegal) = self.get_mem_address_and_set_state(
                            memory_len_in_cells,
                            entry,
                            &mut mem_addr,
                            &global_state,
                        );
                        if !illegal {
                            let mem_addr = usize::try_from(mem_addr_64).unwrap();
                            if cell_len < 32 {
                                memory[mem_addr >> (5 - cell_len_bits)].fetch_and(
                                    !(cell_mask
                                        << ((mem_addr & mem_addr_cell_mask) << cell_len_bits)),
                                    atomic::Ordering::SeqCst,
                                );
                            } else if cell_len == 32 {
                                memory[mem_addr].store(0, atomic::Ordering::SeqCst);
                            } else {
                                let mem_addr = mem_addr << (cell_len_bits - 5);
                                let mem_addr_p1 = mem_addr + (1 << (cell_len_bits - 5));
                                for cell in &memory[mem_addr..mem_addr_p1] {
                                    cell.store(0, atomic::Ordering::SeqCst);
                                }
                            }
                        }
                    }
                }
            });
        buffer[0] = global_state.load(atomic::Ordering::SeqCst);
        Ok(())
    }

    fn execute_write(
        &self,
        buffer: &mut CPUDataHolder,
        memory: &mut CPUDataHolder,
    ) -> Result<(), Infallible> {
        let mut buffer = buffer.get_mut();
        let buffer = buffer.get_mut();
        let mut memory = memory.get_mut();
        let memory = memory.get_mut();
        let memory = get_atomic_slice(memory);
        let dc = &self.data_config;
        let entry_len = dc.len();
        let cell_len_bits = dc.config().cell_len_bits;
        let cell_len = 1 << cell_len_bits;
        let mem_address_len = dc.mem_address_len() >> 5;
        let mem_addr_cell_mask = if cell_len_bits < 5 {
            (1usize << (5 - cell_len_bits)) - 1
        } else {
            0
        };
        assert_eq!((buffer.len() - 64) % entry_len, 0);
        let global_state = Arc::new(AtomicU32::new(buffer[0]));
        let memory_len_in_cells = std::cmp::min(
            u64::try_from(((memory.len() as u128) << 5) >> cell_len_bits).unwrap(),
            self.max_mem_size_in_cells,
        );
        buffer[64..]
            .chunks(entry_len * CYCLE_STAGE_CHUNK_LEN)
            .par_bridge()
            .for_each(|chunk| {
                let mut mem_addr = vec![0; mem_address_len];
                let mut mem_cell = vec![0; (cell_len + 31) >> 5];
                for entry in chunk.chunks(entry_len) {
                    if dc.mem_write(entry) {
                        let (mem_addr_64, illegal) = self.get_mem_address_and_set_state(
                            memory_len_in_cells,
                            entry,
                            &mut mem_addr,
                            &global_state,
                        );
                        if !illegal {
                            let mem_addr = usize::try_from(mem_addr_64).unwrap();
                            dc.write_mem_cell_slice(entry, &mut mem_cell);
                            if cell_len < 32 {
                                memory[mem_addr >> (5 - cell_len_bits)].fetch_or(
                                    mem_cell[0]
                                        << ((mem_addr & mem_addr_cell_mask) << cell_len_bits),
                                    atomic::Ordering::SeqCst,
                                );
                            } else if cell_len == 32 {
                                memory[mem_addr].fetch_or(mem_cell[0], atomic::Ordering::SeqCst);
                            } else {
                                let mem_addr = mem_addr << (cell_len_bits - 5);
                                let mem_addr_p1 = mem_addr + (1 << (cell_len_bits - 5));
                                for (i, cell) in memory[mem_addr..mem_addr_p1].iter().enumerate() {
                                    cell.fetch_or(mem_cell[i], atomic::Ordering::SeqCst);
                                }
                            }
                        }
                    }
                }
            });
        buffer[0] = global_state.load(atomic::Ordering::SeqCst);
        Ok(())
    }
}

const OPENCL_CYCLE_STAGE_KERNELS_CODE: &str = r##"
#define MEM_READ(E) ((((E)[REST_POS]) & (1 << 22)) != 0)
#define MEM_WRITE(E) ((((E)[REST_POS]) & (2 << 22)) != 0)

#if MEM_ADDRESS_LEN == 1
#define MEM_ADDRESS(E, MA, I, MEMORY_LEN_IN_CELLS) { \
    (MA) = (E)[0]; \
    (I) = ((MA) >= (MEMORY_LEN_IN_CELLS)) ? GLOBAL_STATE_ILLEGAL : 0; \
}
#endif
#if MEM_ADDRESS_LEN == 2
#define MEM_ADDRESS(E, MA, I, MEMORY_LEN_IN_CELLS) { \
    (MA) = (((ulong)((E)[0])) | (((ulong)((E)[1])) << 32)); \
    (I) = ((MA) >= (MEMORY_LEN_IN_CELLS)) ? GLOBAL_STATE_ILLEGAL : 0; \
}
#endif
#if MEM_ADDRESS_LEN > 2
#define MEM_ADDRESS(E, MA, I, MEMORY_LEN_IN_CELLS) { \
    uint i; \
    (MA) = (((ulong)((E)[0])) | (((ulong)((E)[1])) << 32)); \
    (I) = ((MA) >= (MEMORY_LEN_IN_CELLS)) ? GLOBAL_STATE_ILLEGAL : 0; \
    for (i = 0; i < MEM_ADDRESS_LEN - 2; i++) \
        (I) |= (((E)[2 + i] != 0) ? GLOBAL_STATE_ILLEGAL : 0); \
}
#endif

#if CELL_LEN_BITS < 5
#define MEM_ADDR_CELL_MASK ((1U << (5 - CELL_LEN_BITS)) - 1)
#define CELL_MASK ((1U << (1U << CELL_LEN_BITS)) - 1)

#define READ_MEM_CELL(E, MEM, MA) { \
    (E)[MEM_CELL_POS] = ((E)[MEM_CELL_POS] & 0xffff0000U) | \
        ((((MEM)[(MA) >> (5 - CELL_LEN_BITS)] >> \
            (((MA) & MEM_ADDR_CELL_MASK) << CELL_LEN_BITS))) & CELL_MASK); \
}
#define CLEAR_READ_MEM_CELL(E) { \
    (E)[MEM_CELL_POS] = ((E)[MEM_CELL_POS] & 0xffff0000U); \
}
#define CLEAR_MEM_CELL(MEM, MA) { \
    atomic_and(&(MEM)[(MA) >> (5 - CELL_LEN_BITS)], ~(CELL_MASK << \
        (((MA) & MEM_ADDR_CELL_MASK) << CELL_LEN_BITS))); \
}
#define WRITE_MEM_CELL(E, MEM, MA) { \
    atomic_or(&((MEM)[(MA) >> (5 - CELL_LEN_BITS)]), \
        ((((E)[MEM_CELL_POS] >> 16) & CELL_MASK) << \
            (((MA) & MEM_ADDR_CELL_MASK) << CELL_LEN_BITS))); \
}
#endif

#if CELL_LEN_BITS == 5
#define READ_MEM_CELL(E, MEM, MA) { (E)[MEM_CELL_POS] = (MEM)[(MA)]; }
#define CLEAR_READ_MEM_CELL(E) { (E)[MEM_CELL_POS] = 0; }
#define CLEAR_MEM_CELL(MEM, MA) { (MEM)[(MA)] = 0; }
#define WRITE_MEM_CELL(E, MEM, MA) { atomic_or(&((MEM)[(MA)]), (E)[WRITE_MEM_CELL_POS]); }
#endif

#if CELL_LEN_BITS == 6
#define READ_MEM_CELL(E, MEM, MA) { \
    (E)[MEM_CELL_POS] = (MEM)[2*(MA)]; \
    (E)[MEM_CELL_POS + 1] = (MEM)[2*(MA) + 1]; \
}
#define CLEAR_READ_MEM_CELL(E) { \
    (E)[MEM_CELL_POS] = 0; \
    (E)[MEM_CELL_POS + 1] = 0; \
}
#define CLEAR_MEM_CELL(MEM, MA) { (MEM)[2*(MA)] = 0; (MEM)[2*(MA) + 1] = 0; }
#define WRITE_MEM_CELL(E, MEM, MA) { \
    atomic_or(&((MEM)[2*(MA)]), (E)[WRITE_MEM_CELL_POS]); \
    atomic_or(&((MEM)[2*(MA) + 1]), (E)[WRITE_MEM_CELL_POS + 1]); \
}
#endif

#if CELL_LEN_BITS > 6
#define READ_MEM_CELL(E, MEM, MA) { \
    size_t i; \
    const size_t tempma = (MA) << (CELL_LEN_BITS - 5); \
    for (i = 0; i < (1U << (CELL_LEN_BITS - 5)); i++) \
        (E)[MEM_CELL_POS + i] = (MEM)[tempma + i]; \
}
#define CLEAR_READ_MEM_CELL(E) { \
    size_t i; \
    for (i = 0; i < (1U << (CELL_LEN_BITS - 5)); i++) \
        (E)[MEM_CELL_POS + i] = 0; \
}
#define CLEAR_MEM_CELL(MEM, MA) { \
    size_t i; \
    const size_t tempma = (MA) << (CELL_LEN_BITS - 5); \
    for (i = 0; i < (1U << (CELL_LEN_BITS - 5)); i++) \
        (MEM)[tempma + i] = 0; \
}
#define WRITE_MEM_CELL(E, MEM, MA) { \
    size_t i; \
    const size_t tempma = (MA) << (CELL_LEN_BITS - 5); \
    for (i = 0; i < (1U << (CELL_LEN_BITS - 5)); i++) \
        atomic_or(&(MEM)[tempma + i], (E)[WRITE_MEM_CELL_POS + i]); \
}
#endif

kernel void cycle_stage_execute_read(ulong n, ulong buf_shift, ulong memory_len_in_cells,
        global uint* buffer, const global uint* mem) {
    const size_t id = get_global_id(0);
    const size_t lid = get_local_id(0);
    if (id >= n) return;
    global uint* entry = buffer + buf_shift + ENTRY_LEN*id;
    uint illegal = 0;
    local uint illegal_loc;
    if (MEM_READ(entry)) {
        size_t ma;
        MEM_ADDRESS(entry, ma, illegal, memory_len_in_cells);
        if (!illegal)
            READ_MEM_CELL(entry, mem, ma);
    } else {
        CLEAR_READ_MEM_CELL(entry);
    }
    if (lid == 0)
        illegal_loc = 0;
    barrier(CLK_LOCAL_MEM_FENCE);
    atomic_or(&illegal_loc, illegal);
    barrier(CLK_LOCAL_MEM_FENCE);
    if (lid == 0)
        atomic_or(buffer, illegal_loc);
}

kernel void cycle_stage_execute_clear(ulong n, ulong buf_shift,  ulong memory_len_in_cells,
        global uint* buffer, global uint* mem) {
    const size_t id = get_global_id(0);
    const size_t lid = get_local_id(0);
    if (id >= n) return;
    const global uint* entry = buffer + buf_shift + ENTRY_LEN*id;
    uint illegal = 0;
    local uint illegal_loc;
    if (MEM_WRITE(entry)) {
        size_t ma;
        MEM_ADDRESS(entry, ma, illegal, memory_len_in_cells);
        if (!illegal)
            CLEAR_MEM_CELL(mem, ma);
    }
    if (lid == 0)
        illegal_loc = 0;
    barrier(CLK_LOCAL_MEM_FENCE);
    atomic_or(&illegal_loc, illegal);
    barrier(CLK_LOCAL_MEM_FENCE);
    if (lid == 0)
        atomic_or(buffer, illegal_loc);
}

kernel void cycle_stage_execute_write(ulong n, ulong buf_shift,  ulong memory_len_in_cells,
        global uint* buffer, global uint* mem) {
    const size_t id = get_global_id(0);
    const size_t lid = get_local_id(0);
    if (id >= n) return;
    const global uint* entry = buffer + buf_shift + ENTRY_LEN*id;
    uint illegal = 0;
    local uint illegal_loc;
    if (MEM_WRITE(entry)) {
        size_t ma;
        MEM_ADDRESS(entry, ma, illegal, memory_len_in_cells);
        if (!illegal)
            WRITE_MEM_CELL(entry, mem, ma);
    }
    if (lid == 0)
        illegal_loc = 0;
    barrier(CLK_LOCAL_MEM_FENCE);
    atomic_or(&illegal_loc, illegal);
    barrier(CLK_LOCAL_MEM_FENCE);
    if (lid == 0)
        atomic_or(buffer, illegal_loc);
}
"##;

/// OpenCL implementation for cycle stage executor.
pub struct OpenCLCycleStage {
    cmd_queue: Arc<CommandQueue>,
    group_len: usize,
    read_kernel: Kernel,
    clear_kernel: Kernel,
    write_kernel: Kernel,
    data_config: ProcIntDataConfig,
    max_mem_size_in_cells: u64,
}

impl OpenCLCycleStage {
    fn new(
        context: Arc<Context>,
        cmd_queue: Arc<CommandQueue>,
        config: InfParMachineConfig,
        env_config: InfParEnvConfig,
    ) -> Self {
        let device = Device::new(context.devices()[0]);
        let group_len = usize::try_from(device.max_work_group_size().unwrap()).unwrap();
        let cell_len_bits = config.cell_len_bits;
        let cell_len = 1 << cell_len_bits;
        let data_config = ProcIntDataConfig::new(config, env_config);
        let defs = format!(
            concat!(
                "-DENTRY_LEN=({}) ",
                "-DREST_POS=({}) ",
                "-DMEM_CELL_POS=({}) ",
                "-DWRITE_MEM_CELL_POS=({}) ",
                "-DMEM_ADDRESS_LEN=({}) ",
                "-DCELL_LEN_BITS=({}) ",
                "-DGLOBAL_STATE_ILLEGAL=({})"
            ),
            data_config.len(),
            data_config.rest_pos(),
            data_config.mem_cell_pos(),
            // this same as previous if cell_len<16
            data_config.mem_cell_pos() + (cell_len >> 5),
            data_config.mem_address_len() >> 5,
            cell_len_bits,
            GLOBAL_STATE_ILLEGAL
        );
        let program =
            Program::create_and_build_from_source(&context, OPENCL_CYCLE_STAGE_KERNELS_CODE, &defs)
                .unwrap();
        let read_kernel = Kernel::create(&program, "cycle_stage_execute_read").unwrap();
        let clear_kernel = Kernel::create(&program, "cycle_stage_execute_clear").unwrap();
        let write_kernel = Kernel::create(&program, "cycle_stage_execute_write").unwrap();

        let max_mem_size = env_config.max_mem_size.unwrap_or(u64::MAX);
        let max_mem_size_in_cells =
            u64::try_from(((max_mem_size as u128) << 3) >> config.cell_len_bits).unwrap();
        Self {
            cmd_queue,
            group_len,
            read_kernel,
            clear_kernel,
            write_kernel,
            data_config,
            max_mem_size_in_cells,
        }
    }

    fn execute_kernel(
        &self,
        buffer: &mut OpenCLDataHolder,
        memory: &OpenCLDataHolder,
        kernel: &Kernel,
    ) -> Result<(), ClError> {
        let entry_len = self.data_config.len();
        let buffer_len = buffer.len();
        let cell_len_bits = self.data_config.config().cell_len_bits;
        assert_eq!((buffer_len - 64) % entry_len, 0);
        let memory_len_in_cells = std::cmp::min(
            u64::try_from(((memory.len() as u128) << 5) >> cell_len_bits).unwrap(),
            self.max_mem_size_in_cells,
        );
        let n = (buffer_len - 64) / entry_len;
        let cl_n = cl_ulong::try_from(n).unwrap();
        let cl_bufshift = cl_ulong::try_from(64).unwrap();
        let cl_memory_len_in_cells = cl_ulong::try_from(memory_len_in_cells).unwrap();
        unsafe {
            ExecuteKernel::new(kernel)
                .set_arg(&cl_n)
                .set_arg(&cl_bufshift)
                .set_arg(&cl_memory_len_in_cells)
                .set_arg(buffer.buffer())
                .set_arg(memory.buffer())
                .set_local_work_size(self.group_len)
                .set_global_work_size(((n + self.group_len - 1) / self.group_len) * self.group_len)
                .enqueue_nd_range(&self.cmd_queue)?;
            self.cmd_queue.finish()?;
        }
        Ok(())
    }
}

impl<'a>
    CycleStage<'a, OpenCLDataReader<'a>, OpenCLDataWriter<'a>, OpenCLDataHolder, OpenCLExecutor>
    for OpenCLCycleStage
{
    type ErrorType = ClError;

    fn new_from_executor(
        exec: &OpenCLExecutor,
        config: InfParMachineConfig,
        env_config: InfParEnvConfig,
    ) -> Self {
        Self::new(
            unsafe { exec.context().clone() },
            unsafe { exec.command_queue().clone() },
            config,
            env_config,
        )
    }

    fn execute_read(
        &self,
        buffer: &mut OpenCLDataHolder,
        memory: &OpenCLDataHolder,
    ) -> Result<(), ClError> {
        self.execute_kernel(buffer, memory, &self.read_kernel)
    }

    fn execute_clear(
        &self,
        buffer: &mut OpenCLDataHolder,
        memory: &mut OpenCLDataHolder,
    ) -> Result<(), ClError> {
        self.execute_kernel(buffer, memory, &self.clear_kernel)
    }

    fn execute_write(
        &self,
        buffer: &mut OpenCLDataHolder,
        memory: &mut OpenCLDataHolder,
    ) -> Result<(), ClError> {
        self.execute_kernel(buffer, memory, &self.write_kernel)
    }
}

#[cfg(test)]
mod tests {
    use super::*;

    use opencl3::device::{get_all_devices, CL_DEVICE_TYPE_GPU};
    use opencl3::memory::CL_MEM_READ_WRITE;
    use opencl3::types::CL_BLOCKING;

    struct CPUDataHolderHelper {}

    impl CPUDataHolderHelper {
        fn new(_: (), data: Vec<u32>) -> CPUDataHolder {
            CPUDataHolder::new(data)
        }
    }

    struct CPUCycleStageHelper {}

    impl CPUCycleStageHelper {
        fn new(_: (), config: InfParMachineConfig, env_config: InfParEnvConfig) -> CPUCycleStage {
            CPUCycleStage::new(config, env_config)
        }
    }

    struct OpenCLDataHolderHelper {}

    impl OpenCLDataHolderHelper {
        fn new(
            (context, cmd_queue): &(Arc<Context>, Arc<CommandQueue>),
            data: Vec<u32>,
        ) -> OpenCLDataHolder {
            let mut dh = OpenCLDataHolder::new(
                data.len(),
                (*context).clone(),
                (*cmd_queue).clone(),
                CL_MEM_READ_WRITE,
            );
            unsafe {
                cmd_queue
                    .enqueue_write_buffer(dh.buffer_mut(), CL_BLOCKING, 0, &data, &[])
                    .unwrap();
            }
            dh
        }
    }

    struct OpenCLCycleStageHelper {}

    impl OpenCLCycleStageHelper {
        fn new(
            (context, cmd_queue): &(Arc<Context>, Arc<CommandQueue>),
            config: InfParMachineConfig,
            env_config: InfParEnvConfig,
        ) -> OpenCLCycleStage {
            OpenCLCycleStage::new((*context).clone(), (*cmd_queue).clone(), config, env_config)
        }
    }

    macro_rules! test_cycle_stage_execute_xxx_macro {
        ($name:ident, $dh_type:ty, $dh_arg_type:ty, $cs_type:ty) => {
            // Form of tests: Test reading, clearing and writing all cells in specified dword
            // in memory. If cell is greater than dword then test 1 cell in aligned address.
            fn $name(arg: $dh_arg_type) {
                let memory_len = 40;
                let memory = <$dh_type>::new(
                    arg,
                    (0..u32::try_from(memory_len).unwrap())
                        .map(|x| x.overflowing_mul(0x34ca7a1).0 ^ 0x2d44891b)
                        .collect::<Vec<_>>(),
                );
                for ma_len in [1, 2] {
                    let env_config = InfParEnvConfig {
                        proc_num: 4,
                        flat_memory: true,
                        max_temp_buffer_len: 24,
                        max_mem_size: Some(if ma_len == 1 { 128 * 4 } else { 1u64 << 48 }),
                    };
                    for clb in 0..=7 {
                        let config = InfParMachineConfig {
                            state_len: 68,
                            data_part_len: 7,
                            cell_len_bits: clb,
                        };
                        let pic = ProcIntDataConfig::new(config, env_config);
                        println!("MemAddrLen: {} {}", pic.mem_address_len() >> 5, clb);
                        let cells_in_dword = if clb <= 5 { 1u32 << (5 - clb) } else { 1 };
                        let cell_mask = if clb < 5 {
                            (1u32 << (1 << clb)) - 1
                        } else {
                            u32::MAX
                        };
                        let cycle_stage = <$cs_type>::new(arg, config, env_config);
                        let dwords_in_cell = if clb >= 5 { 1 << (clb - 5) } else { 1 };
                        let dirty_mem_cell = vec![0xfada; dwords_in_cell];
                        for i in 0..cells_in_dword {
                            let addrs = [
                                (11 * cells_in_dword + i, true),
                                (17 * cells_in_dword + i, false),
                                (35 * cells_in_dword + i, true),
                                (26 * cells_in_dword + i, true),
                                ((memory_len + 37) * cells_in_dword + i, false),
                            ];
                            let mut entries = <$dh_type>::new(
                                arg,
                                std::iter::repeat(0u32)
                                    .take(64)
                                    .chain(
                                        addrs
                                            .into_iter()
                                            .map(|(ma, r)| {
                                                let mut entry = vec![0u32; pic.len()];
                                                // handle scale memory address if cell is greater
                                                // than 32-bit
                                                entry[0] =
                                                    if clb <= 5 { ma } else { ma >> (clb - 5) };
                                                entry[pic.len() - 1] |= u32::from(r) << 22;
                                                pic.set_read_mem_cell_all(
                                                    &mut entry,
                                                    &dirty_mem_cell,
                                                );
                                                entry
                                            })
                                            .flatten(),
                                    )
                                    .collect::<Vec<_>>(),
                            );
                            cycle_stage.execute_read(&mut entries, &memory).unwrap();
                            {
                                let memory = memory.get();
                                let memory = memory.get();
                                let entries = entries.get();
                                let entries = entries.get();
                                for (j, (ma, r)) in addrs.into_iter().enumerate() {
                                    let ma = if clb <= 5 {
                                        (ma >> (5 - clb)) as usize
                                    } else {
                                        // if cell is greater than 32-bit then cut lower bits
                                        // because in entries memory address was divided.
                                        (ma as usize) & !(dwords_in_cell - 1)
                                    };
                                    println!("MALX {} {} {} {} {}", clb, i, j, ma, cell_mask);
                                    let entry =
                                        &entries[64 + pic.len() * j..64 + pic.len() * (j + 1)];
                                    if r {
                                        if clb <= 5 {
                                            assert_eq!(
                                                ((memory[ma] >> (i << clb)) & cell_mask) as u64,
                                                pic.read_mem_cell(entry),
                                                "0: {} {} {}",
                                                clb,
                                                i,
                                                j
                                            );
                                        } else {
                                            assert_eq!(
                                                &memory[ma..ma + dwords_in_cell],
                                                pic.read_mem_cell_all(entry),
                                                "1: {} {} {}",
                                                clb,
                                                i,
                                                j
                                            );
                                        }
                                    } else {
                                        if clb <= 5 {
                                            assert_eq!(
                                                0u64,
                                                pic.read_mem_cell(entry),
                                                "2: {} {} {}",
                                                clb,
                                                i,
                                                j
                                            );
                                        } else {
                                            assert_eq!(
                                                vec![0; dwords_in_cell],
                                                pic.read_mem_cell_all(entry),
                                                "3: {} {} {}",
                                                clb,
                                                i,
                                                j
                                            );
                                        }
                                    }
                                }
                                assert_eq!(0, entries[0], "4: {} {}", clb, i);
                            }

                            //
                            // check execute_clear
                            //
                            let mut memory2 = memory.copy();
                            let memory_exp = {
                                let memory = memory.get();
                                let memory = memory.get();
                                let mut memory_exp = memory.to_vec();
                                for (ma, w) in addrs {
                                    let ma = if clb <= 5 {
                                        (ma >> (5 - clb)) as usize
                                    } else {
                                        // if cell is greater than 32-bit then cut lower bits
                                        // because in entries memory address was divided.
                                        (ma as usize) & !(dwords_in_cell - 1)
                                    };
                                    if w {
                                        // clear
                                        if clb <= 5 {
                                            memory_exp[ma] &= !(cell_mask << (i << clb));
                                        } else {
                                            memory_exp[ma..ma + dwords_in_cell].fill(0);
                                        }
                                    }
                                }
                                memory_exp
                            };
                            let mut entries = <$dh_type>::new(
                                arg,
                                std::iter::repeat(0u32)
                                    .take(64)
                                    .chain(
                                        addrs
                                            .into_iter()
                                            .map(|(ma, w)| {
                                                let mut entry = vec![0u32; pic.len()];
                                                // handle scale memory address if cell is greater
                                                // than 32-bit
                                                entry[0] =
                                                    if clb <= 5 { ma } else { ma >> (clb - 5) };
                                                entry[pic.len() - 1] |= u32::from(w) << (22 + 1);
                                                entry
                                            })
                                            .flatten(),
                                    )
                                    .collect::<Vec<_>>(),
                            );
                            cycle_stage
                                .execute_clear(&mut entries, &mut memory2)
                                .unwrap();
                            {
                                let memory2 = memory2.get();
                                let memory2 = memory2.get();
                                for j in 0..memory.len() {
                                    // println!("memory {} {} {:032b}", i, j, memory2[j]);
                                    assert_eq!(memory_exp[j], memory2[j], "5: {} {} {}", clb, i, j);
                                }
                                let entries = entries.get();
                                let entries = entries.get();
                                assert_eq!(0, entries[0], "6: {} {}", clb, i);
                            }

                            //
                            // execute write
                            //
                            let addr_values = [
                                (11 * cells_in_dword + i, 0xda98231c, true),
                                (17 * cells_in_dword + i, 0xa01934fa, false),
                                (35 * cells_in_dword + i, 0x5c0a9d15, true),
                                (26 * cells_in_dword + i, 0x3ca194e9, true),
                                ((memory_len + 37) * cells_in_dword + i, 0x484f30a6, false),
                            ];
                            let mut memory3 = <$dh_type>::new(arg, vec![0; memory.len()]);
                            let memory_exp = {
                                let mut memory_exp = vec![0; memory.len()];
                                for (ma, cell, w) in addr_values {
                                    let ma = if clb <= 5 {
                                        (ma >> (5 - clb)) as usize
                                    } else {
                                        // if cell is greater than 32-bit then cut lower bits
                                        // because in entries memory address was divided.
                                        (ma as usize) & !(dwords_in_cell - 1)
                                    };
                                    if w {
                                        // write
                                        if clb <= 5 {
                                            memory_exp[ma] |= (cell & cell_mask) << (i << clb);
                                        } else {
                                            for k in 0..dwords_in_cell {
                                                memory_exp[ma + k] =
                                                    cell + u32::try_from(k).unwrap();
                                            }
                                        }
                                    }
                                }
                                memory_exp
                            };
                            let mut entries = <$dh_type>::new(
                                arg,
                                std::iter::repeat(0u32)
                                    .take(64)
                                    .chain(
                                        addr_values
                                            .into_iter()
                                            .map(|(ma, cell, w)| {
                                                let mut entry = vec![0u32; pic.len()];
                                                // handle scale memory address if cell is greater
                                                // than 32-bit
                                                entry[0] =
                                                    if clb <= 5 { ma } else { ma >> (clb - 5) };
                                                entry[pic.len() - 1] |= u32::from(w) << (22 + 1);
                                                if clb <= 4 {
                                                    entry[pic.mem_cell_pos()] =
                                                        (cell & cell_mask) << 16;
                                                } else if clb == 5 {
                                                    entry[pic.mem_cell_pos() + 1] = cell;
                                                } else {
                                                    for k in 0..dwords_in_cell {
                                                        entry[pic.mem_cell_pos()
                                                            + dwords_in_cell
                                                            + k] = cell + u32::try_from(k).unwrap();
                                                    }
                                                }
                                                entry
                                            })
                                            .flatten(),
                                    )
                                    .collect::<Vec<_>>(),
                            );
                            cycle_stage
                                .execute_write(&mut entries, &mut memory3)
                                .unwrap();
                            {
                                let memory3 = memory3.get();
                                let memory3 = memory3.get();
                                for j in 0..memory.len() {
                                    // println!("memory {} {} {:032b}", i, j, memory3[j]);
                                    assert_eq!(memory_exp[j], memory3[j], "7: {} {} {}", clb, i, j);
                                }
                                let entries = entries.get();
                                let entries = entries.get();
                                assert_eq!(0, entries[0], "8: {} {}", clb, i);
                            }

                            //
                            // check illegal
                            //
                            let addrs = [
                                (11 * cells_in_dword + i, true),
                                (26 * cells_in_dword + i, true),
                                (memory_len * cells_in_dword + i, true),
                            ];
                            let mut entries = <$dh_type>::new(
                                arg,
                                std::iter::repeat(0u32)
                                    .take(64)
                                    .chain(
                                        addrs
                                            .into_iter()
                                            .map(|(ma, r)| {
                                                let mut entry = vec![0u32; pic.len()];
                                                // handle scale memory address if cell is greater
                                                // than 32-bit
                                                entry[0] =
                                                    if clb <= 5 { ma } else { ma >> (clb - 5) };
                                                entry[pic.len() - 1] |= u32::from(r) << 22;
                                                entry
                                            })
                                            .flatten(),
                                    )
                                    .collect::<Vec<_>>(),
                            );
                            cycle_stage.execute_read(&mut entries, &memory).unwrap();
                            let entries = entries.get();
                            let entries = entries.get();
                            assert_eq!(GLOBAL_STATE_ILLEGAL, entries[0], "9: {} {}", clb, i);

                            //
                            // check illegal for execute_clear
                            //
                            let addrs = [(memory_len * cells_in_dword + i, true)];
                            let mut entries = <$dh_type>::new(
                                arg,
                                std::iter::repeat(0u32)
                                    .take(64)
                                    .chain(
                                        addrs
                                            .into_iter()
                                            .map(|(ma, w)| {
                                                let mut entry = vec![0u32; pic.len()];
                                                // handle scale memory address if cell is greater
                                                // than  32-bit
                                                entry[0] =
                                                    if clb <= 5 { ma } else { ma >> (clb - 5) };
                                                entry[pic.len() - 1] |= u32::from(w) << (22 + 1);
                                                entry
                                            })
                                            .flatten(),
                                    )
                                    .collect::<Vec<_>>(),
                            );
                            cycle_stage
                                .execute_clear(&mut entries, &mut memory2)
                                .unwrap();
                            let entries = entries.get();
                            let entries = entries.get();
                            assert_eq!(GLOBAL_STATE_ILLEGAL, entries[0], "10: {} {}", clb, i);

                            //
                            // check illegal for execute_clear
                            //
                            let addrs = [(memory_len * cells_in_dword + i, true)];
                            let mut entries = <$dh_type>::new(
                                arg,
                                std::iter::repeat(0u32)
                                    .take(64)
                                    .chain(
                                        addrs
                                            .into_iter()
                                            .map(|(ma, w)| {
                                                let mut entry = vec![0u32; pic.len()];
                                                // handle scale memory address if cell is greater
                                                // than 32-bit
                                                entry[0] =
                                                    if clb <= 5 { ma } else { ma >> (clb - 5) };
                                                entry[pic.len() - 1] |= u32::from(w) << (22 + 1);
                                                entry
                                            })
                                            .flatten(),
                                    )
                                    .collect::<Vec<_>>(),
                            );
                            cycle_stage
                                .execute_write(&mut entries, &mut memory3)
                                .unwrap();
                            let entries = entries.get();
                            let entries = entries.get();
                            assert_eq!(GLOBAL_STATE_ILLEGAL, entries[0], "11: {} {}", clb, i);
                        }
                    }
                }
            }
        };
    }

    test_cycle_stage_execute_xxx_macro!(
        test_cycle_stage_execute_xxx_cpu,
        CPUDataHolderHelper,
        (),
        CPUCycleStageHelper
    );

    test_cycle_stage_execute_xxx_macro!(
        test_cycle_stage_execute_xxx_opencl,
        OpenCLDataHolderHelper,
        &(Arc<Context>, Arc<CommandQueue>),
        OpenCLCycleStageHelper
    );

    #[test]
    fn test_cpu_cycle_stage_execute_xxx() {
        test_cycle_stage_execute_xxx_cpu(());
    }

    #[test]
    fn test_opencl_cycle_stage_execute_xxx() {
        let device = Device::new(*get_all_devices(CL_DEVICE_TYPE_GPU).unwrap().get(0).unwrap());
        let context = Arc::new(Context::from_device(&device).unwrap());
        #[allow(deprecated)]
        let cmd_queue =
            Arc::new(unsafe { CommandQueue::create(&context, device.id(), 0).unwrap() });
        test_cycle_stage_execute_xxx_opencl(&(context, cmd_queue));
    }
}