infmachine 0.1.1

use crate::proc_int_data::ProcIntDataConfig;
use crate::*;
use gatenative::*;
use gatesim::*;

use std::fmt::{Debug, Write};
use std::hash::Hash;

// circuit associated code

// requires definitions of:
// HAVE_OPENCL - defined if OPENCL version (set to 1)
// REST_POS - position of rest of data in ProcIntData.
// MEM_CELL_POS - position of read mem cell in ProcIntData.
// WRITE_MEM_CELL_POS - position of write mem cell in ProcIntData.
// DATA_PART_MOVE_DONE_INPUT - circuit input variable for DATA_PART_MOVE_DONE.
// OUTPUT_REST_OUTPUTS - output transform to circuit output's variables that holds rest outputs.
// TEMP_BUFFER_POS - position of temp buffer in ProcIntData.
// DATA_PART_POS_POS - position of data part's positions in ProcIntData.
// POPULATE_READ_MEM_CELL_ALL - routine to populate read mem cell to circuit input:
//    form: v=READ_MEM_CELL_DWORD(entry, 0); INPUT_TRANSFORM_BXX(iXX0,iXX1,iXX2,v); ....
// POPULATE_READ_DATA_PART_ALL - routine to populate read data part to circuit input:
//    form: v=READ_DATA_PART_DWORD(entry, 0); INPUT_TRANSFORM_BXX(iXX0,iXX1,iXX2,v); ....
// AGGREGATE_WRITE_MEM_CELL_ALL - routine to aggregate write mem cell from circuit output:
//    form: OUTPUT_TRANSFORM_BXX(v,oXX0,oXX1,...); SET_WRITE_MEM_CELL_DWORD(entry, 0, v); ...
// AGGREGATE_WRITE_DATA_PART_ALL - routine to aggregate write data_part from circuit output:
//    form: OUTPUT_TRANSFORM_BXX(v,oXX0,oXX1,...); write_dp[k] = v; ...

const CIRCUIT_USER_DEFS_CODE: &str = r##"
#if __OPENCL_VERSION__
#define HAVE_OPENCL (1)
#endif
#ifndef HAVE_OPENCL
typedef uint32_t uint;
typedef uint64_t ulong;
#endif

#define ENTRY_REST_OUTPUT(E) ((E)[(REST_POS)] >> 22)
#define SET_ENTRY_REST_OUTPUT(E, V) { \
    (E)[(REST_POS)] = ((E)[(REST_POS)] & 0x003fffffU) | ((V) << 22); }
#define DATA_PART_READ_VAL(V) (((V) >> 2) & 1)
#define DATA_PART_WRITE_VAL(V) (((V) >> 3) & 1)
#define DATA_PART_MOVE_DIR_VAL(V) (((V) >> 4) & 3)
#define DATA_KIND_VAL(V) (((V) >> 6) & 3)
#define DATA_PART_MOVE_DONE_VAL(V) (((V) >> 9) & 1)
#define SET_DATA_PART_MOVE_DONE_VAL(O, V) { (O) = ((O) & ~(1U<<9)) | (((V) & 1) << 9); }
#define STOP_MACHINE_VAL(V) (((V) >> 8) & 1)

#if CELL_LEN_BITS < 5
#define READ_MEM_CELL_DWORD(E, IDX) \
    (((E)[(MEM_CELL_POS)]) & ((1U << (1U << CELL_LEN_BITS)) - 1U))
#define SET_WRITE_MEM_CELL_DWORD(E, IDX, V) { (E)[(MEM_CELL_POS)] = \
    ((E)[(MEM_CELL_POS)] & 0x0000ffffU) | \
        (((V) & ((1U << (1U << CELL_LEN_BITS)) - 1U)) << 16); }
#endif
#if CELL_LEN_BITS >= 5
#define READ_MEM_CELL_DWORD(E, IDX) (E)[(MEM_CELL_POS) + (IDX)]
#define SET_WRITE_MEM_CELL_DWORD(E, IDX, V) { (E)[(WRITE_MEM_CELL_POS) + (IDX)] = (V); }
#endif

#define READ_DATA_PART_DWORD(E, IDX) (E)[(DATA_PART_POS) + (IDX)]
#define READ_DATA_PART_DWORD_LAST(E) ((E)[(DATA_PART_POS) + (DATA_PART_LEN >> 5)] & \
    ((1U << (DATA_PART_LEN & 31)) - 1U))

#if (DATA_PART_LEN & 31) != 0
#define SET_READ_DATA_PART(E, V) { \
    uint k; \
    for(k = 0; k < (DATA_PART_LEN >> 5); k++) \
        (E)[(DATA_PART_POS) + k] = (V)[k]; \
    if ((DATA_PART_LEN & 31) != 0) { \
        size_t pos = (DATA_PART_POS) + (DATA_PART_LEN >> 5); \
        (E)[pos] = ((E)[pos] & ~((1U << (DATA_PART_LEN & 31)) - 1U)) \
            | ((V)[(DATA_PART_LEN >> 5)] & ((1U << (DATA_PART_LEN & 31)) - 1U)); \
    } \
}
#define CLEAR_READ_DATA_PART(E) { \
    uint k; \
    for(k = 0; k < (DATA_PART_LEN >> 5); k++) \
        (E)[(DATA_PART_POS) + k] = 0; \
    if ((DATA_PART_LEN & 31) != 0) { \
        size_t pos = (DATA_PART_POS) + (DATA_PART_LEN >> 5); \
        (E)[pos] = ((E)[pos] & ~((1U << (DATA_PART_LEN & 31)) - 1U)); \
    } \
}
#else
#define SET_READ_DATA_PART(E, V) { \
    uint k; \
    for(k = 0; k < (DATA_PART_LEN >> 5); k++) \
        (E)[(DATA_PART_POS) + k] = (V)[k]; \
}
#define CLEAR_READ_DATA_PART(E) { \
    uint k; \
    for(k = 0; k < (DATA_PART_LEN >> 5); k++) \
        (E)[(DATA_PART_POS) + k] = 0; \
}
#endif

#define MEM_ADDRESS_POS(E) ((E)[(DATA_PART_POS_POS)] & 0xff)
#define TEMP_BUFFER_POS(E) (((E)[(DATA_PART_POS_POS)] >> 8) & 0xffff)
#define PROC_ID_POS(E) (((E)[(DATA_PART_POS_POS)] >> 24) & 0xff)
#define SET_MEM_ADDRESS_POS(E, V) { (E)[(DATA_PART_POS_POS)] = \
    ((E)[(DATA_PART_POS_POS)] & 0xffffff00U) | ((V) & 0xff); }
#define SET_TEMP_BUFFER_POS(E, V) { (E)[(DATA_PART_POS_POS)] = \
    ((E)[(DATA_PART_POS_POS)] & 0xff0000ffU) | (((V) & 0xffff) << 8); }
#define SET_PROC_ID_POS(E, V) { (E)[(DATA_PART_POS_POS)] = \
    ((E)[(DATA_PART_POS_POS)] & 0x00ffffffU) | (((V) & 0xff) << 24); }

// data part routines

#ifdef HAVE_OPENCL
#define DP_INLINE inline
#define DP_PRIV_MOD private
#define DP_DATA_MOD global
#define DP_DP_MOD private
#else
#define DP_INLINE inline
#define DP_PRIV_MOD
#define DP_DATA_MOD
#define DP_DP_MOD
#endif

#if (DATA_PART_LEN == 1) || (DATA_PART_LEN == 2) || (DATA_PART_LEN == 4) || \
        (DATA_PART_LEN == 8) || (DATA_PART_LEN == 16)
#if DATA_PART_LEN == 1
#  define DATA_PART_LEN_BITS (0)
#endif
#if DATA_PART_LEN == 2
#  define DATA_PART_LEN_BITS (1)
#endif
#if DATA_PART_LEN == 4
#  define DATA_PART_LEN_BITS (2)
#endif
#if DATA_PART_LEN == 8
#  define DATA_PART_LEN_BITS (3)
#endif
#if DATA_PART_LEN == 16
#  define DATA_PART_LEN_BITS (4)
#endif
#define DATA_PART_STUFF_DEFINED 1
#define DATA_PART_LEN_MASK ((1U << (5 - DATA_PART_LEN_BITS)) - 1U)

// data part routines for DATA_PART_LEN=2**n.
static DP_INLINE void read_data_part_from(const DP_DATA_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    size_t temp_pos = pos >> (5 - DATA_PART_LEN_BITS);
    if (temp_pos < blen)
        val[0] = (bdata[temp_pos] >> ((pos & DATA_PART_LEN_MASK) << DATA_PART_LEN_BITS)) &
                ((1U << DATA_PART_LEN) - 1U);
    else
        val[0] = 0;
}
static DP_INLINE void read_data_part_from_proc_id(const DP_PRIV_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    size_t temp_pos = pos >> (5 - DATA_PART_LEN_BITS);
    if (temp_pos < blen)
        val[0] = (bdata[temp_pos] >> ((pos & DATA_PART_LEN_MASK) << DATA_PART_LEN_BITS)) &
                ((1U << DATA_PART_LEN) - 1U);
    else
        val[0] = 0;
}
static DP_INLINE uint write_data_part_to(DP_DATA_MOD uint* bdata, size_t bitblen,
        const uint* val, size_t pos) {
    size_t temp_pos = pos >> (5 - DATA_PART_LEN_BITS);
    size_t blen = (bitblen + 31) >> 5;
    uint state = 0;
    uint vx = val[0] & ((1U << DATA_PART_LEN) - 1U);
    if ((pos << DATA_PART_LEN_BITS) < bitblen) {
        uint end = 0;
        uint mask = ((1U << DATA_PART_LEN) - 1U);
        if (((pos + 1) << DATA_PART_LEN_BITS) > bitblen) {
            mask = (1U << (bitblen - (pos << DATA_PART_LEN_BITS))) - 1U;
            end = 1;
        }
        if (!end || (vx & ~mask) == 0) {
            uint dword_pos = (pos & DATA_PART_LEN_MASK) << DATA_PART_LEN_BITS;
            mask <<= dword_pos;
            bdata[temp_pos] = (bdata[temp_pos] & ~mask) | ((vx << dword_pos) & mask);
        } else
            state |= GLOBAL_STATE_ILLEGAL;
    } else if (vx != 0)
        state |= GLOBAL_STATE_ILLEGAL;
    return state;
}
#endif

#if !defined(DATA_PART_STUFF_DEFINED) && DATA_PART_LEN < 32

// data part routines for DATA_PART_LEN < 32
static DP_INLINE void read_data_part_from(const DP_DATA_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    size_t temp_pos = (pos * DATA_PART_LEN) >> 5;
    if (temp_pos < blen) {
        uint dword_pos = (pos * DATA_PART_LEN) & 31;
        val[0] = (bdata[temp_pos] >> dword_pos) &
            ((1U << DATA_PART_LEN) - 1U);
        if (dword_pos + DATA_PART_LEN > 32 && temp_pos + 1 < blen)
            val[0] |= (bdata[temp_pos + 1] << (32 - dword_pos)) &
                ((1U << DATA_PART_LEN) - 1U);
    } else
        val[0] = 0;
}
static DP_INLINE void read_data_part_from_proc_id(const DP_PRIV_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    size_t temp_pos = (pos * DATA_PART_LEN) >> 5;
    if (temp_pos < blen) {
        uint dword_pos = (pos * DATA_PART_LEN) & 31;
        val[0] = (bdata[temp_pos] >> dword_pos) &
            ((1U << DATA_PART_LEN) - 1U);
        if (dword_pos + DATA_PART_LEN > 32 && temp_pos + 1 < blen)
            val[0] |= (bdata[temp_pos + 1] << (32 - dword_pos)) &
                ((1U << DATA_PART_LEN) - 1U);
    } else
        val[0] = 0;
}
static DP_INLINE uint write_data_part_to(DP_DATA_MOD uint* bdata, size_t bitblen,
        const uint* val, size_t pos) {
    size_t temp_pos = (pos * DATA_PART_LEN) >> 5;
    size_t blen = (bitblen + 31) >> 5;
    uint state = 0;
    uint vx = val[0] & ((1U << DATA_PART_LEN) - 1U);
    if ((pos * DATA_PART_LEN) < bitblen) {
        uint end = 0;
        uint mask = ((1U << DATA_PART_LEN) - 1U);
        if (((pos + 1) * DATA_PART_LEN) > bitblen) {
            mask = (1U << (bitblen - (pos * DATA_PART_LEN))) - 1U;
            end = 1;
        }
        if (!end || (vx & ~mask) == 0) {
            uint dword_pos = (pos * DATA_PART_LEN) & 31;
            uint mask1 = mask << dword_pos;
            bdata[temp_pos] = (bdata[temp_pos] & ~mask1) | ((vx << dword_pos) & mask1);
            if (dword_pos + DATA_PART_LEN > 32 && temp_pos + 1 < blen) {
                mask = mask >> (32 - dword_pos);
                bdata[temp_pos + 1] = (bdata[temp_pos + 1] & ~mask) |
                    ((vx >> (32 - dword_pos)) & mask);
            }
        } else
            state |= GLOBAL_STATE_ILLEGAL;
    } else if (vx != 0)
        state |= GLOBAL_STATE_ILLEGAL;
    return state;
}
#endif

#if DATA_PART_LEN == 32

// data part routines for DATA_PART_LEN == 32 (dword)
static DP_INLINE void read_data_part_from(const DP_DATA_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    if (pos < blen) {
        val[0] = bdata[pos];
    } else
        val[0] = 0;
}
static DP_INLINE void read_data_part_from_proc_id(const DP_PRIV_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    if (pos < blen) {
        val[0] = bdata[pos];
    } else
        val[0] = 0;
}
static DP_INLINE uint write_data_part_to(DP_DATA_MOD uint* bdata, size_t bitblen,
        const uint* val, size_t pos) {
    size_t blen = (bitblen + 31) >> 5;
    uint state = 0;
    if (pos < blen) {
        if (((pos + 1) << 5) > bitblen) {
            uint mask = (1U << (bitblen - (pos << 5))) - 1U;
            if ((val[0] & ~mask) == 0)
                bdata[pos] = (bdata[pos] & ~mask) | ((val[0]) & mask);
            else
                state |= GLOBAL_STATE_ILLEGAL;
        } else
            bdata[pos] = val[0];
    } else if (val[0] != 0)
        state |= GLOBAL_STATE_ILLEGAL;
    return state;
}
#endif

#if (DATA_PART_LEN > 32) 
#  if (DATA_PART_LEN & 31) == 0
#    if DATA_PART_LEN == 64

// data part routines for DATA_PART_LEN == 64 (double dword)
static DP_INLINE void read_data_part_from(const DP_DATA_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    uint ppos = 2*pos;
    if (ppos < blen) {
        val[0] = bdata[ppos];
        if (ppos + 1 < blen)
            val[1] = bdata[ppos + 1];
        else
            val[1] = 0;
    } else
        val[0] = val[1] = 0;
}
static DP_INLINE void read_data_part_from_proc_id(const DP_PRIV_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    uint ppos = 2*pos;
    if (ppos < blen) {
        val[0] = bdata[ppos];
        if (ppos + 1 < blen)
            val[1] = bdata[ppos + 1];
        else
            val[1] = 0;
    } else
        val[0] = val[1] = 0;
}
static DP_INLINE uint write_data_part_to(DP_DATA_MOD uint* bdata, size_t bitblen,
        const uint* val, size_t pos) {
    size_t blen = (bitblen + 31) >> 5;
    uint ppos = 2*pos;
    uint state = 0;
    if (ppos < blen) {
        if (bitblen - (ppos<<5) < 32) {
            uint mask = (1U << (bitblen - (ppos<<5))) - 1U;
            if ((val[0] & ~mask) == 0) {
                bdata[ppos] = (bdata[ppos] & ~mask) | (val[0] & mask);
            } else
                state |= GLOBAL_STATE_ILLEGAL;
        }  else {
            uint do_it = 1;
            uint mask = 0xffffffffU;
            if (ppos + 1 < blen) {
                if (bitblen - (ppos<<5) < 64) {
                    mask = (1U << (bitblen - (ppos<<5) - 32)) - 1U;
                    if ((val[1] & ~mask) != 0) {
                        do_it = 0;
                        state |= GLOBAL_STATE_ILLEGAL;
                    }
                }
            } else if (val[1] != 0) {
                do_it = 0;
                state |= GLOBAL_STATE_ILLEGAL;
            }
            if (do_it) {
                bdata[ppos] = val[0];
                if (ppos + 1 < blen) {
                    bdata[ppos + 1] = (bdata[ppos + 1] & ~mask) | (val[1] & mask);
                } else if (val[1] != 0)
                    state |= GLOBAL_STATE_ILLEGAL;
            }
        }
    } else if (val[0] != 0 || val[1] != 0)
        state |= GLOBAL_STATE_ILLEGAL;
    return state;
}
#    else

// data part routines for DATA_PART_LEN = 32*n.
static DP_INLINE void read_data_part_from(const DP_DATA_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    size_t i;
    size_t ppos = (pos * (DATA_PART_LEN>>5));
    for (i = 0; i < (DATA_PART_LEN >> 5); i++) {
        if (ppos + i < blen)
            val[i] = bdata[ppos + i];
        else
            val[i] = 0;
    }
}
static DP_INLINE void read_data_part_from_proc_id(const DP_PRIV_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    size_t i;
    size_t ppos = (pos * (DATA_PART_LEN>>5));
    for (i = 0; i < (DATA_PART_LEN >> 5); i++) {
        if (ppos + i < blen)
            val[i] = bdata[ppos + i];
        else
            val[i] = 0;
    }
}
static DP_INLINE uint write_data_part_to(DP_DATA_MOD uint* bdata, size_t bitblen,
        const uint* val, size_t pos) {
    size_t blen = (bitblen + 31) >> 5;
    size_t i;
    uint do_it = 1;
    size_t ppos = (pos * (DATA_PART_LEN>>5));
    uint state = 0;
    if (ppos < blen) {
        if (ppos + (DATA_PART_LEN>>5) >= blen) {
            uint after_end = (bitblen>>5) - ppos;
            if ((bitblen & 31) != 0) {
                if ((val[after_end] & ~((1U << (bitblen & 31)) - 1U)) != 0) {
                    do_it = 0;
                    state |= GLOBAL_STATE_ILLEGAL;
                }
                after_end++;
            }
            for (i = after_end; i < (DATA_PART_LEN >> 5); i++) {
                if (val[i] != 0) {
                    do_it = 0;
                    state |= GLOBAL_STATE_ILLEGAL;
                }
            }
        }
    } else {
        do_it = 0;
        for (i = 0; i < (DATA_PART_LEN >> 5); i++) {
            if (val[i] != 0) {
                state |= GLOBAL_STATE_ILLEGAL;
                break;
            }
        }
    }
    if (do_it)
        for (i = 0; i < (DATA_PART_LEN >> 5); i++) {
            if (ppos + i + 1 == blen && (bitblen & 31) != 0)
                bdata[ppos + i] = val[i] & ((1U << (bitblen & 31)) - 1U);
            else if (ppos + i < blen)
                bdata[ppos + i] = val[i];
        }
    return state;
}
#    endif
#  else

