| |
|
|
| 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; |
| } |
| } |
|
|
