feat(cuda): new fft impl

zama-ai · Aug 29, 2024 · 66a8b9e · 66a8b9e
1 parent 6e2908a
commit 66a8b9e
Show file tree

Hide file tree

Showing 7 changed files with 388 additions and 638 deletions.
diff --git a/backends/tfhe-cuda-backend/cuda/src/crypto/torus.cuh b/backends/tfhe-cuda-backend/cuda/src/crypto/torus.cuh
@@ -19,14 +19,19 @@ __device__ inline void typecast_double_to_torus<uint32_t>(double x,
 template <>
 __device__ inline void typecast_double_to_torus<uint64_t>(double x,
                                                           uint64_t &r) {
-  // The ull intrinsic does not behave in the same way on all architectures and
-  // on some platforms this causes the cmux tree test to fail
-  // Hence the intrinsic is not used here
   uint128 nnnn = make_uint128_from_float(x);
   uint64_t lll = nnnn.lo_;
   r = lll;
 }
 
+template <typename T>
+__device__ inline void typecast_double_round_to_torus(double x, T &r) {
+  double mx = (sizeof(T) == 4) ? 4294967296.0 : 18446744073709551616.0;
+  double frac = x - floor(x);
+  frac *= mx;
+  typecast_double_to_torus(frac, r);
+}
+
 template <typename T>
 __device__ inline T round_to_closest_multiple(T x, uint32_t base_log,
                                               uint32_t level_count) {