// data part routines for DATA_PART_LEN > 32
static DP_INLINE void read_data_part_from(const DP_DATA_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    size_t i;
    size_t bitpos_all = ((pos * (DATA_PART_LEN)));
    size_t ppos = bitpos_all >> 5;
    uint bitposp = bitpos_all & 31;
    for (i = 0; i < ((DATA_PART_LEN + 31) >> 5); i++) {
        if (ppos + i < blen) {
            val[i] = bdata[ppos + i] >> bitposp;
            if (bitposp != 0 && ppos + i + 1 < blen &&
                bitposp + DATA_PART_LEN - (i<<5) > 32)
                val[i] |= (bdata[ppos + i + 1] << (32 - bitposp));
        } else
            val[i] = 0;
    }
    if ((DATA_PART_LEN & 31) != 0)
        val[(DATA_PART_LEN >> 5)] &= (1U << (DATA_PART_LEN & 31)) - 1U;
}
static DP_INLINE void read_data_part_from_proc_id(const DP_PRIV_MOD uint* bdata, size_t blen,
        DP_DP_MOD uint* val, size_t pos) {
    size_t i;
    size_t bitpos_all = ((pos * (DATA_PART_LEN)));
    size_t ppos = bitpos_all >> 5;
    uint bitposp = bitpos_all & 31;
    for (i = 0; i < ((DATA_PART_LEN + 31) >> 5); i++) {
        if (ppos + i < blen) {
            val[i] = bdata[ppos + i] >> bitposp;
            if (bitposp != 0 && ppos + i + 1 < blen &&
                bitposp + DATA_PART_LEN - (i<<5) > 32)
                val[i] |= (bdata[ppos + i + 1] << (32 - bitposp));
        } else
            val[i] = 0;
    }
    if ((DATA_PART_LEN & 31) != 0)
        val[(DATA_PART_LEN >> 5)] &= (1U << (DATA_PART_LEN & 31)) - 1U;
}
static DP_INLINE uint write_data_part_to(DP_DATA_MOD uint* bdata, size_t bitblen,
        const uint* val, size_t pos) {
    size_t blen = (bitblen + 31) >> 5;
    size_t i;
    size_t bitpos_all = ((pos * (DATA_PART_LEN)));
    size_t ppos = bitpos_all >> 5;
    uint bitposp = bitpos_all & 31;
    uint bmask0 = (1U << bitposp) - 1U;
    uint bmask1 = ~((1U << bitposp) - 1U);
    uint do_it = 1;
    uint state = 0;
    if ((pos * DATA_PART_LEN) > bitblen) {
        do_it = 0;
        for (i = 0; i < ((DATA_PART_LEN) >> 5); i++)
            if (val[i] != 0)
                state |= GLOBAL_STATE_ILLEGAL;
        if (((DATA_PART_LEN) & 31) != 0 &&
            (val[DATA_PART_LEN >> 5] & ((1U << (DATA_PART_LEN & 31)) - 1U)) != 0)
            state |= GLOBAL_STATE_ILLEGAL;
    } else if (((pos + 1) * DATA_PART_LEN) > bitblen) {
        uint after_end = bitblen - (pos * DATA_PART_LEN);
        if ((after_end & 31) != 0) {
            uint vx = ((DATA_PART_LEN & 31) != 0 &&
                (DATA_PART_LEN >> 5) == (after_end >> 5)) ?
                (val[after_end >> 5] & ((1U << (DATA_PART_LEN & 31)) - 1U)) :
                    val[after_end >> 5];
            if (((vx) & ~((1U << (after_end & 31)) - 1U)) != 0) {
                do_it = 0;
                state |= GLOBAL_STATE_ILLEGAL;
            }
            after_end += ((after_end + 32) & ~31);
        }
        if (do_it) {
            for (i = after_end >> 5; i < ((DATA_PART_LEN + 31) >> 5); i++) {
                if ((DATA_PART_LEN & 31) != 0 && (DATA_PART_LEN >> 5) == i) {
                    if ((val[i] & ((1U << (DATA_PART_LEN & 31)) - 1U)) != 0) {
                        do_it = 0;
                        state |= GLOBAL_STATE_ILLEGAL;
                    }
                } else if (val[i] != 0) {
                    do_it = 0;
                    state |= GLOBAL_STATE_ILLEGAL;
                }
            }
        }
    }
    if (do_it)
        for (i = 0; i < ((DATA_PART_LEN + 31) >> 5); i++) {
            if (ppos + i < blen) {
                uint newbmask0 = bmask0;
                uint newbmask1 = bmask1;
                uint range0 = DATA_PART_LEN;
                uint range1 = DATA_PART_LEN;
                uint vx = val[i];
                if (bitpos_all + ((i + 1) << 5) > bitblen)
                    range0 = bitblen - (bitpos_all + (i << 5));
                if ((DATA_PART_LEN & 31) != 0 && (DATA_PART_LEN >> 5) == i) {
                    range1 = DATA_PART_LEN - (i << 5);
                    range0 = range0 < range1 ? range0 : range1;
                }
                if (range0 != DATA_PART_LEN) {
                    uint bbp = range0;
                    vx &= ((1U << range0) - 1U);
                    if (bbp + bitposp < 32) {
                        newbmask0 |= (~((1U << (bbp + bitposp)) - 1U));
                        newbmask1 = 0xffffffffU;
                    } else if (bbp + bitposp >= 32)
                        newbmask1 |= (~((1U << ((bbp + bitposp) - 32)) - 1U));
                }
                bdata[ppos + i] = (bdata[ppos + i] & newbmask0) | (vx << bitposp);
                if (bitposp != 0 && ppos + i + 1 < blen &&
                    bitposp + DATA_PART_LEN - (i<<5) > 32) {
                    bdata[ppos + i + 1] = (bdata[ppos + i + 1] & newbmask1) |
                        (vx >> (32 - bitposp));
                }
            }
        }
    return state;
}
#  endif
#endif
"##;

fn gen_circuit_aggr_pop_defs(
    state_len: usize,
    cell_len_bits: usize,
    data_part_len: usize,
) -> String {
    let cell_len = 1 << cell_len_bits;
    let mut out = String::new();
    // POPULATE_READ_MEM_CELL_ALL - routine to populate read mem cell to circuit input:
    //    form: v=READ_MEM_CELL_DWORD(entry, 0); INPUT_TRANSFORM_BXX(iXX0,iXX1,iXX2,v); ....
    out += r##"#define POPULATE_READ_MEM_CELL_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
"##;
    if cell_len < 32 {
        out += r##"    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0); \
"##;
        writeln!(
            out,
            "    INPUT_TRANSFORM_B{}({}, v); \\",
            cell_len,
            (0..cell_len)
                .map(|j| format!("i{}", state_len + j))
                .collect::<Vec<_>>()
                .join(", ")
        )
        .unwrap();
    } else {
        for i in 0..(cell_len >> 5) {
            write!(
                out,
                r##"    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, {}); \
"##,
                i
            )
            .unwrap();
            writeln!(
                out,
                "    INPUT_TRANSFORM_B32({}, v); \\",
                (0..32)
                    .map(|j| format!("i{}", state_len + 32 * i + j))
                    .collect::<Vec<_>>()
                    .join(", ")
            )
            .unwrap();
        }
    }
    out += "}\n";
    // AGGREGATE_WRITE_MEM_CELL_ALL - routine to aggregate write mem cell from circuit output:
    //    form: OUTPUT_TRANSFORM_BXX(v,oXX0,oXX1,...); SET_WRITE_MEM_CELL_DWORD(entry, 0, v); ...
    out += r##"#define AGGREGATE_WRITE_MEM_CELL_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
"##;
    if cell_len < 32 {
        writeln!(
            out,
            "    OUTPUT_TRANSFORM_B{}(v, {}); \\",
            cell_len,
            (0..cell_len)
                .map(|j| format!("o{}", state_len + j))
                .collect::<Vec<_>>()
                .join(", ")
        )
        .unwrap();
        out += r##"    for (k = 0; k < TYPE_LEN; k++) \
        SET_WRITE_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0, v[k]); \
"##;
    } else {
        for i in 0..(cell_len >> 5) {
            writeln!(
                out,
                "    OUTPUT_TRANSFORM_B32(v, {}); \\",
                (0..32)
                    .map(|j| format!("o{}", state_len + 32 * i + j))
                    .collect::<Vec<_>>()
                    .join(", ")
            )
            .unwrap();
            write!(
                out,
                r##"    for (k = 0; k < TYPE_LEN; k++) \
        SET_WRITE_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, {}, v[k]); \
"##,
                i
            )
            .unwrap();
        }
    }
    out += "}\n";
    // POPULATE_READ_DATA_PART_ALL - routine to populate read data part to circuit input:
    //    form: v=READ_DATA_PART_DWORD(entry, 0); INPUT_TRANSFORM_BXX(iXX0,iXX1,iXX2,v); ....
    out += r##"#define POPULATE_READ_DATA_PART_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
"##;
    for i in 0..(data_part_len >> 5) {
        write!(
            out,
            r##"    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, {}); \
"##,
            i
        )
        .unwrap();
        writeln!(
            out,
            "    INPUT_TRANSFORM_B32({}, v); \\",
            (0..32)
                .map(|j| format!("i{}", state_len + cell_len + 32 * i + j))
                .collect::<Vec<_>>()
                .join(", ")
        )
        .unwrap();
    }
    if (data_part_len & 31) != 0 {
        out += r##"    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD_LAST((E) + (idx*TYPE_LEN + k)*ENTRY_LEN); \
"##;
        writeln!(
            out,
            "    INPUT_TRANSFORM_B{}({}, v); \\",
            data_part_len & 31,
            ((data_part_len & !31)..data_part_len)
                .map(|j| format!("i{}", state_len + cell_len + j))
                .collect::<Vec<_>>()
                .join(", ")
        )
        .unwrap();
    }
    out += "}\n";
    // AGGREGATE_WRITE_DATA_PART_ALL - routine to aggregate write mem cell from circuit output:
    //    form: OUTPUT_TRANSFORM_BXX(v,oXX0,oXX1,...); SET_WRITE_DATA_PART_DWORD(entry, 0, v); ...
    if data_part_len <= 32 {
        out += "#define AGGREGATE_WRITE_DATA_PART_ALL(D) { \\\n";
        writeln!(
            out,
            "    OUTPUT_TRANSFORM_B{}((D), {}); \\",
            data_part_len,
            (0..data_part_len)
                .map(|j| format!("o{}", state_len + cell_len + j))
                .collect::<Vec<_>>()
                .join(", ")
        )
        .unwrap();
    } else {
        out += r##"#define AGGREGATE_WRITE_DATA_PART_ALL(D) { \
    uint k; \
    uint v[TYPE_LEN]; \
"##;
        for i in 0..(data_part_len >> 5) {
            writeln!(
                out,
                "    OUTPUT_TRANSFORM_B32(v, {}); \\",
                (0..32)
                    .map(|j| format!("o{}", state_len + cell_len + 32 * i + j))
                    .collect::<Vec<_>>()
                    .join(", ")
            )
            .unwrap();
            write!(
                out,
                r##"    for (k = 0; k < TYPE_LEN; k++) \
        (D)[k*((DATA_PART_LEN+31)>>5) + {}] = v[k]; \
"##,
                i
            )
            .unwrap();
        }
        if (data_part_len & 31) != 0 {
            writeln!(
                out,
                "    OUTPUT_TRANSFORM_B{}(v, {}); \\",
                data_part_len & 31,
                ((data_part_len & !31)..data_part_len)
                    .map(|j| format!("o{}", state_len + cell_len + j))
                    .collect::<Vec<_>>()
                    .join(", ")
            )
            .unwrap();
            out += r##"    for (k = 0; k < TYPE_LEN; k++) \
        (D)[k*((DATA_PART_LEN+31)>>5) + (DATA_PART_LEN>>5)] = v[k]; \
"##;
        }
    }
    out += "}\n";
    out
}

fn gen_circuit_setup_defs(pic: &ProcIntDataConfig) -> String {
    let mut out = String::new();
    let config = pic.config();
    writeln!(out, "#define DATA_PART_LEN ({})", config.data_part_len).unwrap();
    writeln!(out, "#define CELL_LEN_BITS ({})", config.cell_len_bits).unwrap();
    writeln!(out, "#define GLOBAL_STATE_STOP ({})", GLOBAL_STATE_STOP).unwrap();
    writeln!(
        out,
        "#define GLOBAL_STATE_ILLEGAL ({})",
        GLOBAL_STATE_ILLEGAL
    )
    .unwrap();
    writeln!(out, "#define ENTRY_LEN ({})", pic.len()).unwrap();
    writeln!(
        out,
        "#define MEM_ADDRESS_LEN ({})",
        pic.mem_address_len() >> 5
    )
    .unwrap();
    writeln!(
        out,
        "#define TEMP_BUFFER_POS_IN_ENTRY ({})",
        pic.mem_address_len() >> 5
    )
    .unwrap();
    writeln!(
        out,
        "#define TEMP_BUFFER_LEN ({})",
        pic.temp_buffer_len_in_bits()
    )
    .unwrap();
    writeln!(out, "#define MEM_CELL_POS ({})", pic.mem_cell_pos()).unwrap();
    writeln!(
        out,
        "#define WRITE_MEM_CELL_POS ({})",
        pic.write_mem_cell_pos(),
    )
    .unwrap();
    writeln!(out, "#define DATA_PART_POS ({})", pic.data_part_pos()).unwrap();
    writeln!(
        out,
        "#define DATA_PART_POS_POS ({})",
        pic.data_part_pos_pos(),
    )
    .unwrap();
    writeln!(out, "#define REST_POS ({})", pic.rest_pos()).unwrap();
    writeln!(
        out,
        "#define DATA_PART_MOVE_DONE_INPUT i{}",
        config.state_len + (1 << config.cell_len_bits) + config.data_part_len
    )
    .unwrap();
    writeln!(
        out,
        "#define OUTPUT_REST_OUTPUTS(V) OUTPUT_TRANSFORM_B9((V), {})",
        (0..9)
            .map(|i| format!(
                "o{}",
                config.state_len + (1 << config.cell_len_bits) + config.data_part_len + i
            ))
            .collect::<Vec<_>>()
            .join(", ")
    )
    .unwrap();
    out
}

const CIRCUIT_INIT_CODE: &str = r##"
    DP_DATA_MOD uint* global_state = (DP_DATA_MOD uint*)buffer;
#ifdef HAVE_OPENCL
    size_t lidx = get_local_id(0);
    local uint local_state;
#endif
    DP_DATA_MOD uint* entries = ((DP_DATA_MOD uint*)buffer) + 64;
#ifdef HAVE_OPENCL
    if (lidx == 0)
        local_state = 0;
#endif
"##;

const CIRCUIT_POPULATE_CODE: &str = r##"{
    POPULATE_READ_MEM_CELL_ALL(entries);
    POPULATE_READ_DATA_PART_ALL(entries);
    {
        uint temp[TYPE_LEN];
        uint k;
        for (k = 0; k < TYPE_LEN; k++) {
            uint restval = ENTRY_REST_OUTPUT(entries + (idx*TYPE_LEN + k)*ENTRY_LEN);
            temp[k] = DATA_PART_MOVE_DONE_VAL(restval);
        }
        INPUT_TRANSFORM_B1(DATA_PART_MOVE_DONE_INPUT, temp);
    }
}
"##;

const CIRCUIT_AGGREGATE_CODE: &str = r##"{
    uint state = 0;
    uint k;
    uint rest_values_tbl[TYPE_LEN];
    uint write_dp_tbl[TYPE_LEN * ((DATA_PART_LEN + 31) >> 5)];
    uint* write_dps = (uint*)write_dp_tbl;
    OUTPUT_REST_OUTPUTS(rest_values_tbl);
    AGGREGATE_WRITE_DATA_PART_ALL(write_dps);

    for (k = 0; k < TYPE_LEN; k++) {
        DP_DATA_MOD uint* entry = entries + (idx*TYPE_LEN + k)*ENTRY_LEN;
        uint do_it = 1;
        uint rest_values = rest_values_tbl[k];
        uint data_kind, dp_pos, dp_move;
        uint dp_move_done = 0;
        uint max_dp_pos;
        uint read_dp[(DATA_PART_LEN + 31) >> 5];
        uint* write_dp = write_dps + k*((DATA_PART_LEN + 31) >> 5);
        uint do_dp_read = DATA_PART_READ_VAL(rest_values);
        if (STOP_MACHINE_VAL(rest_values))
            state |= GLOBAL_STATE_STOP;
        // get data kind
        data_kind = DATA_KIND_VAL(rest_values);
        dp_move = DATA_PART_MOVE_DIR_VAL(rest_values);
        // get data part position, read data part and write data part
        switch (data_kind) {
            case 0: {   // memory address
                dp_pos = MEM_ADDRESS_POS(entry);
                if (do_dp_read) {
                    read_data_part_from(entry, MEM_ADDRESS_LEN, read_dp, dp_pos);
                    SET_READ_DATA_PART(entry, read_dp);
                }
                if (DATA_PART_WRITE_VAL(rest_values)) {
                    uint entry_state = write_data_part_to(entry, (MEM_ADDRESS_LEN<<5),
                            write_dp, dp_pos);
                    if ((entry_state & GLOBAL_STATE_ILLEGAL) != 0)
                        do_it = 0;
                    state |= entry_state;
                }
                max_dp_pos = 255;
                break;
            }
            case 1: {   // temp buffer
                dp_pos = TEMP_BUFFER_POS(entry);
                if (do_dp_read) {
                    read_data_part_from(entry + TEMP_BUFFER_POS_IN_ENTRY,
                            (TEMP_BUFFER_LEN + 31) >> 5,
                            read_dp, dp_pos);
                    SET_READ_DATA_PART(entry, read_dp);
                }
                if (DATA_PART_WRITE_VAL(rest_values)) {
                    uint entry_state = write_data_part_to(entry + TEMP_BUFFER_POS_IN_ENTRY,
                            TEMP_BUFFER_LEN, write_dp, dp_pos);
                    if ((entry_state & GLOBAL_STATE_ILLEGAL) != 0)
                        do_it = 0;
                    state |= entry_state;
                }
                max_dp_pos = 0xffff;
                break;
            }
            default: {   // proc_id
                ulong proc_id = idx*TYPE_LEN + k;
                uint proc_id_buffer[2];
                proc_id_buffer[0] = proc_id & 0xffffffffU;
                proc_id_buffer[1] = proc_id >> 32;
                dp_pos = PROC_ID_POS(entry);
                if (do_dp_read) {
                    read_data_part_from_proc_id(proc_id_buffer, 2, read_dp, dp_pos);
                    SET_READ_DATA_PART(entry, read_dp);
                }
                max_dp_pos = 255;
                break;
            }
        }
        if (!do_dp_read)
            CLEAR_READ_DATA_PART(entry);
        // move dp pos
        if (do_it)
            switch (dp_move) {
                case 1:
                    if (dp_pos >= max_dp_pos) {
                        state |= GLOBAL_STATE_ILLEGAL; 
                        do_it = 0;
                    } else {
                        dp_pos++;
                        dp_move_done = 1;
                    }
                    break;
                case 2:
                case 3:
                    if (dp_pos != 0) {
                        dp_pos--;
                        dp_move_done = 1;
                    }
                    break;
                default:
                    break;
            }
        // update dp pos
        if (dp_move_done) {
            switch (data_kind) {
                case 0:     // memory address
                    SET_MEM_ADDRESS_POS(entry, dp_pos);
                    break;
                case 1:     // temp buffer
                    SET_TEMP_BUFFER_POS(entry, dp_pos);
                    break;
                default:     // proc_id
                    SET_PROC_ID_POS(entry, dp_pos);
                    break;
            }
        }
        SET_DATA_PART_MOVE_DONE_VAL(rest_values, dp_move_done);
        SET_ENTRY_REST_OUTPUT(entry, rest_values);
    }

    AGGREGATE_WRITE_MEM_CELL_ALL(entries);
