Skip to content

Commit

Permalink
chore(gpu): init_decomposer_state_inplace calls init_decomposer_state
Browse files Browse the repository at this point in the history
  • Loading branch information
pdroalves committed Nov 4, 2024
1 parent 2136742 commit 7d07fee
Showing 1 changed file with 1 addition and 16 deletions.
17 changes: 1 addition & 16 deletions backends/tfhe-cuda-backend/cuda/src/polynomial/functions.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -148,22 +148,7 @@ __device__ void init_decomposer_state_inplace(T *rotated_acc, int base_log,
int tid = threadIdx.x;
for (int i = 0; i < elems_per_thread; i++) {
T x_acc = rotated_acc_slice[tid];

const T rep_bit_count = level_count * base_log;
const T non_rep_bit_count = sizeof(T) * 8 - rep_bit_count;
T res_acc = x_acc >> (non_rep_bit_count - 1);
T rounding_bit = res_acc & T(1);
res_acc += (T)(1);
res_acc = res_acc >> 1;
T torus_max = scalar_max<T>();
T mod_mask = torus_max >> non_rep_bit_count;
res_acc = res_acc & mod_mask;
T shifted_random = rounding_bit << (rep_bit_count - 1);
T need_balance = (((res_acc - (T)(1)) | shifted_random) & res_acc) >>
(rep_bit_count - 1);
res_acc = res_acc - (need_balance << rep_bit_count);

rotated_acc_slice[tid] = res_acc;
rotated_acc_slice[tid] = init_decomposer_state(x_acc, base_log, level_count);
tid = tid + block_size;
}
}
Expand Down

0 comments on commit 7d07fee

Please sign in to comment.