File size: 1,189 Bytes
1c59946 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 | // 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;
}
}
|