#ifdef HAVE_OPENCL
    barrier(CLK_LOCAL_MEM_FENCE);
    atomic_or(&local_state, state);
    barrier(CLK_LOCAL_MEM_FENCE);
    if (lidx == 0)
        atomic_or(global_state, local_state);
#else
    *global_state |= state;
#endif
}
"##;

pub(crate) fn build_circuit<'a, DR, DW, D, E, B, T>(
    mut builder: B,
    circuit: Circuit<T>,
    pic: &ProcIntDataConfig,
    proc_num: usize,
) -> Result<E, B::ErrorType>
where
    DR: DataReader,
    DW: DataWriter,
    D: DataHolder<'a, DR, DW>,
    E: Executor<'a, DR, DW, D>,
    B: Builder<'a, DR, DW, D, E>,
    T: Clone + Copy + Ord + PartialEq + Eq + Hash,
    T: Default + TryFrom<usize>,
    <T as TryFrom<usize>>::Error: Debug,
    usize: TryFrom<T>,
    <usize as TryFrom<T>>::Error: Debug,
{
    builder.transform_helpers();
    builder.user_defs(&gen_circuit_setup_defs(&pic));
    builder.user_defs(CIRCUIT_USER_DEFS_CODE);
    let config = pic.config();
    let circuit_input_len = usize::try_from(circuit.input_len()).unwrap();
    let circuit_output_len = circuit.outputs().len();
    let state_len =
        circuit_input_len - (1 << config.cell_len_bits) - config.data_part_len as usize - 1;
    builder.user_defs(&gen_circuit_aggr_pop_defs(
        state_len,
        config.cell_len_bits as usize,
        config.data_part_len as usize,
    ));
    builder.add_with_config(
        "machine",
        circuit,
        CodeConfig::new()
            .single_buffer(true)
            .init_code(Some(&CIRCUIT_INIT_CODE))
            .pop_input_code(Some(&CIRCUIT_POPULATE_CODE))
            .pop_from_buffer(Some(&(state_len..circuit_input_len).collect::<Vec<_>>()))
            .pop_input_len(Some(64 + proc_num * pic.len()))
            .aggr_output_code(Some(&CIRCUIT_AGGREGATE_CODE))
            .aggr_only_to_buffer(Some(&(state_len..circuit_output_len).collect::<Vec<_>>()))
            .aggr_output_len(Some(64 + proc_num * pic.len())),
    );
    builder.build().map(|mut e| e.remove(0))
}

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

    use gatenative::cpu_build_exec::*;
    use gatenative::opencl_build_exec::*;
    use libloading::{Library, Symbol};

    fn circuit_user_defs_test_data_part(data_part_len: usize, code: &str) -> Library {
        let shared_lib = SharedLib::new_with_cpu_ext(CPUExtension::NoExtension);
        let test_code = r##"#include <stddef.h>
#include <stdint.h>
#include <stdio.h>
#include <unistd.h>
#define GLOBAL_STATE_ILLEGAL (2)
"##
        .to_string()
            + &format!("#define DATA_PART_LEN ({})", data_part_len)
            + CIRCUIT_USER_DEFS_CODE
            + code;
        shared_lib.build(test_code.as_bytes()).unwrap()
    }

    #[test]
    fn test_circuit_user_defs_code() {
        let test_code = r##"
#define MY_ASSERT(A) { if (!(A)) return __LINE__; }
#define MY_ASSERT_EQ(A, B) { \
    if ((A) != (B)) { \
        printf("AssertEq: %u!=%u\n", (A), (B)); \
        return __LINE__; \
    } \
}
#define MY_ASSERT_EQ_ARRAY(A, B, LEN) { \
    size_t i; \
    for (i = 0; i < (LEN); i++) \
        if ((A)[i] != (B)[i]) { \
            printf("AssertEqArray: %lu: %u!=%u\n", (ulong)i, (A)[i], (B)[i]); \
            return __LINE__; \
        } \
}

void simple_read_data_part_from(size_t data_part_len, const uint* barray,
        size_t blen, uint* v, size_t pos) {
    size_t i;
    size_t dp_len_dwords = (data_part_len + 31) >> 5;
    // clear
    for (i = 0; i < dp_len_dwords; i++)
        v[i] = 0;
    // read from barray to v
    for (i = 0; i < data_part_len; i++) {
        size_t src = pos * data_part_len + i;
        if ((src >> 5) < blen) {
            uint bit = (barray[src >> 5] >> (src & 31)) & 1;
            v[i >> 5] |= bit << (i & 31);
        }
    }
}

uint simple_write_data_part_to(size_t data_part_len, uint* barray,
        size_t bitblen, const uint* v, size_t pos) {
    size_t i;
    // write from v to barray
    uint illegal = 0;
    for (i = 0; i < data_part_len; i++) {
        size_t dest = pos * data_part_len + i;
        uint bit = ((v[i >> 5] >> (i & 31)) & 1);
        if (dest >= bitblen && bit != 0)
            illegal = 1;
    }
    if (!illegal)
        for (i = 0; i < data_part_len; i++) {
            size_t dest = pos * data_part_len + i;
            uint bit = ((v[i >> 5] >> (i & 31)) & 1);
            if (dest < bitblen) {
                barray[dest >> 5] &= ~(1U << (dest & 31));
                barray[dest >> 5] |= bit << (dest & 31);
            }
        }
    return illegal;
}

