// TM reset-per-step kernel.

extern "C" __global__
void tm_reset_step(
    unsigned int * __restrict__ cell_active_bits,
    unsigned int * __restrict__ cell_winner_bits,
    unsigned int * __restrict__ cell_predictive_bits,
    unsigned int * __restrict__ prev_active_bits,
    unsigned int * __restrict__ prev_winner_bits,
    unsigned char * __restrict__ col_predicted,
    unsigned int * __restrict__ unpredicted_count,
    unsigned int * __restrict__ burst_cols_count,
    unsigned int * __restrict__ col_best_match,
    unsigned int   bits_words,
    unsigned int   n_cols
) {
    unsigned int tid_global = blockIdx.x * blockDim.x + threadIdx.x;

    if (tid_global < bits_words) {
        prev_active_bits[tid_global] = cell_active_bits[tid_global];
        prev_winner_bits[tid_global] = cell_winner_bits[tid_global];
        cell_active_bits[tid_global] = 0u;
        cell_winner_bits[tid_global] = 0u;
        cell_predictive_bits[tid_global] = 0u;
    }

    if (tid_global < n_cols) {
        col_predicted[tid_global] = 0;
        col_best_match[tid_global] = 0u;
    }

    if (tid_global == 0) {
        unpredicted_count[0] = 0u;
        burst_cols_count[0] = 0u;
    }
}