// return 0 if no fail. return number of line if fail
uint32_t test1(void) {
    printf("StartLineNo: %u\n", __LINE__);
    size_t blen, bitblen, pos, k;
    const uint data[12 + 8] = {
        0xcd0a3b56U, 0x7d19d24eU, 0xda093bf1U, 0x3a0589c6U,
        0x7829cda0U, 0xc0d93e16U, 0xef9a371dU, 0xea0ca392U,
        0x14f059a4U, 0x3a0e50a7U, 0xe569b121U, 0xba0ea352U,
        0, 0, 0, 0, 0, 0, 0, 0
    };
    uint exp_data[12 + 8];
    uint res_data[12 + 8];
    uint expv[8];
    uint resv[8];
    uint state = 0;
    size_t dp_len_dwords = (DATA_PART_LEN + 31) >> 5;
    for (blen = 1; blen <= 12; blen++) {
        for (pos = 0; pos * DATA_PART_LEN < 12 * 32; pos++) {
            printf("TestRead: %u %lu %lu\n", DATA_PART_LEN,
                (unsigned long)blen, (unsigned long)pos);
            simple_read_data_part_from(DATA_PART_LEN, data, blen, expv, pos);
            for (k = 0; k < dp_len_dwords; k++)
                resv[k] = 0;
            read_data_part_from(data, blen, resv, pos);
            MY_ASSERT_EQ_ARRAY(expv, resv, dp_len_dwords);
            for (k = 0; k < dp_len_dwords; k++)
                resv[k] = 0;
            read_data_part_from_proc_id(data, blen, resv, pos);
            MY_ASSERT_EQ_ARRAY(expv, resv, dp_len_dwords);
        }
    }
    for (bitblen = 3; bitblen <= 12 * 32; bitblen += 3) {
        for (pos = 0; pos * DATA_PART_LEN < 12 * 32; pos++) {
            uint do_illegal_n = (pos + 1) * DATA_PART_LEN - 1 >= bitblen ? 2 : 1;
            uint t = 0;
            blen = (bitblen + 31) >> 5;
            for (t = 0; t < do_illegal_n; t++) {
                // t=0 only legal, t=1 only illegal
                state = 0;
                // copy to datas
                for (k = 0; k < 12 + 8; k++) {
                    exp_data[k] = 0;
                    res_data[k] = 0;
                }
                for (k = 0; k < bitblen; k++) {
                    exp_data[k >> 5] |= data[k >> 5] & (1U << (k & 31));
                    res_data[k >> 5] |= data[k >> 5] & (1U << (k & 31));
                }
                // read value
                simple_read_data_part_from(DATA_PART_LEN, exp_data, blen, expv, pos);
                // negate read value
                for (k = 0; k < DATA_PART_LEN; k++) {
                    if (t || (pos * DATA_PART_LEN + k) < bitblen)
                        // if illegal or not in bound of data
                        expv[k >> 5] ^= (1U << (k & 31));
                    else
                        // clear if out of bound
                        expv[k >> 5] &= ~(1U << (k & 31));
                }
                printf("TestWrite: %u %lu %lu %u %u\n", DATA_PART_LEN,
                    (unsigned long)bitblen, (unsigned long)pos, t, expv[0]);
                // make write
                simple_write_data_part_to(DATA_PART_LEN, exp_data, bitblen, expv, pos);
                // make write real
                state = write_data_part_to(res_data, bitblen, expv, pos);
                if (!t) {
                    MY_ASSERT_EQ_ARRAY(exp_data, res_data, 12 + 8);
                    MY_ASSERT_EQ(0, state);
                } else {
                    MY_ASSERT_EQ_ARRAY(exp_data, res_data, 12 + 8);
                    MY_ASSERT_EQ(GLOBAL_STATE_ILLEGAL, state);
                }
            }
        }
    }
    return 0;
}
"##;
        // for dp_len in (1..160).chain((5..9).map(|i| 32 * i)) {
        for dp_len in [
            1, 2, 3, 4, 7, 8, 12, 16, 23, 32, 33, 37, 38, 46, 58, 64, 74, 80, 96, 97, 104, 111,
            128, 142, 147, 160, 192, 200, 224, 256,
        ] {
            let lib = circuit_user_defs_test_data_part(dp_len, &test_code);
            let test_func: Symbol<unsafe extern "C" fn() -> u32> =
                unsafe { lib.get(b"test1").unwrap() };
            assert_eq!(unsafe { test_func() }, 0, "dp_len={}", dp_len);
        }
    }

    #[test]
    fn test_gen_circuit_aggr_pop_defs() {
        assert_eq!(
            r##"#define POPULATE_READ_MEM_CELL_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0); \
    INPUT_TRANSFORM_B4(i36, i37, i38, i39, v); \
}
#define AGGREGATE_WRITE_MEM_CELL_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    OUTPUT_TRANSFORM_B4(v, o36, o37, o38, o39); \
    for (k = 0; k < TYPE_LEN; k++) \
        SET_WRITE_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0, v[k]); \
}
#define POPULATE_READ_DATA_PART_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD_LAST((E) + (idx*TYPE_LEN + k)*ENTRY_LEN); \
    INPUT_TRANSFORM_B5(i40, i41, i42, i43, i44, v); \
}
#define AGGREGATE_WRITE_DATA_PART_ALL(D) { \
    OUTPUT_TRANSFORM_B5((D), o40, o41, o42, o43, o44); \
}
"##,
            gen_circuit_aggr_pop_defs(36, 2, 5)
        );
        assert_eq!(
            r##"#define POPULATE_READ_MEM_CELL_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0); \
    INPUT_TRANSFORM_B16(i40, i41, i42, i43, i44, i45, i46, i47, i48, i49, i50, i51, i52, i53, i54, i55, v); \
}
#define AGGREGATE_WRITE_MEM_CELL_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    OUTPUT_TRANSFORM_B16(v, o40, o41, o42, o43, o44, o45, o46, o47, o48, o49, o50, o51, o52, o53, o54, o55); \
    for (k = 0; k < TYPE_LEN; k++) \
        SET_WRITE_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0, v[k]); \
}
#define POPULATE_READ_DATA_PART_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD_LAST((E) + (idx*TYPE_LEN + k)*ENTRY_LEN); \
    INPUT_TRANSFORM_B27(i56, i57, i58, i59, i60, i61, i62, i63, i64, i65, i66, i67, i68, i69, i70, i71, i72, i73, i74, i75, i76, i77, i78, i79, i80, i81, i82, v); \
}
#define AGGREGATE_WRITE_DATA_PART_ALL(D) { \
    OUTPUT_TRANSFORM_B27((D), o56, o57, o58, o59, o60, o61, o62, o63, o64, o65, o66, o67, o68, o69, o70, o71, o72, o73, o74, o75, o76, o77, o78, o79, o80, o81, o82); \
}
"##,
            gen_circuit_aggr_pop_defs(40, 4, 27)
        );
        assert_eq!(
            r##"#define POPULATE_READ_MEM_CELL_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0); \
    INPUT_TRANSFORM_B32(i40, i41, i42, i43, i44, i45, i46, i47, i48, i49, i50, i51, i52, i53, i54, i55, i56, i57, i58, i59, i60, i61, i62, i63, i64, i65, i66, i67, i68, i69, i70, i71, v); \
}
#define AGGREGATE_WRITE_MEM_CELL_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    OUTPUT_TRANSFORM_B32(v, o40, o41, o42, o43, o44, o45, o46, o47, o48, o49, o50, o51, o52, o53, o54, o55, o56, o57, o58, o59, o60, o61, o62, o63, o64, o65, o66, o67, o68, o69, o70, o71); \
    for (k = 0; k < TYPE_LEN; k++) \
        SET_WRITE_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0, v[k]); \
}
#define POPULATE_READ_DATA_PART_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0); \
    INPUT_TRANSFORM_B32(i72, i73, i74, i75, i76, i77, i78, i79, i80, i81, i82, i83, i84, i85, i86, i87, i88, i89, i90, i91, i92, i93, i94, i95, i96, i97, i98, i99, i100, i101, i102, i103, v); \
}
#define AGGREGATE_WRITE_DATA_PART_ALL(D) { \
    OUTPUT_TRANSFORM_B32((D), o72, o73, o74, o75, o76, o77, o78, o79, o80, o81, o82, o83, o84, o85, o86, o87, o88, o89, o90, o91, o92, o93, o94, o95, o96, o97, o98, o99, o100, o101, o102, o103); \
}
"##,
            gen_circuit_aggr_pop_defs(40, 5, 32)
        );
        assert_eq!(
            r##"#define POPULATE_READ_MEM_CELL_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0); \
    INPUT_TRANSFORM_B32(i100, i101, i102, i103, i104, i105, i106, i107, i108, i109, i110, i111, i112, i113, i114, i115, i116, i117, i118, i119, i120, i121, i122, i123, i124, i125, i126, i127, i128, i129, i130, i131, v); \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 1); \
    INPUT_TRANSFORM_B32(i132, i133, i134, i135, i136, i137, i138, i139, i140, i141, i142, i143, i144, i145, i146, i147, i148, i149, i150, i151, i152, i153, i154, i155, i156, i157, i158, i159, i160, i161, i162, i163, v); \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 2); \
    INPUT_TRANSFORM_B32(i164, i165, i166, i167, i168, i169, i170, i171, i172, i173, i174, i175, i176, i177, i178, i179, i180, i181, i182, i183, i184, i185, i186, i187, i188, i189, i190, i191, i192, i193, i194, i195, v); \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 3); \
    INPUT_TRANSFORM_B32(i196, i197, i198, i199, i200, i201, i202, i203, i204, i205, i206, i207, i208, i209, i210, i211, i212, i213, i214, i215, i216, i217, i218, i219, i220, i221, i222, i223, i224, i225, i226, i227, v); \
}
#define AGGREGATE_WRITE_MEM_CELL_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    OUTPUT_TRANSFORM_B32(v, o100, o101, o102, o103, o104, o105, o106, o107, o108, o109, o110, o111, o112, o113, o114, o115, o116, o117, o118, o119, o120, o121, o122, o123, o124, o125, o126, o127, o128, o129, o130, o131); \
    for (k = 0; k < TYPE_LEN; k++) \
        SET_WRITE_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0, v[k]); \
    OUTPUT_TRANSFORM_B32(v, o132, o133, o134, o135, o136, o137, o138, o139, o140, o141, o142, o143, o144, o145, o146, o147, o148, o149, o150, o151, o152, o153, o154, o155, o156, o157, o158, o159, o160, o161, o162, o163); \
    for (k = 0; k < TYPE_LEN; k++) \
        SET_WRITE_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 1, v[k]); \
    OUTPUT_TRANSFORM_B32(v, o164, o165, o166, o167, o168, o169, o170, o171, o172, o173, o174, o175, o176, o177, o178, o179, o180, o181, o182, o183, o184, o185, o186, o187, o188, o189, o190, o191, o192, o193, o194, o195); \
    for (k = 0; k < TYPE_LEN; k++) \
        SET_WRITE_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 2, v[k]); \
    OUTPUT_TRANSFORM_B32(v, o196, o197, o198, o199, o200, o201, o202, o203, o204, o205, o206, o207, o208, o209, o210, o211, o212, o213, o214, o215, o216, o217, o218, o219, o220, o221, o222, o223, o224, o225, o226, o227); \
    for (k = 0; k < TYPE_LEN; k++) \
        SET_WRITE_MEM_CELL_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 3, v[k]); \
}
#define POPULATE_READ_DATA_PART_ALL(E) { \
    uint k; \
    uint v[TYPE_LEN]; \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 0); \
    INPUT_TRANSFORM_B32(i228, i229, i230, i231, i232, i233, i234, i235, i236, i237, i238, i239, i240, i241, i242, i243, i244, i245, i246, i247, i248, i249, i250, i251, i252, i253, i254, i255, i256, i257, i258, i259, v); \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 1); \
    INPUT_TRANSFORM_B32(i260, i261, i262, i263, i264, i265, i266, i267, i268, i269, i270, i271, i272, i273, i274, i275, i276, i277, i278, i279, i280, i281, i282, i283, i284, i285, i286, i287, i288, i289, i290, i291, v); \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 2); \
    INPUT_TRANSFORM_B32(i292, i293, i294, i295, i296, i297, i298, i299, i300, i301, i302, i303, i304, i305, i306, i307, i308, i309, i310, i311, i312, i313, i314, i315, i316, i317, i318, i319, i320, i321, i322, i323, v); \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 3); \
    INPUT_TRANSFORM_B32(i324, i325, i326, i327, i328, i329, i330, i331, i332, i333, i334, i335, i336, i337, i338, i339, i340, i341, i342, i343, i344, i345, i346, i347, i348, i349, i350, i351, i352, i353, i354, i355, v); \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD((E) + (idx*TYPE_LEN + k)*ENTRY_LEN, 4); \
    INPUT_TRANSFORM_B32(i356, i357, i358, i359, i360, i361, i362, i363, i364, i365, i366, i367, i368, i369, i370, i371, i372, i373, i374, i375, i376, i377, i378, i379, i380, i381, i382, i383, i384, i385, i386, i387, v); \
    for (k = 0; k < TYPE_LEN; k++) \
        v[k] = READ_DATA_PART_DWORD_LAST((E) + (idx*TYPE_LEN + k)*ENTRY_LEN); \
    INPUT_TRANSFORM_B7(i388, i389, i390, i391, i392, i393, i394, v); \
}
#define AGGREGATE_WRITE_DATA_PART_ALL(D) { \
    uint k; \
    uint v[TYPE_LEN]; \
    OUTPUT_TRANSFORM_B32(v, o228, o229, o230, o231, o232, o233, o234, o235, o236, o237, o238, o239, o240, o241, o242, o243, o244, o245, o246, o247, o248, o249, o250, o251, o252, o253, o254, o255, o256, o257, o258, o259); \
    for (k = 0; k < TYPE_LEN; k++) \
        (D)[k*((DATA_PART_LEN+31)>>5) + 0] = v[k]; \
    OUTPUT_TRANSFORM_B32(v, o260, o261, o262, o263, o264, o265, o266, o267, o268, o269, o270, o271, o272, o273, o274, o275, o276, o277, o278, o279, o280, o281, o282, o283, o284, o285, o286, o287, o288, o289, o290, o291); \
    for (k = 0; k < TYPE_LEN; k++) \
        (D)[k*((DATA_PART_LEN+31)>>5) + 1] = v[k]; \
    OUTPUT_TRANSFORM_B32(v, o292, o293, o294, o295, o296, o297, o298, o299, o300, o301, o302, o303, o304, o305, o306, o307, o308, o309, o310, o311, o312, o313, o314, o315, o316, o317, o318, o319, o320, o321, o322, o323); \
    for (k = 0; k < TYPE_LEN; k++) \
        (D)[k*((DATA_PART_LEN+31)>>5) + 2] = v[k]; \
    OUTPUT_TRANSFORM_B32(v, o324, o325, o326, o327, o328, o329, o330, o331, o332, o333, o334, o335, o336, o337, o338, o339, o340, o341, o342, o343, o344, o345, o346, o347, o348, o349, o350, o351, o352, o353, o354, o355); \
    for (k = 0; k < TYPE_LEN; k++) \
        (D)[k*((DATA_PART_LEN+31)>>5) + 3] = v[k]; \
    OUTPUT_TRANSFORM_B32(v, o356, o357, o358, o359, o360, o361, o362, o363, o364, o365, o366, o367, o368, o369, o370, o371, o372, o373, o374, o375, o376, o377, o378, o379, o380, o381, o382, o383, o384, o385, o386, o387); \
    for (k = 0; k < TYPE_LEN; k++) \
        (D)[k*((DATA_PART_LEN+31)>>5) + 4] = v[k]; \
    OUTPUT_TRANSFORM_B7(v, o388, o389, o390, o391, o392, o393, o394); \
    for (k = 0; k < TYPE_LEN; k++) \
        (D)[k*((DATA_PART_LEN+31)>>5) + (DATA_PART_LEN>>5)] = v[k]; \
}
"##,
            gen_circuit_aggr_pop_defs(100, 7, 167)
        );
    }

    #[test]
    fn test_gen_circuit_setup_defs() {
        let config = InfParMachineConfig {
            state_len: 100,
            data_part_len: 80,
            cell_len_bits: 6,
        };
        let env_config = InfParEnvConfig {
            proc_num: 1 << 8,
            flat_memory: true,
            max_temp_buffer_len: 24,
            max_mem_size: Some(20 * (1 << 20)),
        };
        let pic = ProcIntDataConfig::new(config, env_config);
        assert_eq!(
            r##"#define DATA_PART_LEN (80)
#define CELL_LEN_BITS (6)
#define GLOBAL_STATE_STOP (1)
#define GLOBAL_STATE_ILLEGAL (2)
#define ENTRY_LEN (10)
#define MEM_ADDRESS_LEN (1)
#define TEMP_BUFFER_POS_IN_ENTRY (1)
#define TEMP_BUFFER_LEN (24)
#define MEM_CELL_POS (3)
#define WRITE_MEM_CELL_POS (5)
#define DATA_PART_POS (7)
#define DATA_PART_POS_POS (2)
#define REST_POS (9)
#define DATA_PART_MOVE_DONE_INPUT i244
#define OUTPUT_REST_OUTPUTS(V) OUTPUT_TRANSFORM_B9((V), o244, o245, o246, o247, o248, o249, o250, o251, o252)
"##,
            gen_circuit_setup_defs(&pic)
        );
    }

    use gatenative::clang_writer::*;
    use opencl3::device::{get_all_devices, Device, CL_DEVICE_TYPE_GPU};

    fn cpu_builder_creator() -> CPUBuilder<'static> {
        CPUBuilder::new_with_cpu_ext_and_clang_config(
            CPUExtension::NoExtension,
            &CLANG_WRITER_U64,
            None,
        )
    }

    fn opencl_builder_creator() -> OpenCLBuilder<'static> {
        let device = Device::new(*get_all_devices(CL_DEVICE_TYPE_GPU).unwrap().get(0).unwrap());
        OpenCLBuilder::new(&device, None)
    }

    macro_rules! test_build_circuit_xxx {
        ($name:ident, $bc:ident) => {
            fn $name() {
                //
                // 32-bits
                //
                let builder = $bc();
                let config = InfParMachineConfig {
                    state_len: 20,
                    data_part_len: 32,
                    cell_len_bits: 5,
                };
                let env_config = InfParEnvConfig {
                    proc_num: 1024,
                    flat_memory: true,
                    max_temp_buffer_len: 96,
                    max_mem_size: Some(20 * (1 << 41)),
                };
                let pic = ProcIntDataConfig::new(config, env_config);
                let mem_cell_data_part_len = (1 << config.cell_len_bits) + config.data_part_len;
                let circuit_input_len = config.state_len + mem_cell_data_part_len + 1;
                // negate mem cell to write and data part to write to distinguish values from
                // read mem cell and read data part
                let circuit = Circuit::<u32>::new(
                    circuit_input_len,
                    [Gate::new_nor(0, circuit_input_len - 1)],
                    std::iter::once((circuit_input_len, true)).chain(
                        (1..circuit_input_len - 1)
                            .chain(0..9)
                            .map(|i| (i, i >= config.state_len && i < circuit_input_len - 1)),
                    ),
                )
                .unwrap();
                let mut exec = build_circuit(
                    builder,
                    circuit,
                    &pic,
                    usize::try_from(env_config.proc_num).unwrap(),
                )
                .unwrap();
                // execution test 1
                let input_state_rest_data = std::iter::repeat(0)
                    .take(100)
                    .chain(
                        [
                            //DPMD,STOP,KIND,DPMOVE,DPRW,MEMRW
                            0b00_00_00_00_01, // 0
                            0b00_00_00_01_10, // 1
                            0b00_00_00_10_11, // 2
                            0b00_00_00_11_00, // 3
                            0b00_01_00_00_01, // 4
                            0b00_01_00_01_10, // 5
                            0b00_01_00_10_11, // 6
                            0b00_01_00_11_00, // 7
                            0b00_10_00_00_01, // 8
                            0b00_10_00_01_10, // 9
                            0b00_10_00_10_11, // 10
                            0b00_10_00_11_00, // 11
                            0b00_11_00_00_01, // 12
                            0b00_11_00_01_10, // 13
                            0b00_11_00_10_11, // 14
                            0b00_11_00_11_00, // 15
                            0b00_00_01_00_01, // 16
                            0b00_00_01_01_10, // 17
                            0b00_00_01_10_11, // 18
                            0b00_00_01_11_00, // 19
                            0b00_01_01_00_01, // 20
                            0b00_01_01_01_10, // 21
                            0b00_01_01_10_11, // 22
                            0b00_01_01_11_00, // 23
                            0b00_10_01_00_01, // 24
                            0b00_10_01_01_10, // 25
                            0b00_10_01_10_11, // 26
                            0b00_10_01_11_00, // 27
                            0b00_11_01_00_01, // 28
                            0b00_11_01_01_10, // 29
                            0b00_11_01_10_11, // 30
                            0b00_11_01_11_00, // 31
                            0b00_00_10_00_01, // 0
                            0b00_00_10_01_10, // 1
                            0b00_00_10_10_11, // 2
                            0b00_00_10_11_00, // 3
                            0b00_01_10_00_01, // 4
                            0b00_01_10_01_10, // 5
                            0b00_01_10_10_11, // 6
                            0b00_01_10_11_00, // 7
                            0b00_10_10_00_01, // 8
                            0b00_10_10_01_10, // 9
                            0b00_10_10_10_11, // 10
                            0b00_10_10_11_00, // 11
                            0b00_11_10_00_01, // 12
                            0b00_11_10_01_10, // 13
                            0b00_11_10_10_11, // 14
                            0b00_11_10_11_00, // 15
                            0b00_00_11_00_01, // 16
                            0b00_00_11_01_10, // 17
                            0b00_00_11_10_11, // 18
                            0b00_00_11_11_00, // 19
                            0b00_01_11_00_01, // 20
                            0b00_01_11_01_10, // 21
                            0b00_01_11_10_11, // 22
                            0b00_01_11_11_00, // 23
                            0b00_10_11_00_01, // 24
                            0b00_10_11_01_10, // 25
                            0b00_10_11_10_11, // 26
                            0b00_10_11_11_00, // 27
                            0b00_11_11_00_01, // 28
                            0b00_11_11_01_10, // 29
                            0b00_11_11_10_11, // 30
                            0b00_11_11_11_00, // 31
                        ]
                        .into_iter(),
                    )
                    .chain(std::iter::repeat(0).take(1024 - 64 - 100))
                    .collect::<Vec<_>>();
                println!("PIC entry len {}", pic.len());
                let res_output_data = std::iter::repeat(0)
                    .take(64 + 100 * pic.len())
                    .chain(
                        (0..64)
                            .map(|k| {
                                [
                                    // entry 0
                                    0xd0895511,        // memaddr 0
                                    0x844dd61 + 5 * k, // memaddr 1
                                    0xdd17c1 + k,      // temp buffer 0
                                    0x1af1a1,          // temp buffer 1
                                    0x1c141b1 + 4 * k, // temp buffer 2
                                    0x00000201,        // data positions
                                    0x22415b + 3 * k,  // mem read
                                    0xb30401 + k,      // mem write
                                    0xda931 + 2 * k,   // data read
                                    0x7ba00000,        // rest
                                ]
                            })
                            .flatten(),
                    )
                    .chain(
                        // test for loading data_part_move_done
                        [0, 0, 0, 0, 0, 0, 0, 0, 0, 1 << 31].into_iter(),
                    )
                    .chain(std::iter::repeat(0).take((1024 - 64 - 100 - 1) * pic.len()))
                    .collect::<Vec<_>>();
                // execute main test
                let mut it = exec
                    .input_transformer(32, &(0..20).collect::<Vec<_>>())
                    .unwrap();
                let mut ot = exec
                    .output_transformer(32, &(0..20).collect::<Vec<_>>())
                    .unwrap();
                let input_data = exec.new_data_from_vec(input_state_rest_data.clone());
                let mut input_data = it.transform(&input_data).unwrap();
                let mut buffer = exec.new_data_from_vec(res_output_data.clone());
                exec.execute_buffer_single(&mut input_data, 0, &mut buffer)
                    .unwrap();
                {
                    let output_data = ot.transform(&input_data).unwrap();
                    let output_data = output_data.get();
                    let output_data = output_data.get();
                    // if data_part_move_done is 1 then must be 1.
                    assert_eq!(output_data[100 + 64], 1);
                }
                // expected buffer data
                let exp_buffer_data = std::iter::repeat(0)
                    .take(64)
                    .chain(
                        std::iter::repeat([0u32, 0, 0, 0, 0, 0, 0, 0xffffffff, 0, 0])
                            .take(100)
                            .flatten(),
                    )
                    .chain([
                        // 100
                        0xd0895511, 0x0844dd61, 0x00dd17c1, 0x001af1a1, 0x01c141b1, 0x00000201,
                        0x0022415b, 0xffddbea4, 0x00000000, 0x00600000, // 101
                        0xd0895511, 0x0844dd66, 0x00dd17c2, 0x001af1a1, 0x01c141b5, 0x00000201,
                        0x0022415e, 0xffddbea1, 0x0844dd66, 0x01a00000, // 102
                        0xd0895511, 0xfff256ca, 0x00dd17c3, 0x001af1a1, 0x01c141b9, 0x00000201,
                        0x00224161, 0xffddbe9e, 0x00000000, 0x02e00000, // 103
                        0xd0895511, 0xfff256c8, 0x00dd17c4, 0x001af1a1, 0x01c141bd, 0x00000201,
                        0x00224164, 0xffddbe9b, 0x0844dd70, 0x03200000, // 104
                        0xd0895511, 0x0844dd75, 0x00dd17c5, 0x001af1a1, 0x01c141c1, 0x00000201,
                        0x00224167, 0xffddbe98, 0x00000000, 0x10600000, // 105
                        0xd0895511, 0x0844dd7a, 0x00dd17c6, 0x001af1a1, 0x01c141c5, 0x00000201,
                        0x0022416a, 0xffddbe95, 0x01c141c5, 0x11a00000, // 106
                        0xd0895511, 0x0844dd7f, 0x00dd17c7, 0x001af1a1, 0xfff256c2, 0x00000201,
                        0x0022416d, 0xffddbe92, 0x00000000, 0x12e00000, // 107
                        0xd0895511, 0x0844dd84, 0x00dd17c8, 0x001af1a1, 0xfff256c0, 0x00000201,
                        0x00224170, 0xffddbe8f, 0x01c141cd, 0x13200000, // 108
                        0xd0895511, 0x0844dd89, 0x00dd17c9, 0x001af1a1, 0x01c141d1, 0x00000201,
                        0x00224173, 0xffddbe8c, 0x00000000, 0x20600000, // 109
                        0xd0895511, 0x0844dd8e, 0x00dd17ca, 0x001af1a1, 0x01c141d5, 0x00000201,
                        0x00224176, 0xffddbe89, 0x0000006d, 0x21a00000, // 110
                        0xd0895511, 0x0844dd93, 0x00dd17cb, 0x001af1a1, 0x01c141d9, 0x00000201,
                        0x00224179, 0xffddbe86, 0x00000000, 0x22e00000, // 111
                        0xd0895511, 0x0844dd98, 0x00dd17cc, 0x001af1a1, 0x01c141dd, 0x00000201,
                        0x0022417c, 0xffddbe83, 0x0000006f, 0x23200000, // 112
                        0xd0895511, 0x0844dd9d, 0x00dd17cd, 0x001af1a1, 0x01c141e1, 0x00000201,
                        0x0022417f, 0xffddbe80, 0x00000000, 0x30600000, // 113
                        0xd0895511, 0x0844dda2, 0x00dd17ce, 0x001af1a1, 0x01c141e5, 0x00000201,
                        0x00224182, 0xffddbe7d, 0x00000071, 0x31a00000, // 114
                        0xd0895511, 0x0844dda7, 0x00dd17cf, 0x001af1a1, 0x01c141e9, 0x00000201,
                        0x00224185, 0xffddbe7a, 0x00000000, 0x32e00000, // 115
                        0xd0895511, 0x0844ddac, 0x00dd17d0, 0x001af1a1, 0x01c141ed, 0x00000201,
                        0x00224188, 0xffddbe77, 0x00000073, 0x33200000, // 116
                        0xd0895511, 0x0844ddb1, 0x00dd17d1, 0x001af1a1, 0x01c141f1, 0x00000202,
                        0x0022418b, 0xffddbe74, 0x00000000, 0x84600000, // 117
                        0xd0895511, 0x0844ddb6, 0x00dd17d2, 0x001af1a1, 0x01c141f5, 0x00000202,
                        0x0022418e, 0xffddbe71, 0x0844ddb6, 0x85a00000, // 118
                        0xd0895511, 0xfff256aa, 0x00dd17d3, 0x001af1a1, 0x01c141f9, 0x00000202,
                        0x00224191, 0xffddbe6e, 0x00000000, 0x86e00000, // 119
                        0xd0895511, 0xfff256a8, 0x00dd17d4, 0x001af1a1, 0x01c141fd, 0x00000202,
                        0x00224194, 0xffddbe6b, 0x0844ddc0, 0x87200000, // 120
                        0xd0895511, 0x0844ddc5, 0x00dd17d5, 0x001af1a1, 0x01c14201, 0x00000301,
                        0x00224197, 0xffddbe68, 0x00000000, 0x94600000, // 121
                        0xd0895511, 0x0844ddca, 0x00dd17d6, 0x001af1a1, 0x01c14205, 0x00000301,
                        0x0022419a, 0xffddbe65, 0x01c14205, 0x95a00000, // 122
                        0xd0895511, 0x0844ddcf, 0x00dd17d7, 0x001af1a1, 0xfff256a2, 0x00000301,
                        0x0022419d, 0xffddbe62, 0x00000000, 0x96e00000, // 123
                        0xd0895511, 0x0844ddd4, 0x00dd17d8, 0x001af1a1, 0xfff256a0, 0x00000301,
                        0x002241a0, 0xffddbe5f, 0x01c1420d, 0x97200000, // 124
                        0xd0895511, 0x0844ddd9, 0x00dd17d9, 0x001af1a1, 0x01c14211, 0x01000201,
                        0x002241a3, 0xffddbe5c, 0x00000000, 0xa4600000, // 125
                        0xd0895511, 0x0844ddde, 0x00dd17da, 0x001af1a1, 0x01c14215, 0x01000201,
                        0x002241a6, 0xffddbe59, 0x0000007d, 0xa5a00000, // 126
                        0xd0895511, 0x0844dde3, 0x00dd17db, 0x001af1a1, 0x01c14219, 0x01000201,
                        0x002241a9, 0xffddbe56, 0x00000000, 0xa6e00000, // 127
                        0xd0895511, 0x0844dde8, 0x00dd17dc, 0x001af1a1, 0x01c1421d, 0x01000201,
                        0x002241ac, 0xffddbe53, 0x0000007f, 0xa7200000, // 128
                        0xd0895511, 0x0844dded, 0x00dd17dd, 0x001af1a1, 0x01c14221, 0x01000201,
                        0x002241af, 0xffddbe50, 0x00000000, 0xb4600000, // 129
                        0xd0895511, 0x0844ddf2, 0x00dd17de, 0x001af1a1, 0x01c14225, 0x01000201,
                        0x002241b2, 0xffddbe4d, 0x00000081, 0xb5a00000, // 130
                        0xd0895511, 0x0844ddf7, 0x00dd17df, 0x001af1a1, 0x01c14229, 0x01000201,
                        0x002241b5, 0xffddbe4a, 0x00000000, 0xb6e00000, // 131
                        0xd0895511, 0x0844ddfc, 0x00dd17e0, 0x001af1a1, 0x01c1422d, 0x01000201,
                        0x002241b8, 0xffddbe47, 0x00000083, 0xb7200000, // 132
                        0xd0895511, 0x0844de01, 0x00dd17e1, 0x001af1a1, 0x01c14231, 0x00000200,
                        0x002241bb, 0xffddbe44, 0x00000000, 0x88600000, // 133
                        0xd0895511, 0x0844de06, 0x00dd17e2, 0x001af1a1, 0x01c14235, 0x00000200,
                        0x002241be, 0xffddbe41, 0x0844de06, 0x89a00000, // 134
                        0xd0895511, 0xfff2568a, 0x00dd17e3, 0x001af1a1, 0x01c14239, 0x00000200,
                        0x002241c1, 0xffddbe3e, 0x00000000, 0x8ae00000, // 135
                        0xd0895511, 0xfff25688, 0x00dd17e4, 0x001af1a1, 0x01c1423d, 0x00000200,
                        0x002241c4, 0xffddbe3b, 0x0844de10, 0x8b200000, // 136
                        0xd0895511, 0x0844de15, 0x00dd17e5, 0x001af1a1, 0x01c14241, 0x00000101,
                        0x002241c7, 0xffddbe38, 0x00000000, 0x98600000, // 137
                        0xd0895511, 0x0844de1a, 0x00dd17e6, 0x001af1a1, 0x01c14245, 0x00000101,
                        0x002241ca, 0xffddbe35, 0x01c14245, 0x99a00000, // 138
                        0xd0895511, 0x0844de1f, 0x00dd17e7, 0x001af1a1, 0xfff25682, 0x00000101,
                        0x002241cd, 0xffddbe32, 0x00000000, 0x9ae00000, // 139
                        0xd0895511, 0x0844de24, 0x00dd17e8, 0x001af1a1, 0xfff25680, 0x00000101,
                        0x002241d0, 0xffddbe2f, 0x01c1424d, 0x9b200000, // 140
                        0xd0895511, 0x0844de29, 0x00dd17e9, 0x001af1a1, 0x01c14251, 0x00000201,
                        0x002241d3, 0xffddbe2c, 0x00000000, 0x28600000, // 141
                        0xd0895511, 0x0844de2e, 0x00dd17ea, 0x001af1a1, 0x01c14255, 0x00000201,
                        0x002241d6, 0xffddbe29, 0x0000008d, 0x29a00000, // 142
                        0xd0895511, 0x0844de33, 0x00dd17eb, 0x001af1a1, 0x01c14259, 0x00000201,
                        0x002241d9, 0xffddbe26, 0x00000000, 0x2ae00000, // 143
                        0xd0895511, 0x0844de38, 0x00dd17ec, 0x001af1a1, 0x01c1425d, 0x00000201,
                        0x002241dc, 0xffddbe23, 0x0000008f, 0x2b200000, // 144
                        0xd0895511, 0x0844de3d, 0x00dd17ed, 0x001af1a1, 0x01c14261, 0x00000201,
                        0x002241df, 0xffddbe20, 0x00000000, 0x38600000, // 145
                        0xd0895511, 0x0844de42, 0x00dd17ee, 0x001af1a1, 0x01c14265, 0x00000201,
                        0x002241e2, 0xffddbe1d, 0x00000091, 0x39a00000, // 146
                        0xd0895511, 0x0844de47, 0x00dd17ef, 0x001af1a1, 0x01c14269, 0x00000201,
                        0x002241e5, 0xffddbe1a, 0x00000000, 0x3ae00000, // 147
                        0xd0895511, 0x0844de4c, 0x00dd17f0, 0x001af1a1, 0x01c1426d, 0x00000201,
                        0x002241e8, 0xffddbe17, 0x00000093, 0x3b200000, // 148
                        0xd0895511, 0x0844de51, 0x00dd17f1, 0x001af1a1, 0x01c14271, 0x00000200,
                        0x002241eb, 0xffddbe14, 0x00000000, 0x8c600000, // 149
                        0xd0895511, 0x0844de56, 0x00dd17f2, 0x001af1a1, 0x01c14275, 0x00000200,
                        0x002241ee, 0xffddbe11, 0x0844de56, 0x8da00000, // 150
                        0xd0895511, 0xfff2566a, 0x00dd17f3, 0x001af1a1, 0x01c14279, 0x00000200,
                        0x002241f1, 0xffddbe0e, 0x00000000, 0x8ee00000, // 151
                        0xd0895511, 0xfff25668, 0x00dd17f4, 0x001af1a1, 0x01c1427d, 0x00000200,
                        0x002241f4, 0xffddbe0b, 0x0844de60, 0x8f200000, // 152
                        0xd0895511, 0x0844de65, 0x00dd17f5, 0x001af1a1, 0x01c14281, 0x00000101,
                        0x002241f7, 0xffddbe08, 0x00000000, 0x9c600000, // 153
                        0xd0895511, 0x0844de6a, 0x00dd17f6, 0x001af1a1, 0x01c14285, 0x00000101,
                        0x002241fa, 0xffddbe05, 0x01c14285, 0x9da00000, // 154
                        0xd0895511, 0x0844de6f, 0x00dd17f7, 0x001af1a1, 0xfff25662, 0x00000101,
                        0x002241fd, 0xffddbe02, 0x00000000, 0x9ee00000, // 155
                        0xd0895511, 0x0844de74, 0x00dd17f8, 0x001af1a1, 0xfff25660, 0x00000101,
                        0x00224200, 0xffddbdff, 0x01c1428d, 0x9f200000, // 156
                        0xd0895511, 0x0844de79, 0x00dd17f9, 0x001af1a1, 0x01c14291, 0x00000201,
                        0x00224203, 0xffddbdfc, 0x00000000, 0x2c600000, // 157
                        0xd0895511, 0x0844de7e, 0x00dd17fa, 0x001af1a1, 0x01c14295, 0x00000201,
                        0x00224206, 0xffddbdf9, 0x0000009d, 0x2da00000, // 158
                        0xd0895511, 0x0844de83, 0x00dd17fb, 0x001af1a1, 0x01c14299, 0x00000201,
                        0x00224209, 0xffddbdf6, 0x00000000, 0x2ee00000, // 159
                        0xd0895511, 0x0844de88, 0x00dd17fc, 0x001af1a1, 0x01c1429d, 0x00000201,
                        0x0022420c, 0xffddbdf3, 0x0000009f, 0x2f200000, // 160
                        0xd0895511, 0x0844de8d, 0x00dd17fd, 0x001af1a1, 0x01c142a1, 0x00000201,
                        0x0022420f, 0xffddbdf0, 0x00000000, 0x3c600000, // 161
                        0xd0895511, 0x0844de92, 0x00dd17fe, 0x001af1a1, 0x01c142a5, 0x00000201,
                        0x00224212, 0xffddbded, 0x000000a1, 0x3da00000, // 162
                        0xd0895511, 0x0844de97, 0x00dd17ff, 0x001af1a1, 0x01c142a9, 0x00000201,
                        0x00224215, 0xffddbdea, 0x00000000, 0x3ee00000, // 163
                        0xd0895511, 0x0844de9c, 0x00dd1800, 0x001af1a1, 0x01c142ad, 0x00000201,
                        0x00224218, 0xffddbde7, 0x000000a3, 0x3f200000,
                    ])
                    .chain(
                        std::iter::repeat([0u32, 0, 0, 0, 0, 0, 0, 0xffffffff, 0, 0])
                            .take(1024 - 64 - 100)
                            .flatten(),
                    )
                    .collect::<Vec<_>>();
                {
                    let buffer = buffer.get();
                    let buffer = buffer.get();
                    for i in 0..buffer.len() {
                        assert_eq!(exp_buffer_data[i], buffer[i], "0: {}", i);
                    }
                }
                // stop machine
                let input_data = {
                    let mut stop_input_state_rest_data = input_state_rest_data.clone();
                    stop_input_state_rest_data[100 + 41] |= 1 << 8;
                    exec.new_data_from_vec(stop_input_state_rest_data.clone())
                };
                let mut input_data = it.transform(&input_data).unwrap();
                let mut buffer = exec.new_data_from_vec(res_output_data.clone());
                let stop_exp_buffer_data = {
                    let mut stop_exp_buffer_data = exp_buffer_data.clone();
                    stop_exp_buffer_data[0] = GLOBAL_STATE_STOP;
                    let entry =
                        &mut stop_exp_buffer_data[64 + 141 * pic.len()..64 + 142 * pic.len()];
                    entry[pic.rest_pos()] |= 1 << (22 + 8); // set stop machine
                    stop_exp_buffer_data
                };
                exec.execute_buffer_single(&mut input_data, 0, &mut buffer)
                    .unwrap();
                {
                    let buffer = buffer.get();
                    let buffer = buffer.get();
                    for i in 0..buffer.len() {
                        assert_eq!(stop_exp_buffer_data[i], buffer[i], "1: {}", i);
                    }
                }
                // illegal data part move
                for (data_kind, value) in [(0, 0x000002ff), (1, 0x00ffff01), (2, 0xff000201)] {
                    let input_data = exec.new_data_from_vec(input_state_rest_data.clone());
                    let mut input_data = it.transform(&input_data).unwrap();
                    let ill_res_output_data = {
                        let mut ill_res_output_data = res_output_data.clone();
                        let entry_id = 100 + 16 + (data_kind << 2);
                        let entry = &mut ill_res_output_data
                            [64 + entry_id * pic.len()..64 + (entry_id + 1) * pic.len()];
                        entry[pic.data_part_pos_pos()] = value;
                        ill_res_output_data
                    };
                    let mut buffer = exec.new_data_from_vec(ill_res_output_data);
                    let ill_exp_buffer_data = {
                        let mut ill_exp_buffer_data = exp_buffer_data.clone();
                        ill_exp_buffer_data[0] = GLOBAL_STATE_ILLEGAL;
                        let entry_id = 100 + 16 + (data_kind << 2);
                        let entry = &mut ill_exp_buffer_data
                            [64 + entry_id * pic.len()..64 + (entry_id + 1) * pic.len()];
                        entry[pic.data_part_pos_pos()] = value;
                        entry[pic.rest_pos()] &= !(1 << (22 + 9));
                        ill_exp_buffer_data
                    };
                    exec.execute_buffer_single(&mut input_data, 0, &mut buffer)
                        .unwrap();
                    {
                        let buffer = buffer.get();
                        let buffer = buffer.get();
                        for i in 0..buffer.len() {
                            assert_eq!(
                                ill_exp_buffer_data[i], buffer[i],
                                "2: {}: {}",
                                data_kind, i
                            );
                        }
                    }
                }
                // illegal data part write
                for (data_kind, value) in [(0, 0x00000202), (1, 0x00000501)] {
                    let input_data = exec.new_data_from_vec(input_state_rest_data.clone());
                    let mut input_data = it.transform(&input_data).unwrap();
                    let ill_res_output_data = {
                        let mut ill_res_output_data = res_output_data.clone();
                        let entry_id = 100 + 16 + (data_kind << 2) + 2;
                        let entry = &mut ill_res_output_data
                            [64 + entry_id * pic.len()..64 + (entry_id + 1) * pic.len()];
                        entry[pic.data_part_pos_pos()] = value;
                        ill_res_output_data
                    };
                    let mut buffer = exec.new_data_from_vec(ill_res_output_data);
                    let ill_exp_buffer_data = {
                        let mut ill_exp_buffer_data = exp_buffer_data.clone();
                        ill_exp_buffer_data[0] = GLOBAL_STATE_ILLEGAL;
                        let entry_id = 100 + 16 + (data_kind << 2) + 2;
                        let orig_entry = &res_output_data
                            [64 + entry_id * pic.len()..64 + (entry_id + 1) * pic.len()];
                        let entry = &mut ill_exp_buffer_data
                            [64 + entry_id * pic.len()..64 + (entry_id + 1) * pic.len()];
                        if data_kind == 0 {
                            entry[1] = orig_entry[1];
                        } else {
                            entry[4] = orig_entry[4];
                        }
                        entry[pic.data_part_pos_pos()] = value;
                        entry[pic.rest_pos()] &= !(1 << (22 + 9));
                        ill_exp_buffer_data
                    };
                    exec.execute_buffer_single(&mut input_data, 0, &mut buffer)
                        .unwrap();
                    {
                        let buffer = buffer.get();
                        let buffer = buffer.get();
                        for i in 0..buffer.len() {
                            assert_eq!(
                                ill_exp_buffer_data[i], buffer[i],
                                "3: {}: {}",
                                data_kind, i
                            );
                        }
                    }
                }
                //
                // 16-bit and 12-bit
                //
                let builder = $bc();
                let config = InfParMachineConfig {
                    state_len: 20,
                    data_part_len: 12,
                    cell_len_bits: 4,
                };
                let env_config = InfParEnvConfig {
                    proc_num: 1024,
                    flat_memory: true,
                    max_temp_buffer_len: 90,
                    max_mem_size: Some(20 * (1 << 42)),
                };
                let pic = ProcIntDataConfig::new(config, env_config);
                println!("PIC2 entry len {}", pic.len());
                let mem_cell_data_part_len = (1 << config.cell_len_bits) + config.data_part_len;
                let circuit_input_len = config.state_len + mem_cell_data_part_len + 1;
                // negate mem cell to write and data part to write to distinguish values from
                // read mem cell and read data part
                let circuit = Circuit::<u32>::new(
                    circuit_input_len,
                    [Gate::new_nor(0, circuit_input_len - 1)],
                    std::iter::once((circuit_input_len, true)).chain(
                        (1..circuit_input_len - 1)
                            .chain(0..9)
                            .map(|i| (i, i >= config.state_len && i < circuit_input_len - 1)),
                    ),
                )
                .unwrap();
                let mut exec = build_circuit(
                    builder,
                    circuit,
                    &pic,
                    usize::try_from(env_config.proc_num).unwrap(),
                )
                .unwrap();
                // res output data
                let res_output_data = std::iter::repeat(0)
                    .take(64 + 100 * pic.len())
                    .chain(
                        (0..64)
                            .map(|k| {
                                [
                                    // entry 0
                                    0xd0895511,                   // memaddr 0
                                    0x844dd61 + 5 * k,            // memaddr 1
                                    0xdd17c1 + k,                 // temp buffer 0
                                    0x1af1a1,                     // temp buffer 1
                                    0x1c141b1 + 4 * k,            // temp buffer 2
                                    0x00000201,                   // data positions
                                    0x6da1bc03 + (0x10001) * k,   // mem read, mem write
                                    (0x931 + 2 * k) | 0x98700000, // data read, rest
                                ]
                            })
                            .flatten(),
                    )
                    .chain(std::iter::repeat(0).take((1024 - 64 - 100) * pic.len()))
                    .collect::<Vec<_>>();
                // execute main test
                let mut it = exec
                    .input_transformer(32, &(0..20).collect::<Vec<_>>())
                    .unwrap();
                let input_data = exec.new_data_from_vec(input_state_rest_data.clone());
                let mut input_data = it.transform(&input_data).unwrap();
                let mut buffer = exec.new_data_from_vec(res_output_data.clone());
                exec.execute_buffer_single(&mut input_data, 0, &mut buffer)
                    .unwrap();
                // expected buffer data
                let exp_buffer_data = std::iter::repeat(0)
                    .take(64)
                    .chain(
                        std::iter::repeat([0u32, 0, 0, 0, 0, 0, 0xffff0000, 0])
                            .take(100)
                            .flatten(),
                    )
                    .chain([
                        // 100
                        0xd0895511, 0x0844dd61, 0x00dd17c1, 0x001af1a1, 0x01c141b1, 0x00000201,
                        0x43fcbc03, 0x00700000, // 101
                        0xd0895511, 0x0844dd66, 0x00dd17c2, 0x001af1a1, 0x01c141b5, 0x00000201,
                        0x43fbbc04, 0x01b00895, // 102
                        0xd06ca511, 0x0844dd6b, 0x00dd17c3, 0x001af1a1, 0x01c141b9, 0x00000201,
                        0x43fabc05, 0x02f00000, // 103
                        0xd06c8511, 0x0844dd70, 0x00dd17c4, 0x001af1a1, 0x01c141bd, 0x00000201,
                        0x43f9bc06, 0x03300895, // 104
                        0xd0895511, 0x0844dd75, 0x00dd17c5, 0x001af1a1, 0x01c141c1, 0x00000201,
                        0x43f8bc07, 0x10700000, // 105
                        0xd0895511, 0x0844dd7a, 0x00dd17c6, 0x001af1a1, 0x01c141c5, 0x00000201,
                        0x43f7bc08, 0x11b00100, // 106
                        0xd0895511, 0x0844dd7f, 0xc2dd17c7, 0x001af1a6, 0x01c141c9, 0x00000201,
                        0x43f6bc09, 0x12f00000, // 107
                        0xd0895511, 0x0844dd84, 0xc0dd17c8, 0x001af1a6, 0x01c141cd, 0x00000201,
                        0x43f5bc0a, 0x13300100, // 108
                        0xd0895511, 0x0844dd89, 0x00dd17c9, 0x001af1a1, 0x01c141d1, 0x00000201,
                        0x43f4bc0b, 0x20700000, // 109
                        0xd0895511, 0x0844dd8e, 0x00dd17ca, 0x001af1a1, 0x01c141d5, 0x00000201,
                        0x43f3bc0c, 0x21b0006d, // 110
                        0xd0895511, 0x0844dd93, 0x00dd17cb, 0x001af1a1, 0x01c141d9, 0x00000201,
                        0x43f2bc0d, 0x22f00000, // 111
                        0xd0895511, 0x0844dd98, 0x00dd17cc, 0x001af1a1, 0x01c141dd, 0x00000201,
                        0x43f1bc0e, 0x2330006f, // 112
                        0xd0895511, 0x0844dd9d, 0x00dd17cd, 0x001af1a1, 0x01c141e1, 0x00000201,
                        0x43f0bc0f, 0x30700000, // 113
                        0xd0895511, 0x0844dda2, 0x00dd17ce, 0x001af1a1, 0x01c141e5, 0x00000201,
                        0x43efbc10, 0x31b00071, // 114
                        0xd0895511, 0x0844dda7, 0x00dd17cf, 0x001af1a1, 0x01c141e9, 0x00000201,
                        0x43eebc11, 0x32f00000, // 115
                        0xd0895511, 0x0844ddac, 0x00dd17d0, 0x001af1a1, 0x01c141ed, 0x00000201,
                        0x43edbc12, 0x33300073, // 116
                        0xd0895511, 0x0844ddb1, 0x00dd17d1, 0x001af1a1, 0x01c141f1, 0x00000202,
                        0x43ecbc13, 0x84700000, // 117
                        0xd0895511, 0x0844ddb6, 0x00dd17d2, 0x001af1a1, 0x01c141f5, 0x00000202,
                        0x43ebbc14, 0x85b00895, // 118
                        0xd06aa511, 0x0844ddbb, 0x00dd17d3, 0x001af1a1, 0x01c141f9, 0x00000202,
                        0x43eabc15, 0x86f00000, // 119
                        0xd06a8511, 0x0844ddc0, 0x00dd17d4, 0x001af1a1, 0x01c141fd, 0x00000202,
                        0x43e9bc16, 0x87300895, // 120
                        0xd0895511, 0x0844ddc5, 0x00dd17d5, 0x001af1a1, 0x01c14201, 0x00000301,
                        0x43e8bc17, 0x94700000, // 121
                        0xd0895511, 0x0844ddca, 0x00dd17d6, 0x001af1a1, 0x01c14205, 0x00000301,
                        0x43e7bc18, 0x95b00100, // 122
                        0xd0895511, 0x0844ddcf, 0xa2dd17d7, 0x001af1a6, 0x01c14209, 0x00000301,
                        0x43e6bc19, 0x96f00000, // 123
                        0xd0895511, 0x0844ddd4, 0xa0dd17d8, 0x001af1a6, 0x01c1420d, 0x00000301,
                        0x43e5bc1a, 0x97300100, // 124
                        0xd0895511, 0x0844ddd9, 0x00dd17d9, 0x001af1a1, 0x01c14211, 0x01000201,
                        0x43e4bc1b, 0xa4700000, // 125
                        0xd0895511, 0x0844ddde, 0x00dd17da, 0x001af1a1, 0x01c14215, 0x01000201,
                        0x43e3bc1c, 0xa5b0007d, // 126
                        0xd0895511, 0x0844dde3, 0x00dd17db, 0x001af1a1, 0x01c14219, 0x01000201,
                        0x43e2bc1d, 0xa6f00000, // 127
                        0xd0895511, 0x0844dde8, 0x00dd17dc, 0x001af1a1, 0x01c1421d, 0x01000201,
                        0x43e1bc1e, 0xa730007f, // 128
                        0xd0895511, 0x0844dded, 0x00dd17dd, 0x001af1a1, 0x01c14221, 0x01000201,
                        0x43e0bc1f, 0xb4700000, // 129
                        0xd0895511, 0x0844ddf2, 0x00dd17de, 0x001af1a1, 0x01c14225, 0x01000201,
                        0x43dfbc20, 0xb5b00081, // 130
                        0xd0895511, 0x0844ddf7, 0x00dd17df, 0x001af1a1, 0x01c14229, 0x01000201,
                        0x43debc21, 0xb6f00000, // 131
                        0xd0895511, 0x0844ddfc, 0x00dd17e0, 0x001af1a1, 0x01c1422d, 0x01000201,
                        0x43ddbc22, 0xb7300083, // 132
                        0xd0895511, 0x0844de01, 0x00dd17e1, 0x001af1a1, 0x01c14231, 0x00000200,
                        0x43dcbc23, 0x88700000, // 133
                        0xd0895511, 0x0844de06, 0x00dd17e2, 0x001af1a1, 0x01c14235, 0x00000200,
                        0x43dbbc24, 0x89b00895, // 134
                        0xd068a511, 0x0844de0b, 0x00dd17e3, 0x001af1a1, 0x01c14239, 0x00000200,
                        0x43dabc25, 0x8af00000, // 135
                        0xd0688511, 0x0844de10, 0x00dd17e4, 0x001af1a1, 0x01c1423d, 0x00000200,
                        0x43d9bc26, 0x8b300895, // 136
                        0xd0895511, 0x0844de15, 0x00dd17e5, 0x001af1a1, 0x01c14241, 0x00000101,
                        0x43d8bc27, 0x98700000, // 137
                        0xd0895511, 0x0844de1a, 0x00dd17e6, 0x001af1a1, 0x01c14245, 0x00000101,
                        0x43d7bc28, 0x99b00100, // 138
                        0xd0895511, 0x0844de1f, 0x82dd17e7, 0x001af1a6, 0x01c14249, 0x00000101,
                        0x43d6bc29, 0x9af00000, // 139
                        0xd0895511, 0x0844de24, 0x80dd17e8, 0x001af1a6, 0x01c1424d, 0x00000101,
                        0x43d5bc2a, 0x9b300100, // 140
                        0xd0895511, 0x0844de29, 0x00dd17e9, 0x001af1a1, 0x01c14251, 0x00000201,
                        0x43d4bc2b, 0x28700000, // 141
                        0xd0895511, 0x0844de2e, 0x00dd17ea, 0x001af1a1, 0x01c14255, 0x00000201,
                        0x43d3bc2c, 0x29b0008d, // 142
                        0xd0895511, 0x0844de33, 0x00dd17eb, 0x001af1a1, 0x01c14259, 0x00000201,
                        0x43d2bc2d, 0x2af00000, // 143
                        0xd0895511, 0x0844de38, 0x00dd17ec, 0x001af1a1, 0x01c1425d, 0x00000201,
                        0x43d1bc2e, 0x2b30008f, // 144
                        0xd0895511, 0x0844de3d, 0x00dd17ed, 0x001af1a1, 0x01c14261, 0x00000201,
                        0x43d0bc2f, 0x38700000, // 145
                        0xd0895511, 0x0844de42, 0x00dd17ee, 0x001af1a1, 0x01c14265, 0x00000201,
                        0x43cfbc30, 0x39b00091, // 146
                        0xd0895511, 0x0844de47, 0x00dd17ef, 0x001af1a1, 0x01c14269, 0x00000201,
                        0x43cebc31, 0x3af00000, // 147
                        0xd0895511, 0x0844de4c, 0x00dd17f0, 0x001af1a1, 0x01c1426d, 0x00000201,
                        0x43cdbc32, 0x3b300093, // 148
                        0xd0895511, 0x0844de51, 0x00dd17f1, 0x001af1a1, 0x01c14271, 0x00000200,
                        0x43ccbc33, 0x8c700000, // 149
                        0xd0895511, 0x0844de56, 0x00dd17f2, 0x001af1a1, 0x01c14275, 0x00000200,
                        0x43cbbc34, 0x8db00895, // 150
                        0xd066a511, 0x0844de5b, 0x00dd17f3, 0x001af1a1, 0x01c14279, 0x00000200,
                        0x43cabc35, 0x8ef00000, // 151
                        0xd0668511, 0x0844de60, 0x00dd17f4, 0x001af1a1, 0x01c1427d, 0x00000200,
                        0x43c9bc36, 0x8f300895, // 152
                        0xd0895511, 0x0844de65, 0x00dd17f5, 0x001af1a1, 0x01c14281, 0x00000101,
                        0x43c8bc37, 0x9c700000, // 153
                        0xd0895511, 0x0844de6a, 0x00dd17f6, 0x001af1a1, 0x01c14285, 0x00000101,
                        0x43c7bc38, 0x9db00100, // 154
                        0xd0895511, 0x0844de6f, 0x62dd17f7, 0x001af1a6, 0x01c14289, 0x00000101,
                        0x43c6bc39, 0x9ef00000, // 155
                        0xd0895511, 0x0844de74, 0x60dd17f8, 0x001af1a6, 0x01c1428d, 0x00000101,
                        0x43c5bc3a, 0x9f300100, // 156
                        0xd0895511, 0x0844de79, 0x00dd17f9, 0x001af1a1, 0x01c14291, 0x00000201,
                        0x43c4bc3b, 0x2c700000, // 157
                        0xd0895511, 0x0844de7e, 0x00dd17fa, 0x001af1a1, 0x01c14295, 0x00000201,
                        0x43c3bc3c, 0x2db0009d, // 158
                        0xd0895511, 0x0844de83, 0x00dd17fb, 0x001af1a1, 0x01c14299, 0x00000201,
                        0x43c2bc3d, 0x2ef00000, // 159
                        0xd0895511, 0x0844de88, 0x00dd17fc, 0x001af1a1, 0x01c1429d, 0x00000201,
                        0x43c1bc3e, 0x2f30009f, // 160
                        0xd0895511, 0x0844de8d, 0x00dd17fd, 0x001af1a1, 0x01c142a1, 0x00000201,
                        0x43c0bc3f, 0x3c700000, // 161
                        0xd0895511, 0x0844de92, 0x00dd17fe, 0x001af1a1, 0x01c142a5, 0x00000201,
                        0x43bfbc40, 0x3db000a1, // 162
                        0xd0895511, 0x0844de97, 0x00dd17ff, 0x001af1a1, 0x01c142a9, 0x00000201,
                        0x43bebc41, 0x3ef00000, // 163
                        0xd0895511, 0x0844de9c, 0x00dd1800, 0x001af1a1, 0x01c142ad, 0x00000201,
                        0x43bdbc42, 0x3f3000a3,
                    ])
                    .chain(
                        std::iter::repeat([0u32, 0, 0, 0, 0, 0, 0xffff0000, 0])
                            .take(1024 - 64 - 100)
                            .flatten(),
                    )
                    .collect::<Vec<_>>();
                //
                {
                    let buffer = buffer.get();
                    let buffer = buffer.get();
                    for i in 0..buffer.len() {
                        assert_eq!(exp_buffer_data[i], buffer[i], "4: {}", i);
                    }
                }
                //
                // 8-bit and 4-bit
                //
                let builder = $bc();
                let config = InfParMachineConfig {
                    state_len: 20,
                    data_part_len: 4,
                    cell_len_bits: 3,
                };
                let env_config = InfParEnvConfig {
                    proc_num: 1024,
                    flat_memory: true,
                    max_temp_buffer_len: 32,
                    max_mem_size: Some(20 * (1 << 42)),
                };
                let pic = ProcIntDataConfig::new(config, env_config);
                println!("PIC3 entry len {}", pic.len());
                let mem_cell_data_part_len = (1 << config.cell_len_bits) + config.data_part_len;
                let circuit_input_len = config.state_len + mem_cell_data_part_len + 1;
                // negate mem cell to write and data part to write to distinguish values from
                // read mem cell and read data part
                let circuit = Circuit::<u32>::new(
                    circuit_input_len,
                    [Gate::new_nor(0, circuit_input_len - 1)],
                    std::iter::once((circuit_input_len, true)).chain(
                        (1..circuit_input_len - 1)
                            .chain(0..9)
                            .map(|i| (i, i >= config.state_len && i < circuit_input_len - 1)),
                    ),
                )
                .unwrap();
                let mut exec = build_circuit(
                    builder,
                    circuit,
                    &pic,
                    usize::try_from(env_config.proc_num).unwrap(),
                )
                .unwrap();
                // res output data
                let res_output_data = std::iter::repeat(0)
                    .take(64 + 100 * pic.len())
                    .chain(
                        (0..64)
                            .map(|k| {
                                [
                                    // entry 0
                                    0xd0895511,                   // memaddr 0
                                    0x844dd61 + 5 * k,            // memaddr 1
                                    0xdd17c1 + k,                 // temp buffer 0
                                    0x01000201,                   // data positions
                                    0x6da1bc03 + (0x10001) * k,   // mem read, mem write
                                    (0x931 + 2 * k) | 0x98700000, // data read, rest
                                ]
                            })
                            .flatten(),
                    )
                    .chain(std::iter::repeat(0).take((1024 - 64 - 100) * pic.len()))
                    .collect::<Vec<_>>();
                // execute main test
                let mut it = exec
                    .input_transformer(32, &(0..20).collect::<Vec<_>>())
                    .unwrap();
                let input_data = exec.new_data_from_vec(input_state_rest_data.clone());
                let mut input_data = it.transform(&input_data).unwrap();
                let mut buffer = exec.new_data_from_vec(res_output_data.clone());
                exec.execute_buffer_single(&mut input_data, 0, &mut buffer)
                    .unwrap();
                // expected buffer data
                let exp_buffer_data = std::iter::repeat(0)
                    .take(64)
                    .chain(
                        std::iter::repeat([0u32, 0, 0, 0, 0x00ff0000, 0])
                            .take(100)
                            .flatten(),
                    )
                    .chain([
                        // 100
                        0xd0895511, 0x0844dd61, 0x00dd17c1, 0x01000201, 0x00fcbc03, 0x00700930,
                        // 101
                        0xd0895511, 0x0844dd66, 0x00dd17c2, 0x01000201, 0x00fbbc04, 0x01b00931,
                        // 102
                        0xd08955a1, 0x0844dd6b, 0x00dd17c3, 0x01000201, 0x00fabc05, 0x02f00930,
                        // 103
                        0xd0895581, 0x0844dd70, 0x00dd17c4, 0x01000201, 0x00f9bc06, 0x03300931,
                        // 104
                        0xd0895511, 0x0844dd75, 0x00dd17c5, 0x01000201, 0x00f8bc07, 0x10700930,
                        // 105
                        0xd0895511, 0x0844dd7a, 0x00dd17c6, 0x01000201, 0x00f7bc08, 0x11b00937,
                        // 106
                        0xd0895511, 0x0844dd7f, 0x00dd12c7, 0x01000201, 0x00f6bc09, 0x12f00930,
                        // 107
                        0xd0895511, 0x0844dd84, 0x00dd10c8, 0x01000201, 0x00f5bc0a, 0x13300937,
                        // 108
                        0xd0895511, 0x0844dd89, 0x00dd17c9, 0x01000201, 0x00f4bc0b, 0x20700940,
                        // 109
                        0xd0895511, 0x0844dd8e, 0x00dd17ca, 0x01000201, 0x00f3bc0c, 0x21b00946,
                        // 110
                        0xd0895511, 0x0844dd93, 0x00dd17cb, 0x01000201, 0x00f2bc0d, 0x22f00940,
                        // 111
                        0xd0895511, 0x0844dd98, 0x00dd17cc, 0x01000201, 0x00f1bc0e, 0x23300946,
                        // 112
                        0xd0895511, 0x0844dd9d, 0x00dd17cd, 0x01000201, 0x00f0bc0f, 0x30700940,
                        // 113
                        0xd0895511, 0x0844dda2, 0x00dd17ce, 0x01000201, 0x00efbc10, 0x31b00947,
                        // 114
                        0xd0895511, 0x0844dda7, 0x00dd17cf, 0x01000201, 0x00eebc11, 0x32f00940,
                        // 115
                        0xd0895511, 0x0844ddac, 0x00dd17d0, 0x01000201, 0x00edbc12, 0x33300947,
                        // 116
                        0xd0895511, 0x0844ddb1, 0x00dd17d1, 0x01000202, 0x00ecbc13, 0x84700950,
                        // 117
                        0xd0895511, 0x0844ddb6, 0x00dd17d2, 0x01000202, 0x00ebbc14, 0x85b00951,
                        // 118
                        0xd08955a1, 0x0844ddbb, 0x00dd17d3, 0x01000202, 0x00eabc15, 0x86f00950,
                        // 119
                        0xd0895581, 0x0844ddc0, 0x00dd17d4, 0x01000202, 0x00e9bc16, 0x87300951,
                        // 120
                        0xd0895511, 0x0844ddc5, 0x00dd17d5, 0x01000301, 0x00e8bc17, 0x94700950,
                        // 121
                        0xd0895511, 0x0844ddca, 0x00dd17d6, 0x01000301, 0x00e7bc18, 0x95b00957,
                        // 122
                        0xd0895511, 0x0844ddcf, 0x00dd12d7, 0x01000301, 0x00e6bc19, 0x96f00950,
                        // 123
                        0xd0895511, 0x0844ddd4, 0x00dd10d8, 0x01000301, 0x00e5bc1a, 0x97300957,
                        // 124
                        0xd0895511, 0x0844ddd9, 0x00dd17d9, 0x02000201, 0x00e4bc1b, 0xa4700960,
                        // 125
                        0xd0895511, 0x0844ddde, 0x00dd17da, 0x02000201, 0x00e3bc1c, 0xa5b00967,
                        // 126
                        0xd0895511, 0x0844dde3, 0x00dd17db, 0x02000201, 0x00e2bc1d, 0xa6f00960,
                        // 127
                        0xd0895511, 0x0844dde8, 0x00dd17dc, 0x02000201, 0x00e1bc1e, 0xa7300967,
                        // 128
                        0xd0895511, 0x0844dded, 0x00dd17dd, 0x02000201, 0x00e0bc1f, 0xb4700960,
                        // 129
                        0xd0895511, 0x0844ddf2, 0x00dd17de, 0x02000201, 0x00dfbc20, 0xb5b00968,
                        // 130
                        0xd0895511, 0x0844ddf7, 0x00dd17df, 0x02000201, 0x00debc21, 0xb6f00960,
                        // 131
                        0xd0895511, 0x0844ddfc, 0x00dd17e0, 0x02000201, 0x00ddbc22, 0xb7300968,
                        // 132
                        0xd0895511, 0x0844de01, 0x00dd17e1, 0x01000200, 0x00dcbc23, 0x88700970,
                        // 133
                        0xd0895511, 0x0844de06, 0x00dd17e2, 0x01000200, 0x00dbbc24, 0x89b00971,
                        // 134
                        0xd08955a1, 0x0844de0b, 0x00dd17e3, 0x01000200, 0x00dabc25, 0x8af00970,
                        // 135
                        0xd0895581, 0x0844de10, 0x00dd17e4, 0x01000200, 0x00d9bc26, 0x8b300971,
                        // 136
                        0xd0895511, 0x0844de15, 0x00dd17e5, 0x01000101, 0x00d8bc27, 0x98700970,
                        // 137
                        0xd0895511, 0x0844de1a, 0x00dd17e6, 0x01000101, 0x00d7bc28, 0x99b00977,
                        // 138
                        0xd0895511, 0x0844de1f, 0x00dd12e7, 0x01000101, 0x00d6bc29, 0x9af00970,
                        // 139
                        0xd0895511, 0x0844de24, 0x00dd10e8, 0x01000101, 0x00d5bc2a, 0x9b300977,
                        // 140
                        0xd0895511, 0x0844de29, 0x00dd17e9, 0x00000201, 0x00d4bc2b, 0xa8700980,
                        // 141
                        0xd0895511, 0x0844de2e, 0x00dd17ea, 0x00000201, 0x00d3bc2c, 0xa9b00988,
                        // 142
                        0xd0895511, 0x0844de33, 0x00dd17eb, 0x00000201, 0x00d2bc2d, 0xaaf00980,
                        // 143
                        0xd0895511, 0x0844de38, 0x00dd17ec, 0x00000201, 0x00d1bc2e, 0xab300988,
                        // 144
                        0xd0895511, 0x0844de3d, 0x00dd17ed, 0x00000201, 0x00d0bc2f, 0xb8700980,
                        // 145
                        0xd0895511, 0x0844de42, 0x00dd17ee, 0x00000201, 0x00cfbc30, 0xb9b00989,
                        // 146
                        0xd0895511, 0x0844de47, 0x00dd17ef, 0x00000201, 0x00cebc31, 0xbaf00980,
                        // 147
                        0xd0895511, 0x0844de4c, 0x00dd17f0, 0x00000201, 0x00cdbc32, 0xbb300989,
                        // 148
                        0xd0895511, 0x0844de51, 0x00dd17f1, 0x01000200, 0x00ccbc33, 0x8c700990,
                        // 149
                        0xd0895511, 0x0844de56, 0x00dd17f2, 0x01000200, 0x00cbbc34, 0x8db00991,
                        // 150
                        0xd08955a1, 0x0844de5b, 0x00dd17f3, 0x01000200, 0x00cabc35, 0x8ef00990,
                        // 151
                        0xd0895581, 0x0844de60, 0x00dd17f4, 0x01000200, 0x00c9bc36, 0x8f300991,
                        // 152
                        0xd0895511, 0x0844de65, 0x00dd17f5, 0x01000101, 0x00c8bc37, 0x9c700990,
                        // 153
                        0xd0895511, 0x0844de6a, 0x00dd17f6, 0x01000101, 0x00c7bc38, 0x9db00997,
                        // 154
                        0xd0895511, 0x0844de6f, 0x00dd12f7, 0x01000101, 0x00c6bc39, 0x9ef00990,
                        // 155
                        0xd0895511, 0x0844de74, 0x00dd10f8, 0x01000101, 0x00c5bc3a, 0x9f300997,
                        // 156
                        0xd0895511, 0x0844de79, 0x00dd17f9, 0x00000201, 0x00c4bc3b, 0xac7009a0,
                        // 157
                        0xd0895511, 0x0844de7e, 0x00dd17fa, 0x00000201, 0x00c3bc3c, 0xadb009a9,
                        // 158
                        0xd0895511, 0x0844de83, 0x00dd17fb, 0x00000201, 0x00c2bc3d, 0xaef009a0,
                        // 159
                        0xd0895511, 0x0844de88, 0x00dd17fc, 0x00000201, 0x00c1bc3e, 0xaf3009a9,
                        // 160
                        0xd0895511, 0x0844de8d, 0x00dd17fd, 0x00000201, 0x00c0bc3f, 0xbc7009a0,
                        // 161
                        0xd0895511, 0x0844de92, 0x00dd17fe, 0x00000201, 0x00bfbc40, 0xbdb009aa,
                        // 162
                        0xd0895511, 0x0844de97, 0x00dd17ff, 0x00000201, 0x00bebc41, 0xbef009a0,
                        // 163
                        0xd0895511, 0x0844de9c, 0x00dd1800, 0x00000201, 0x00bdbc42, 0xbf3009aa,
                    ])
                    .chain(
                        std::iter::repeat([0u32, 0, 0, 0, 0x00ff0000, 0])
                            .take(1024 - 64 - 100)
                            .flatten(),
                    )
                    .collect::<Vec<_>>();
                //
                {
                    let buffer = buffer.get();
                    let buffer = buffer.get();
                    for i in 0..buffer.len() {
                        assert_eq!(exp_buffer_data[i], buffer[i], "5: {}", i);
                    }
                }
                //
                // 128-bit and 48-bit
                //
                let builder = $bc();
                let config = InfParMachineConfig {
                    state_len: 20,
                    data_part_len: 48,
                    cell_len_bits: 7,
                };
                let env_config = InfParEnvConfig {
                    proc_num: 1024,
                    flat_memory: true,
                    max_temp_buffer_len: 160,
                    max_mem_size: Some(20 * (1 << 48)),
                };
                let pic = ProcIntDataConfig::new(config, env_config);
                println!("PIC4 entry len {}", pic.len());
                let mem_cell_data_part_len = (1 << config.cell_len_bits) + config.data_part_len;
                let circuit_input_len = config.state_len + mem_cell_data_part_len + 1;
                // negate mem cell to write and data part to write to distinguish values from
                // read mem cell and read data part
                let circuit = Circuit::<u32>::new(
                    circuit_input_len,
                    [Gate::new_nor(0, circuit_input_len - 1)],
                    std::iter::once((circuit_input_len, true)).chain(
                        (1..circuit_input_len - 1)
                            .chain(0..9)
                            .map(|i| (i, i >= config.state_len && i < circuit_input_len - 1)),
                    ),
                )
                .unwrap();
                let mut exec = build_circuit(
                    builder,
                    circuit,
                    &pic,
                    usize::try_from(env_config.proc_num).unwrap(),
                )
                .unwrap();
                // res output data
                let res_output_data = std::iter::repeat(0)
                    .take(64 + 100 * pic.len())
                    .chain(
                        (0..64)
                            .map(|k| {
                                [
                                    // entry 0
                                    0xd0895511,                        // memaddr 0
                                    0x4844dd61 + 5 * k,                // memaddr 1
                                    0x1add17c1 + k,                    // temp buffer 0
                                    0x6611176e + k,                    // temp buffer 1
                                    0x17601913 + 2 * k,                // temp buffer 2
                                    0x6afd17cd + k,                    // temp buffer 3
                                    0xad0d17c1 + 3 * k,                // temp buffer 4
                                    0x00000201,                        // data positions
                                    0xa9048941 + k,                    // mem read 0
                                    0xa954dda5 + k,                    // mem read 1
                                    0xd5668941 + k,                    // mem read 2
                                    0xc6645555 + k,                    // mem read 3
                                    0x11189410 + k,                    // mem write 0
                                    0x1222dda5 + k,                    // mem write 1
                                    0x23389410 + k,                    // mem write 2
                                    0x33337755 + k,                    // mem write 3
                                    0x10ca0546 + k,                    // data read 0
                                    (0x00005ced + 1 * k) | 0x98700000, // data read 1, rest
                                ]
                            })
                            .flatten(),
                    )
                    .chain(std::iter::repeat(0).take((1024 - 64 - 100) * pic.len()))
                    .collect::<Vec<_>>();
                // execute main test
                let mut it = exec
                    .input_transformer(32, &(0..20).collect::<Vec<_>>())
                    .unwrap();
                let input_data = exec.new_data_from_vec(input_state_rest_data.clone());
                let mut input_data = it.transform(&input_data).unwrap();
                let mut buffer = exec.new_data_from_vec(res_output_data.clone());
                exec.execute_buffer_single(&mut input_data, 0, &mut buffer)
                    .unwrap();
                // expected buffer data
                let exp_buffer_data = std::iter::once(2)
                    .chain(std::iter::repeat(0).take(64 - 1))
                    .chain(
                        std::iter::repeat([
                            0u32, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0xffffffff, 0xffffffff,
                            0xffffffff, 0xffffffff, 0, 0,
                        ])
                        .take(100)
                        .flatten(),
                    )
                    .chain([
                        // 100
                        0xd0895511, 0x4844dd61, 0x1add17c1, 0x6611176e, 0x17601913, 0x6afd17cd,
                        0xad0d17c1, 0x00000201, 0xa9048941, 0xa954dda5, 0xd5668941, 0xc6645555,
                        0x56fb76be, 0x56ab225a, 0x2a9976be, 0x399baaaa, 0x00000000,
                        0x00700000, // 101
                        0xd0895511, 0x4844dd66, 0x1add17c2, 0x6611176f, 0x17601915, 0x6afd17ce,
                        0xad0d17c4, 0x00000201, 0xa9048942, 0xa954dda6, 0xd5668942, 0xc6645556,
                        0x56fb76bd, 0x56ab2259, 0x2a9976bd, 0x399baaa9, 0x00004844,
                        0x01b00000, // 102
                        0xd0895511, 0x4844dd6b, 0x1add17c3, 0x66111770, 0x17601917, 0x6afd17cf,
                        0xad0d17c7, 0x00000201, 0xa9048943, 0xa954dda7, 0xd5668943, 0xc6645557,
                        0x56fb76bc, 0x56ab2258, 0x2a9976bc, 0x399baaa8, 0x00000000,
                        0x02f00000, // 103
                        0xd0895511, 0x4844dd70, 0x1add17c4, 0x66111771, 0x17601919, 0x6afd17d0,
                        0xad0d17ca, 0x00000201, 0xa9048944, 0xa954dda8, 0xd5668944, 0xc6645558,
                        0x56fb76bb, 0x56ab2257, 0x2a9976bb, 0x399baaa7, 0x00004844,
                        0x03300000, // 104
                        0xd0895511, 0x4844dd75, 0x1add17c5, 0x66111772, 0x1760191b, 0x6afd17d1,
                        0xad0d17cd, 0x00000201, 0xa9048945, 0xa954dda9, 0xd5668945, 0xc6645559,
                        0x56fb76ba, 0x56ab2256, 0x2a9976ba, 0x399baaa6, 0x00000000,
                        0x10700000, // 105
                        0xd0895511, 0x4844dd7a, 0x1add17c6, 0x66111773, 0x1760191d, 0x6afd17d2,
                        0xad0d17d0, 0x00000201, 0xa9048946, 0xa954ddaa, 0xd5668946, 0xc664555a,
                        0x56fb76b9, 0x56ab2255, 0x2a9976b9, 0x399baaa5, 0x6afd17d2,
                        0x11b017d0, // 106
                        0xd0895511, 0x4844dd7f, 0x1add17c7, 0x66111774, 0x1760191f, 0xef35fab3,
                        0xad0da30c, 0x00000201, 0xa9048947, 0xa954ddab, 0xd5668947, 0xc664555b,
                        0x56fb76b8, 0x56ab2254, 0x2a9976b8, 0x399baaa4, 0x00000000,
                        0x12f00000, // 107
                        0xd0895511, 0x4844dd84, 0x1add17c8, 0x66111775, 0x17601921, 0xef35fab2,
                        0xad0da30b, 0x00000201, 0xa9048948, 0xa954ddac, 0xd5668948, 0xc664555c,
                        0x56fb76b7, 0x56ab2253, 0x2a9976b7, 0x399baaa3, 0x6afd17d4,
                        0x133017d6, // 108
                        0xd0895511, 0x4844dd89, 0x1add17c9, 0x66111776, 0x17601923, 0x6afd17d5,
                        0xad0d17d9, 0x00000201, 0xa9048949, 0xa954ddad, 0xd5668949, 0xc664555d,
                        0x56fb76b6, 0x56ab2252, 0x2a9976b6, 0x399baaa2, 0x00000000,
                        0x20700000, // 109
                        0xd0895511, 0x4844dd8e, 0x1add17ca, 0x66111777, 0x17601925, 0x6afd17d6,
                        0xad0d17dc, 0x00000201, 0xa904894a, 0xa954ddae, 0xd566894a, 0xc664555e,
                        0x56fb76b5, 0x56ab2251, 0x2a9976b5, 0x399baaa1, 0x0000006d,
                        0x21b00000, // 110
                        0xd0895511, 0x4844dd93, 0x1add17cb, 0x66111778, 0x17601927, 0x6afd17d7,
                        0xad0d17df, 0x00000201, 0xa904894b, 0xa954ddaf, 0xd566894b, 0xc664555f,
                        0x56fb76b4, 0x56ab2250, 0x2a9976b4, 0x399baaa0, 0x00000000,
                        0x22f00000, // 111
                        0xd0895511, 0x4844dd98, 0x1add17cc, 0x66111779, 0x17601929, 0x6afd17d8,
                        0xad0d17e2, 0x00000201, 0xa904894c, 0xa954ddb0, 0xd566894c, 0xc6645560,
                        0x56fb76b3, 0x56ab224f, 0x2a9976b3, 0x399baa9f, 0x0000006f,
                        0x23300000, // 112
                        0xd0895511, 0x4844dd9d, 0x1add17cd, 0x6611177a, 0x1760192b, 0x6afd17d9,
                        0xad0d17e5, 0x00000201, 0xa904894d, 0xa954ddb1, 0xd566894d, 0xc6645561,
                        0x56fb76b2, 0x56ab224e, 0x2a9976b2, 0x399baa9e, 0x00000000,
                        0x30700000, // 113
                        0xd0895511, 0x4844dda2, 0x1add17ce, 0x6611177b, 0x1760192d, 0x6afd17da,
                        0xad0d17e8, 0x00000201, 0xa904894e, 0xa954ddb2, 0xd566894e, 0xc6645562,
                        0x56fb76b1, 0x56ab224d, 0x2a9976b1, 0x399baa9d, 0x00000071,
                        0x31b00000, // 114
                        0xd0895511, 0x4844dda7, 0x1add17cf, 0x6611177c, 0x1760192f, 0x6afd17db,
                        0xad0d17eb, 0x00000201, 0xa904894f, 0xa954ddb3, 0xd566894f, 0xc6645563,
                        0x56fb76b0, 0x56ab224c, 0x2a9976b0, 0x399baa9c, 0x00000000,
                        0x32f00000, // 115
                        0xd0895511, 0x4844ddac, 0x1add17d0, 0x6611177d, 0x17601931, 0x6afd17dc,
                        0xad0d17ee, 0x00000201, 0xa9048950, 0xa954ddb4, 0xd5668950, 0xc6645564,
                        0x56fb76af, 0x56ab224b, 0x2a9976af, 0x399baa9b, 0x00000073,
                        0x33300000, // 116
                        0xd0895511, 0x4844ddb1, 0x1add17d1, 0x6611177e, 0x17601933, 0x6afd17dd,
                        0xad0d17f1, 0x00000202, 0xa9048951, 0xa954ddb5, 0xd5668951, 0xc6645565,
                        0x56fb76ae, 0x56ab224a, 0x2a9976ae, 0x399baa9a, 0x00000000,
                        0x84700000, // 117
                        0xd0895511, 0x4844ddb6, 0x1add17d2, 0x6611177f, 0x17601935, 0x6afd17de,
                        0xad0d17f4, 0x00000202, 0xa9048952, 0xa954ddb6, 0xd5668952, 0xc6645566,
                        0x56fb76ad, 0x56ab2249, 0x2a9976ad, 0x399baa99, 0x00004844,
                        0x85b00000, // 118
                        0xd0895511, 0x4844ddbb, 0x1add17d3, 0x66111780, 0x17601937, 0x6afd17df,
                        0xad0d17f7, 0x00000201, 0xa9048953, 0xa954ddb7, 0xd5668953, 0xc6645567,
                        0x56fb76ac, 0x56ab2248, 0x2a9976ac, 0x399baa98, 0x00000000,
                        0x06f00000, // 119
                        0xd0895511, 0x4844ddc0, 0x1add17d4, 0x66111781, 0x17601939, 0x6afd17e0,
                        0xad0d17fa, 0x00000201, 0xa9048954, 0xa954ddb8, 0xd5668954, 0xc6645568,
                        0x56fb76ab, 0x56ab2247, 0x2a9976ab, 0x399baa97, 0x00004844,
                        0x07300000, // 120
                        0xd0895511, 0x4844ddc5, 0x1add17d5, 0x66111782, 0x1760193b, 0x6afd17e1,
                        0xad0d17fd, 0x00000301, 0xa9048955, 0xa954ddb9, 0xd5668955, 0xc6645569,
                        0x56fb76aa, 0x56ab2246, 0x2a9976aa, 0x399baa96, 0x00000000,
                        0x94700000, // 121
                        0xd0895511, 0x4844ddca, 0x1add17d6, 0x66111783, 0x1760193d, 0x6afd17e2,
                        0xad0d1800, 0x00000301, 0xa9048956, 0xa954ddba, 0xd5668956, 0xc664556a,
                        0x56fb76a9, 0x56ab2245, 0x2a9976a9, 0x399baa95, 0x6afd17e2,
                        0x95b01800, // 122
                        0xd0895511, 0x4844ddcf, 0x1add17d7, 0x66111784, 0x1760193f, 0xef35faa3,
                        0xad0da2fc, 0x00000301, 0xa9048957, 0xa954ddbb, 0xd5668957, 0xc664556b,
                        0x56fb76a8, 0x56ab2244, 0x2a9976a8, 0x399baa94, 0x00000000,
                        0x96f00000, // 123
                        0xd0895511, 0x4844ddd4, 0x1add17d8, 0x66111785, 0x17601941, 0xef35faa2,
                        0xad0da2fb, 0x00000301, 0xa9048958, 0xa954ddbc, 0xd5668958, 0xc664556c,
                        0x56fb76a7, 0x56ab2243, 0x2a9976a7, 0x399baa93, 0x6afd17e4,
                        0x97301806, // 124
                        0xd0895511, 0x4844ddd9, 0x1add17d9, 0x66111786, 0x17601943, 0x6afd17e5,
                        0xad0d1809, 0x01000201, 0xa9048959, 0xa954ddbd, 0xd5668959, 0xc664556d,
                        0x56fb76a6, 0x56ab2242, 0x2a9976a6, 0x399baa92, 0x00000000,
                        0xa4700000, // 125
                        0xd0895511, 0x4844ddde, 0x1add17da, 0x66111787, 0x17601945, 0x6afd17e6,
                        0xad0d180c, 0x01000201, 0xa904895a, 0xa954ddbe, 0xd566895a, 0xc664556e,
                        0x56fb76a5, 0x56ab2241, 0x2a9976a5, 0x399baa91, 0x0000007d,
                        0xa5b00000, // 126
                        0xd0895511, 0x4844dde3, 0x1add17db, 0x66111788, 0x17601947, 0x6afd17e7,
                        0xad0d180f, 0x01000201, 0xa904895b, 0xa954ddbf, 0xd566895b, 0xc664556f,
                        0x56fb76a4, 0x56ab2240, 0x2a9976a4, 0x399baa90, 0x00000000,
                        0xa6f00000, // 127
                        0xd0895511, 0x4844dde8, 0x1add17dc, 0x66111789, 0x17601949, 0x6afd17e8,
                        0xad0d1812, 0x01000201, 0xa904895c, 0xa954ddc0, 0xd566895c, 0xc6645570,
                        0x56fb76a3, 0x56ab223f, 0x2a9976a3, 0x399baa8f, 0x0000007f,
                        0xa7300000, // 128
                        0xd0895511, 0x4844dded, 0x1add17dd, 0x6611178a, 0x1760194b, 0x6afd17e9,
                        0xad0d1815, 0x01000201, 0xa904895d, 0xa954ddc1, 0xd566895d, 0xc6645571,
                        0x56fb76a2, 0x56ab223e, 0x2a9976a2, 0x399baa8e, 0x00000000,
                        0xb4700000, // 129
                        0xd0895511, 0x4844ddf2, 0x1add17de, 0x6611178b, 0x1760194d, 0x6afd17ea,
                        0xad0d1818, 0x01000201, 0xa904895e, 0xa954ddc2, 0xd566895e, 0xc6645572,
                        0x56fb76a1, 0x56ab223d, 0x2a9976a1, 0x399baa8d, 0x00000081,
                        0xb5b00000, // 130
                        0xd0895511, 0x4844ddf7, 0x1add17df, 0x6611178c, 0x1760194f, 0x6afd17eb,
                        0xad0d181b, 0x01000201, 0xa904895f, 0xa954ddc3, 0xd566895f, 0xc6645573,
                        0x56fb76a0, 0x56ab223c, 0x2a9976a0, 0x399baa8c, 0x00000000,
                        0xb6f00000, // 131
                        0xd0895511, 0x4844ddfc, 0x1add17e0, 0x6611178d, 0x17601951, 0x6afd17ec,
                        0xad0d181e, 0x01000201, 0xa9048960, 0xa954ddc4, 0xd5668960, 0xc6645574,
                        0x56fb769f, 0x56ab223b, 0x2a99769f, 0x399baa8b, 0x00000083,
                        0xb7300000, // 132
                        0xd0895511, 0x4844de01, 0x1add17e1, 0x6611178e, 0x17601953, 0x6afd17ed,
                        0xad0d1821, 0x00000200, 0xa9048961, 0xa954ddc5, 0xd5668961, 0xc6645575,
                        0x56fb769e, 0x56ab223a, 0x2a99769e, 0x399baa8a, 0x00000000,
                        0x88700000, // 133
                        0xd0895511, 0x4844de06, 0x1add17e2, 0x6611178f, 0x17601955, 0x6afd17ee,
                        0xad0d1824, 0x00000200, 0xa9048962, 0xa954ddc6, 0xd5668962, 0xc6645576,
                        0x56fb769d, 0x56ab2239, 0x2a99769d, 0x399baa89, 0x00004844,
                        0x89b00000, // 134
                        0xd0895511, 0x4844de0b, 0x1add17e3, 0x66111790, 0x17601957, 0x6afd17ef,
                        0xad0d1827, 0x00000201, 0xa9048963, 0xa954ddc7, 0xd5668963, 0xc6645577,
                        0x56fb769c, 0x56ab2238, 0x2a99769c, 0x399baa88, 0x00000000,
                        0x0af00000, // 135
                        0xd0895511, 0x4844de10, 0x1add17e4, 0x66111791, 0x17601959, 0x6afd17f0,
                        0xad0d182a, 0x00000201, 0xa9048964, 0xa954ddc8, 0xd5668964, 0xc6645578,
                        0x56fb769b, 0x56ab2237, 0x2a99769b, 0x399baa87, 0x00004844,
                        0x0b300000, // 136
                        0xd0895511, 0x4844de15, 0x1add17e5, 0x66111792, 0x1760195b, 0x6afd17f1,
                        0xad0d182d, 0x00000101, 0xa9048965, 0xa954ddc9, 0xd5668965, 0xc6645579,
                        0x56fb769a, 0x56ab2236, 0x2a99769a, 0x399baa86, 0x00000000,
                        0x98700000, // 137
                        0xd0895511, 0x4844de1a, 0x1add17e6, 0x66111793, 0x1760195d, 0x6afd17f2,
                        0xad0d1830, 0x00000101, 0xa9048966, 0xa954ddca, 0xd5668966, 0xc664557a,
                        0x56fb7699, 0x56ab2235, 0x2a997699, 0x399baa85, 0x6afd17f2,
                        0x99b01830, // 138
                        0xd0895511, 0x4844de1f, 0x1add17e7, 0x66111794, 0x1760195f, 0xef35fa93,
                        0xad0da2ec, 0x00000101, 0xa9048967, 0xa954ddcb, 0xd5668967, 0xc664557b,
                        0x56fb7698, 0x56ab2234, 0x2a997698, 0x399baa84, 0x00000000,
                        0x9af00000, // 139
                        0xd0895511, 0x4844de24, 0x1add17e8, 0x66111795, 0x17601961, 0xef35fa92,
                        0xad0da2eb, 0x00000101, 0xa9048968, 0xa954ddcc, 0xd5668968, 0xc664557c,
                        0x56fb7697, 0x56ab2233, 0x2a997697, 0x399baa83, 0x6afd17f4,
                        0x9b301836, // 140
                        0xd0895511, 0x4844de29, 0x1add17e9, 0x66111796, 0x17601963, 0x6afd17f5,
                        0xad0d1839, 0x00000201, 0xa9048969, 0xa954ddcd, 0xd5668969, 0xc664557d,
                        0x56fb7696, 0x56ab2232, 0x2a997696, 0x399baa82, 0x00000000,
                        0x28700000, // 141
                        0xd0895511, 0x4844de2e, 0x1add17ea, 0x66111797, 0x17601965, 0x6afd17f6,
                        0xad0d183c, 0x00000201, 0xa904896a, 0xa954ddce, 0xd566896a, 0xc664557e,
                        0x56fb7695, 0x56ab2231, 0x2a997695, 0x399baa81, 0x0000008d,
                        0x29b00000, // 142
                        0xd0895511, 0x4844de33, 0x1add17eb, 0x66111798, 0x17601967, 0x6afd17f7,
                        0xad0d183f, 0x00000201, 0xa904896b, 0xa954ddcf, 0xd566896b, 0xc664557f,
                        0x56fb7694, 0x56ab2230, 0x2a997694, 0x399baa80, 0x00000000,
                        0x2af00000, // 143
                        0xd0895511, 0x4844de38, 0x1add17ec, 0x66111799, 0x17601969, 0x6afd17f8,
                        0xad0d1842, 0x00000201, 0xa904896c, 0xa954ddd0, 0xd566896c, 0xc6645580,
                        0x56fb7693, 0x56ab222f, 0x2a997693, 0x399baa7f, 0x0000008f,
                        0x2b300000, // 144
                        0xd0895511, 0x4844de3d, 0x1add17ed, 0x6611179a, 0x1760196b, 0x6afd17f9,
                        0xad0d1845, 0x00000201, 0xa904896d, 0xa954ddd1, 0xd566896d, 0xc6645581,
                        0x56fb7692, 0x56ab222e, 0x2a997692, 0x399baa7e, 0x00000000,
                        0x38700000, // 145
                        0xd0895511, 0x4844de42, 0x1add17ee, 0x6611179b, 0x1760196d, 0x6afd17fa,
                        0xad0d1848, 0x00000201, 0xa904896e, 0xa954ddd2, 0xd566896e, 0xc6645582,
                        0x56fb7691, 0x56ab222d, 0x2a997691, 0x399baa7d, 0x00000091,
                        0x39b00000, // 146
                        0xd0895511, 0x4844de47, 0x1add17ef, 0x6611179c, 0x1760196f, 0x6afd17fb,
                        0xad0d184b, 0x00000201, 0xa904896f, 0xa954ddd3, 0xd566896f, 0xc6645583,
                        0x56fb7690, 0x56ab222c, 0x2a997690, 0x399baa7c, 0x00000000,
                        0x3af00000, // 147
                        0xd0895511, 0x4844de4c, 0x1add17f0, 0x6611179d, 0x17601971, 0x6afd17fc,
                        0xad0d184e, 0x00000201, 0xa9048970, 0xa954ddd4, 0xd5668970, 0xc6645584,
                        0x56fb768f, 0x56ab222b, 0x2a99768f, 0x399baa7b, 0x00000093,
                        0x3b300000, // 148
                        0xd0895511, 0x4844de51, 0x1add17f1, 0x6611179e, 0x17601973, 0x6afd17fd,
                        0xad0d1851, 0x00000200, 0xa9048971, 0xa954ddd5, 0xd5668971, 0xc6645585,
                        0x56fb768e, 0x56ab222a, 0x2a99768e, 0x399baa7a, 0x00000000,
                        0x8c700000, // 149
                        0xd0895511, 0x4844de56, 0x1add17f2, 0x6611179f, 0x17601975, 0x6afd17fe,
                        0xad0d1854, 0x00000200, 0xa9048972, 0xa954ddd6, 0xd5668972, 0xc6645586,
                        0x56fb768d, 0x56ab2229, 0x2a99768d, 0x399baa79, 0x00004844,
                        0x8db00000, // 150
                        0xd0895511, 0x4844de5b, 0x1add17f3, 0x661117a0, 0x17601977, 0x6afd17ff,
                        0xad0d1857, 0x00000201, 0xa9048973, 0xa954ddd7, 0xd5668973, 0xc6645587,
                        0x56fb768c, 0x56ab2228, 0x2a99768c, 0x399baa78, 0x00000000,
                        0x0ef00000, // 151
                        0xd0895511, 0x4844de60, 0x1add17f4, 0x661117a1, 0x17601979, 0x6afd1800,
                        0xad0d185a, 0x00000201, 0xa9048974, 0xa954ddd8, 0xd5668974, 0xc6645588,
                        0x56fb768b, 0x56ab2227, 0x2a99768b, 0x399baa77, 0x00004844,
                        0x0f300000, // 152
                        0xd0895511, 0x4844de65, 0x1add17f5, 0x661117a2, 0x1760197b, 0x6afd1801,
                        0xad0d185d, 0x00000101, 0xa9048975, 0xa954ddd9, 0xd5668975, 0xc6645589,
                        0x56fb768a, 0x56ab2226, 0x2a99768a, 0x399baa76, 0x00000000,
                        0x9c700000, // 153
                        0xd0895511, 0x4844de6a, 0x1add17f6, 0x661117a3, 0x1760197d, 0x6afd1802,
                        0xad0d1860, 0x00000101, 0xa9048976, 0xa954ddda, 0xd5668976, 0xc664558a,
                        0x56fb7689, 0x56ab2225, 0x2a997689, 0x399baa75, 0x6afd1802,
                        0x9db01860, // 154
                        0xd0895511, 0x4844de6f, 0x1add17f7, 0x661117a4, 0x1760197f, 0xef35fa83,
                        0xad0da2dc, 0x00000101, 0xa9048977, 0xa954dddb, 0xd5668977, 0xc664558b,
                        0x56fb7688, 0x56ab2224, 0x2a997688, 0x399baa74, 0x00000000,
                        0x9ef00000, // 155
                        0xd0895511, 0x4844de74, 0x1add17f8, 0x661117a5, 0x17601981, 0xef35fa82,
                        0xad0da2db, 0x00000101, 0xa9048978, 0xa954dddc, 0xd5668978, 0xc664558c,
                        0x56fb7687, 0x56ab2223, 0x2a997687, 0x399baa73, 0x6afd1804,
                        0x9f301866, // 156
                        0xd0895511, 0x4844de79, 0x1add17f9, 0x661117a6, 0x17601983, 0x6afd1805,
                        0xad0d1869, 0x00000201, 0xa9048979, 0xa954dddd, 0xd5668979, 0xc664558d,
                        0x56fb7686, 0x56ab2222, 0x2a997686, 0x399baa72, 0x00000000,
                        0x2c700000, // 157
                        0xd0895511, 0x4844de7e, 0x1add17fa, 0x661117a7, 0x17601985, 0x6afd1806,
                        0xad0d186c, 0x00000201, 0xa904897a, 0xa954ddde, 0xd566897a, 0xc664558e,
                        0x56fb7685, 0x56ab2221, 0x2a997685, 0x399baa71, 0x0000009d,
                        0x2db00000, // 158
                        0xd0895511, 0x4844de83, 0x1add17fb, 0x661117a8, 0x17601987, 0x6afd1807,
                        0xad0d186f, 0x00000201, 0xa904897b, 0xa954dddf, 0xd566897b, 0xc664558f,
                        0x56fb7684, 0x56ab2220, 0x2a997684, 0x399baa70, 0x00000000,
                        0x2ef00000, // 159
                        0xd0895511, 0x4844de88, 0x1add17fc, 0x661117a9, 0x17601989, 0x6afd1808,
                        0xad0d1872, 0x00000201, 0xa904897c, 0xa954dde0, 0xd566897c, 0xc6645590,
                        0x56fb7683, 0x56ab221f, 0x2a997683, 0x399baa6f, 0x0000009f,
                        0x2f300000, // 160
                        0xd0895511, 0x4844de8d, 0x1add17fd, 0x661117aa, 0x1760198b, 0x6afd1809,
                        0xad0d1875, 0x00000201, 0xa904897d, 0xa954dde1, 0xd566897d, 0xc6645591,
                        0x56fb7682, 0x56ab221e, 0x2a997682, 0x399baa6e, 0x00000000,
                        0x3c700000, // 161
                        0xd0895511, 0x4844de92, 0x1add17fe, 0x661117ab, 0x1760198d, 0x6afd180a,
                        0xad0d1878, 0x00000201, 0xa904897e, 0xa954dde2, 0xd566897e, 0xc6645592,
                        0x56fb7681, 0x56ab221d, 0x2a997681, 0x399baa6d, 0x000000a1,
                        0x3db00000, // 162
                        0xd0895511, 0x4844de97, 0x1add17ff, 0x661117ac, 0x1760198f, 0x6afd180b,
                        0xad0d187b, 0x00000201, 0xa904897f, 0xa954dde3, 0xd566897f, 0xc6645593,
                        0x56fb7680, 0x56ab221c, 0x2a997680, 0x399baa6c, 0x00000000,
                        0x3ef00000, // 163
                        0xd0895511, 0x4844de9c, 0x1add1800, 0x661117ad, 0x17601991, 0x6afd180c,
                        0xad0d187e, 0x00000201, 0xa9048980, 0xa954dde4, 0xd5668980, 0xc6645594,
                        0x56fb767f, 0x56ab221b, 0x2a99767f, 0x399baa6b, 0x000000a3, 0x3f300000,
                    ])
                    .chain(
                        std::iter::repeat([
                            0u32, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0xffffffff, 0xffffffff,
                            0xffffffff, 0xffffffff, 0, 0,
                        ])
                        .take(1024 - 64 - 100)
                        .flatten(),
                    )
                    .collect::<Vec<_>>();
                //
                {
                    let buffer = buffer.get();
                    let buffer = buffer.get();
                    for i in 0..buffer.len() {
                        assert_eq!(exp_buffer_data[i], buffer[i], "6: {}", i);
                    }
                }
                //
                // 128-bit and 200-bit
                //
                let builder = $bc();
                let config = InfParMachineConfig {
                    state_len: 20,
                    data_part_len: 200,
                    cell_len_bits: 7,
                };
                let env_config = InfParEnvConfig {
                    proc_num: 1024,
                    flat_memory: true,
                    max_temp_buffer_len: 350,
                    max_mem_size: Some(20 * (1 << 48)),
                };
                let pic = ProcIntDataConfig::new(config, env_config);
                println!("PIC5 entry len {}", pic.len());
                let mem_cell_data_part_len = (1 << config.cell_len_bits) + config.data_part_len;
                let circuit_input_len = config.state_len + mem_cell_data_part_len + 1;
                // negate mem cell to write and data part to write to distinguish values from
                // read mem cell and read data part
                let circuit = Circuit::<u32>::new(
                    circuit_input_len,
                    [Gate::new_nor(0, circuit_input_len - 1)],
                    std::iter::once((circuit_input_len, true)).chain(
                        (1..circuit_input_len - 1)
                            .chain(0..9)
                            .map(|i| (i, i >= config.state_len && i < circuit_input_len - 1)),
                    ),
                )
                .unwrap();
                build_circuit(
                    builder,
                    circuit,
                    &pic,
                    usize::try_from(env_config.proc_num).unwrap(),
                )
                .unwrap();
            }
        };
    }

    test_build_circuit_xxx!(test_build_circuit_cpu_routine, cpu_builder_creator);
    test_build_circuit_xxx!(test_build_circuit_opencl_routine, opencl_builder_creator);

    #[test]
    fn test_build_circuit_cpu() {
        test_build_circuit_cpu_routine();
    }

    #[test]
    fn test_build_circuit_opencl() {
        test_build_circuit_opencl_routine();
    }
}