diff --git a/constantine.nimble b/constantine.nimble
index 18eeb14e..4138b429 100644
--- a/constantine.nimble
+++ b/constantine.nimble
@@ -1,6 +1,6 @@
 packageName   = "constantine"
 version       = "0.0.1"
-author        = "Status Research & Development GmbH"
+author        = "Mamy Ratsimbazafy"
 description   = "This library provides thoroughly tested and highly-optimized implementations of cryptography protocols."
 license       = "MIT or Apache License 2.0"
 
diff --git a/constantine/ethereum_kzg_polynomial_commitments.nim b/constantine/ethereum_kzg_polynomial_commitments.nim
index 835b1a5d..1f1fce63 100644
--- a/constantine/ethereum_kzg_polynomial_commitments.nim
+++ b/constantine/ethereum_kzg_polynomial_commitments.nim
@@ -50,7 +50,7 @@ const BYTES_PER_FIELD_ELEMENT = 32
 # Presets
 # ------------------------------------------------------------
 
-const FIELD_ELEMENTS_PER_BLOB = 4096
+const FIELD_ELEMENTS_PER_BLOB {.intdefine.} = 4096
 const FIAT_SHAMIR_PROTOCOL_DOMAIN = asBytes"FSBLOBVERIFY_V1_"
 const RANDOM_CHALLENGE_KZG_BATCH_DOMAIN = asBytes"RCKZGBATCH___V1_"
 
diff --git a/constantine/math/arithmetic/limbs_exgcd.nim b/constantine/math/arithmetic/limbs_exgcd.nim
index a6445b70..df109798 100644
--- a/constantine/math/arithmetic/limbs_exgcd.nim
+++ b/constantine/math/arithmetic/limbs_exgcd.nim
@@ -742,7 +742,7 @@ func batchedDivsteps_vartime(
 
   while true:
     # Count zeros up to bitsLeft and process a batch of divsteps up to that number
-    let zeros = (g.BaseType or (1.BaseType shl bitsLeft)).countTrailingZeroBits_vartime()
+    let zeros = (BaseType(g) or (BaseType(1) shl bitsLeft)).countTrailingZeroBits_vartime()
     g = g shr zeros
     u = u shl zeros
     v = v shl zeros
diff --git a/constantine/math/elliptic/ec_shortweierstrass_jacobian.nim b/constantine/math/elliptic/ec_shortweierstrass_jacobian.nim
index 8693aa72..ca9f4067 100644
--- a/constantine/math/elliptic/ec_shortweierstrass_jacobian.nim
+++ b/constantine/math/elliptic/ec_shortweierstrass_jacobian.nim
@@ -644,10 +644,17 @@ func `-=`*(P: var ECP_ShortW_Jac, Q: ECP_ShortW_Aff) {.inline.} =
   nQ.neg(Q)
   P.madd(P, nQ)
 
+# Conversions
+# -----------
+
 template affine*[F, G](_: type ECP_ShortW_Jac[F, G]): typedesc =
   ## Returns the affine type that corresponds to the Jacobian type input
   ECP_ShortW_Aff[F, G]
 
+template jacobian*[F, G](_: type ECP_ShortW_Aff[F, G]): typedesc =
+  ## Returns the jacobian type that corresponds to the affine type input
+  ECP_ShortW_Jac[F, G]
+
 func affine*[F; G](
        aff: var ECP_ShortW_Aff[F, G],
        jac: ECP_ShortW_Jac[F, G]) {.meter.} =
diff --git a/constantine/math/elliptic/ec_shortweierstrass_projective.nim b/constantine/math/elliptic/ec_shortweierstrass_projective.nim
index 425f50a5..dea8fb95 100644
--- a/constantine/math/elliptic/ec_shortweierstrass_projective.nim
+++ b/constantine/math/elliptic/ec_shortweierstrass_projective.nim
@@ -439,6 +439,11 @@ template affine*[F, G](_: type ECP_ShortW_Prj[F, G]): typedesc =
   ## Returns the affine type that corresponds to the Jacobian type input
   ECP_ShortW_Aff[F, G]
 
+template projective*[F, G](_: type ECP_ShortW_Aff[F, G]): typedesc =
+  ## Returns the projective type that corresponds to the affine type input
+  ECP_ShortW_Prj[F, G]
+
+
 func affine*[F, G](
        aff: var ECP_ShortW_Aff[F, G],
        proj: ECP_ShortW_Prj[F, G]) {.meter.} =
diff --git a/constantine/math/polynomials/fft.nim b/constantine/math/polynomials/fft.nim
new file mode 100644
index 00000000..e16137b9
--- /dev/null
+++ b/constantine/math/polynomials/fft.nim
@@ -0,0 +1,311 @@
+# Constantine
+# Copyright (c) 2018-2019    Status Research & Development GmbH
+# Copyright (c) 2020-Present Mamy André-Ratsimbazafy
+# Licensed and distributed under either of
+#   * MIT license (license terms in the root directory or at http://opensource.org/licenses/MIT).
+#   * Apache v2 license (license terms in the root directory or at http://www.apache.org/licenses/LICENSE-2.0).
+# at your option. This file may not be copied, modified, or distributed except according to those terms.
+
+import
+  ../config/curves,
+  ../arithmetic,
+  ../ec_shortweierstrass,
+  ../elliptic/ec_scalar_mul_vartime,
+  ../../platforms/[abstractions, allocs, views]
+
+# ############################################################
+#
+#               Fast Fourier Transform
+#
+# ############################################################
+
+# Elliptic curve Fast Fourier Transform
+# ----------------------------------------------------------------
+
+type
+  FFTStatus = enum
+    FFTS_Success
+    FFTS_TooManyValues = "Input length greater than the field 2-adicity (number of roots of unity)"
+    FFTS_SizeNotPowerOfTwo = "Input must be of a power of 2 length"
+
+  ECFFT_Descriptor*[EC] = object
+    ## Metadata for FFT on Elliptic Curve
+    order*: int
+    rootsOfUnity*: ptr UncheckedArray[matchingOrderBigInt(EC.F.C)]
+      ## domain, starting and ending with 1, length cardinality+1
+      ## This allows FFT and inverse FFT to use the same buffer for roots.
+
+func computeRootsOfUnity[EC](ctx: var ECFFT_Descriptor[EC], generatorRootOfUnity: auto) =
+  static: doAssert typeof(generatorRootOfUnity) is Fr[EC.F.C]
+
+  ctx.rootsOfUnity[0].setOne()
+
+  var cur = generatorRootOfUnity
+  for i in 1 .. ctx.order:
+    ctx.rootsOfUnity[i].fromField(cur)
+    cur *= generatorRootOfUnity
+
+  doAssert ctx.rootsOfUnity[ctx.order].isOne().bool()
+
+func new*(T: type ECFFT_Descriptor, order: int, generatorRootOfUnity: auto): T =
+  result.order = order
+  result.rootsOfUnity = allocHeapArrayAligned(matchingOrderBigInt(T.EC.F.C), order+1, alignment = 64)
+
+  result.computeRootsOfUnity(generatorRootOfUnity)
+
+func simpleFT[EC; bits: static int](
+       output: var StridedView[EC],
+       vals: StridedView[EC],
+       rootsOfUnity: StridedView[BigInt[bits]]) =
+  # FFT is a recursive algorithm
+  # This is the base-case using a O(n²) algorithm
+
+  # TODO: endomorphism acceleration for windowed-NAF
+
+  let L = output.len
+  var last {.noInit.}, v {.noInit.}: EC
+
+  for i in 0 ..< L:
+    last = vals[0]
+    last.scalarMul_minHammingWeight_windowed_vartime(rootsOfUnity[0], window = 5)
+    for j in 1 ..< L:
+      v = vals[j]
+      v.scalarMul_minHammingWeight_windowed_vartime(rootsOfUnity[(i*j) mod L], window = 5)
+      last += v
+    output[i] = last
+
+func fft_internal[EC; bits: static int](
+       output: var StridedView[EC],
+       vals: StridedView[EC],
+       rootsOfUnity: StridedView[BigInt[bits]]
+     ) =
+  if output.len <= 4:
+    simpleFT(output, vals, rootsOfUnity)
+    return
+
+  # Recursive Divide-and-Conquer
+  let (evenVals, oddVals) = vals.splitAlternate()
+  var (outLeft, outRight) = output.splitMiddle()
+  let halfROI = rootsOfUnity.skipHalf()
+
+  fft_internal(outLeft, evenVals, halfROI)
+  fft_internal(outRight, oddVals, halfROI)
+
+  let half = outLeft.len
+  var y_times_root{.noinit.}: EC
+
+  for i in 0 ..< half:
+    # FFT Butterfly
+    y_times_root = output[i+half]
+    y_times_root   .scalarMul_minHammingWeight_windowed_vartime(rootsOfUnity[i], window = 5)
+    output[i+half] .diff(output[i], y_times_root)
+    output[i]      += y_times_root
+
+func fft*[EC](
+       desc: ECFFT_Descriptor[EC],
+       output: var openarray[EC],
+       vals: openarray[EC]): FFT_Status =
+  if vals.len > desc.order:
+    return FFTS_TooManyValues
+  if not vals.len.uint64.isPowerOf2_vartime():
+    return FFTS_SizeNotPowerOfTwo
+
+  let rootz = desc.rootsOfUnity
+                  .toStridedView(desc.order)
+                  .slice(0, desc.order-1, desc.order div vals.len)
+
+  var voutput = output.toStridedView()
+  fft_internal(voutput, vals.toStridedView(), rootz)
+  return FFTS_Success
+
+func ifft*[EC](
+       desc: ECFFT_Descriptor[EC],
+       output: var openarray[EC],
+       vals: openarray[EC]): FFT_Status =
+  ## Inverse FFT
+  if vals.len > desc.order:
+    return FFTS_TooManyValues
+  if not vals.len.uint64.isPowerOf2_vartime():
+    return FFTS_SizeNotPowerOfTwo
+
+  let rootz = desc.rootsOfUnity
+                  .toStridedView(desc.order+1) # Extra 1 at the end so that when reversed the buffer starts with 1
+                  .reversed()
+                  .slice(0, desc.order-1, desc.order div vals.len)
+
+  var voutput = output.toStridedView()
+  fft_internal(voutput, vals.toStridedView(), rootz)
+
+  var invLen {.noInit.}: Fr[EC.F.C]
+  invLen.fromUint(vals.len.uint64)
+  invLen.inv_vartime()
+  let inv = invLen.toBig()
+
+  for i in 0 ..< output.len:
+    output[i].scalarMul_minHammingWeight_windowed_vartime(inv, window = 5)
+
+  return FFTS_Success
+
+# ############################################################
+#
+#                   Bit reversal permutations
+#
+# ############################################################
+# - Towards an Optimal Bit-Reversal Permutation Program
+#   Larry Carter and Kang Su Gatlin, 1998
+#   https://csaws.cs.technion.ac.il/~itai/Courses/Cache/bit.pdf
+#
+# - Practically efficient methods for performing bit-reversed
+#   permutation in C++11 on the x86-64 architecture
+#   Knauth, Adas, Whitfield, Wang, Ickler, Conrad, Serang, 2017
+#   https://arxiv.org/pdf/1708.01873.pdf
+
+func deriveLogTileSize(T: typedesc): int =
+  ## Returns the log of the tile size
+  # `lscpu` can return correct values.
+  # We underestimate modern cache sizes so that performance is good even on older architectures.
+  const cacheLine = 64     # Size of a cache line
+  const l1Size = 32 * 1024 # Size of L1 cache
+  const elems_per_cacheline = max(1, cacheLine div sizeof(T))
+
+  var q = l1Size div sizeof(T)
+  q = q div 2 # use only half of the cache, this limits cache eviction, especially with hyperthreading.
+  q = q.uint32.nextPowerOfTwo_vartime().log2_vartime().int
+  q = q div 2 # 2²𐞥 should be smaller than the cache
+
+  # If the cache line can accomodate spare elements
+  #
+  while 1 shl q < elems_per_cacheline:
+    q += 1
+
+  return
+
+func bit_reversal_permutation[N: static int, T](buf: array[N, T]) =
+  ## Bit reversal permutation using a cache-blocking algorithm
+
+# ############################################################
+#
+#                    Sanity checks
+#
+# ############################################################
+
+when isMainModule:
+
+  import
+    std/[times, monotimes, strformat],
+    ../../../helpers/prng_unsafe,
+    ../constants/zoo_generators,
+    ../io/[io_fields, io_ec]
+
+  const ctt_eth_kzg_fr_pow2_roots_of_unity = [
+    # primitive_root⁽ᵐᵒᵈᵘˡᵘˢ⁻¹⁾/⁽²^ⁱ⁾ for i in [0, 32)
+    # The primitive root chosen is 7
+    Fr[BLS12_381].fromHex"0x1",
+    Fr[BLS12_381].fromHex"0x73eda753299d7d483339d80809a1d80553bda402fffe5bfeffffffff00000000",
+    Fr[BLS12_381].fromHex"0x8d51ccce760304d0ec030002760300000001000000000000",
+    Fr[BLS12_381].fromHex"0x345766f603fa66e78c0625cd70d77ce2b38b21c28713b7007228fd3397743f7a",
+    Fr[BLS12_381].fromHex"0x20b1ce9140267af9dd1c0af834cec32c17beb312f20b6f7653ea61d87742bcce",
+    Fr[BLS12_381].fromHex"0x50e0903a157988bab4bcd40e22f55448bf6e88fb4c38fb8a360c60997369df4e",
+    Fr[BLS12_381].fromHex"0x45af6345ec055e4d14a1e27164d8fdbd2d967f4be2f951558140d032f0a9ee53",
+    Fr[BLS12_381].fromHex"0x6898111413588742b7c68b4d7fdd60d098d0caac87f5713c5130c2c1660125be",
+    Fr[BLS12_381].fromHex"0x4f9b4098e2e9f12e6b368121ac0cf4ad0a0865a899e8deff4935bd2f817f694b",
+    Fr[BLS12_381].fromHex"0x95166525526a65439feec240d80689fd697168a3a6000fe4541b8ff2ee0434e",
+    Fr[BLS12_381].fromHex"0x325db5c3debf77a18f4de02c0f776af3ea437f9626fc085e3c28d666a5c2d854",
+    Fr[BLS12_381].fromHex"0x6d031f1b5c49c83409f1ca610a08f16655ea6811be9c622d4a838b5d59cd79e5",
+    Fr[BLS12_381].fromHex"0x564c0a11a0f704f4fc3e8acfe0f8245f0ad1347b378fbf96e206da11a5d36306",
+    Fr[BLS12_381].fromHex"0x485d512737b1da3d2ccddea2972e89ed146b58bc434906ac6fdd00bfc78c8967",
+    Fr[BLS12_381].fromHex"0x56624634b500a166dc86b01c0d477fa6ae4622f6a9152435034d2ff22a5ad9e1",
+    Fr[BLS12_381].fromHex"0x3291357ee558b50d483405417a0cbe39c8d5f51db3f32699fbd047e11279bb6e",
+    Fr[BLS12_381].fromHex"0x2155379d12180caa88f39a78f1aeb57867a665ae1fcadc91d7118f85cd96b8ad",
+    Fr[BLS12_381].fromHex"0x224262332d8acbf4473a2eef772c33d6cd7f2bd6d0711b7d08692405f3b70f10",
+    Fr[BLS12_381].fromHex"0x2d3056a530794f01652f717ae1c34bb0bb97a3bf30ce40fd6f421a7d8ef674fb",
+    Fr[BLS12_381].fromHex"0x520e587a724a6955df625e80d0adef90ad8e16e84419c750194e8c62ecb38d9d",
+    Fr[BLS12_381].fromHex"0x3e1c54bcb947035a57a6e07cb98de4a2f69e02d265e09d9fece7e0e39898d4b",
+    Fr[BLS12_381].fromHex"0x47c8b5817018af4fc70d0874b0691d4e46b3105f04db5844cd3979122d3ea03a",
+    Fr[BLS12_381].fromHex"0xabe6a5e5abcaa32f2d38f10fbb8d1bbe08fec7c86389beec6e7a6ffb08e3363",
+    Fr[BLS12_381].fromHex"0x73560252aa0655b25121af06a3b51e3cc631ffb2585a72db5616c57de0ec9eae",
+    Fr[BLS12_381].fromHex"0x291cf6d68823e6876e0bcd91ee76273072cf6a8029b7d7bc92cf4deb77bd779c",
+    Fr[BLS12_381].fromHex"0x19fe632fd3287390454dc1edc61a1a3c0ba12bb3da64ca5ce32ef844e11a51e",
+    Fr[BLS12_381].fromHex"0xa0a77a3b1980c0d116168bffbedc11d02c8118402867ddc531a11a0d2d75182",
+    Fr[BLS12_381].fromHex"0x23397a9300f8f98bece8ea224f31d25db94f1101b1d7a628e2d0a7869f0319ed",
+    Fr[BLS12_381].fromHex"0x52dd465e2f09425699e276b571905a7d6558e9e3f6ac7b41d7b688830a4f2089",
+    Fr[BLS12_381].fromHex"0xc83ea7744bf1bee8da40c1ef2bb459884d37b826214abc6474650359d8e211b",
+    Fr[BLS12_381].fromHex"0x2c6d4e4511657e1e1339a815da8b398fed3a181fabb30adc694341f608c9dd56",
+    Fr[BLS12_381].fromHex"0x4b5371495990693fad1715b02e5713b5f070bb00e28a193d63e7cb4906ffc93f"
+  ]
+
+  type EC_G1 = ECP_ShortW_Prj[Fp[BLS12_381], G1]
+
+  proc roundtrip() =
+    let fftDesc = ECFFT_Descriptor[EC_G1].new(order = 1 shl 4, ctt_eth_kzg_fr_pow2_roots_of_unity[4])
+    var data = newSeq[EC_G1](fftDesc.order)
+    data[0].fromAffine(BLS12_381.getGenerator("G1"))
+    for i in 1 ..< fftDesc.order:
+      data[i].madd(data[i-1], BLS12_381.getGenerator("G1"))
+
+    var coefs = newSeq[EC_G1](data.len)
+    let fftOk = fft(fftDesc, coefs, data)
+    doAssert fftOk == FFTS_Success
+    # display("coefs", 0, coefs)
+
+    var res = newSeq[EC_G1](data.len)
+    let ifftOk = ifft(fftDesc, res, coefs)
+    doAssert ifftOk == FFTS_Success
+    # display("res", 0, res)
+
+    for i in 0 ..< res.len:
+      if bool(res[i] != data[i]):
+        echo "Error: expected ", data[i].toHex(), " but got ", res[i].toHex()
+        quit 1
+
+    echo "FFT round-trip check SUCCESS"
+
+  proc warmup() =
+    # Warmup - make sure cpu is on max perf
+    let start = cpuTime()
+    var foo = 123
+    for i in 0 ..< 300_000_000:
+      foo += i*i mod 456
+      foo = foo mod 789
+
+    # Compiler shouldn't optimize away the results as cpuTime rely on sideeffects
+    let stop = cpuTime()
+    echo &"Warmup: {stop - start:>4.4f} s, result {foo} (displayed to avoid compiler optimizing warmup away)\n"
+
+
+  proc bench() =
+    echo "Starting benchmark ..."
+    const NumIters = 3
+
+    var rng: RngState
+    rng.seed 0x1234
+    # TODO: view types complain about mutable borrow
+    # in `random_unsafe` due to pseudo view type LimbsViewMut
+    # (which was views before Nim properly supported them)
+
+    warmup()
+
+    for scale in 4 ..< 10:
+      # Setup
+
+      let fftDesc = ECFFTDescriptor[EC_G1].new(order = 1 shl scale, ctt_eth_kzg_fr_pow2_roots_of_unity[scale])
+      var data = newSeq[EC_G1](fftDesc.order)
+      data[0].fromAffine(BLS12_381.getGenerator("G1"))
+      for i in 1 ..< fftDesc.order:
+        data[i].madd(data[i-1], BLS12_381.getGenerator("G1"))
+
+      var coefsOut = newSeq[EC_G1](data.len)
+
+      # Bench
+      let start = getMonotime()
+      for i in 0 ..< NumIters:
+        let status = fftDesc.fft(coefsOut, data)
+        doAssert status == FFTS_Success
+      let stop = getMonotime()
+
+      let ns = inNanoseconds((stop-start) div NumIters)
+      echo &"FFT scale {scale:>2}     {ns:>8} ns/op"
+
+  roundtrip()
+  warmup()
+  bench()
diff --git a/constantine/platforms/bithacks.nim b/constantine/platforms/bithacks.nim
index 7fcdca98..20a87861 100644
--- a/constantine/platforms/bithacks.nim
+++ b/constantine/platforms/bithacks.nim
@@ -156,4 +156,82 @@ func isPowerOf2_vartime*(n: SomeUnsignedInt): bool {.inline.} =
 func nextPowerOfTwo_vartime*(n: uint32): uint32 {.inline.} =
   ## Returns x if x is a power of 2
   ## or the next biggest power of 2
-  1'u32 shl (log2_vartime(n-1) + 1)
\ No newline at end of file
+  1'u32 shl (log2_vartime(n-1) + 1)
+
+func swapBytes_impl(n: uint32): uint32 {.inline.} =
+  result = n
+  result = ((result shl 8) and 0xff00ff00'u32) or ((result shr 8) and 0x00ff00ff'u32)
+  result = (result shl 16) or (result shr 16)
+
+func swapBytes_impl(n: uint64): uint64 {.inline.} =
+  result = n
+  result = ((result shl 8) and 0xff00ff00ff00ff00'u64) or ((result shr 8) and 0x00ff00ff00ff00ff'u64)
+  result = ((result shl 16) and 0xffff0000ffff0000'u64) or ((result shr 16) and 0x0000ffff0000ffff'u64)
+  result = (result shl 32) or (result shr 32)
+
+func swapBytes*(n: SomeUnsignedInt): SomeUnsignedInt {.inline.} =
+  # Note:
+  #   using the raw Nim implementation:
+  #     - leads to vectorized code if swapping an array
+  #     - leads to builtin swap on modern compilers
+  when nimvm:
+    swapBytes_impl(n)
+  else:
+    swapBytes_c_compiler(n)
+
+func reverseBits*(n: uint32): uint32 {.inline.} =
+  when true:
+    var n = n
+    n = ((n and 0x55555555'u32) shl 1) or ((n and 0xaaaaaaaa'u32) shr 1)
+    n = ((n and 0x33333333'u32) shl 2) or ((n and 0xcccccccc'u32) shr 2)
+    n = ((n and 0x0f0f0f0f'u32) shl 4) or ((n and 0xf0f0f0f0'u32) shr 4)
+
+    # Swap bytes - allow vectorization by using raw Nim impl instead of compiler builtin
+    n = swapBytes_impl(n)
+  else:
+    # Modern compiler should either:
+    # - either hardcode / constant propagate the bitmask m
+    # - or use fused shift+add (with LEA on x86 or add on ARM), if immediates are too costly
+    #
+    # and also vectorize the code if used for arrays
+
+    # Swap bytes - allow vectorization by using raw Nim impl instead of compiler builtin
+    var n = (n shl 16) or (n shr 16)
+    var m = 0x00ff00ff'u32
+    n = ((n shr 8) and m) or ((n shl 8) and not m)
+
+    # Reverse bits
+    m = m xor (m shl 4); n = ((n shr 4) and m) or ((n shl 4) and not m)
+    m = m xor (m shl 2); n = ((n shr 2) and m) or ((n shl 2) and not m)
+    m = m xor (m shl 1); n = ((n shr 1) and m) or ((n shl 1) and not m)
+
+  return n
+
+func reverseBits*(n: uint64): uint64 {.inline.} =
+  when true:
+    var n = n
+    n = ((n and 0x5555555555555555'u64) shl 1) or ((n and 0xaaaaaaaaaaaaaaaa'u64) shr 1)
+    n = ((n and 0x3333333333333333'u64) shl 2) or ((n and 0xcccccccccccccccc'u64) shr 2)
+    n = ((n and 0x0f0f0f0f0f0f0f0f'u64) shl 4) or ((n and 0xf0f0f0f0f0f0f0f0'u64) shr 4)
+
+    # Swap bytes - allow vectorization by using raw Nim impl instead of compiler builtin
+    n = swapBytes_impl(n)
+  else:
+    # Modern compiler should either:
+    # - either hardcode / constant propagate the bitmask m
+    # - or use fused shift+add (with LEA on x86 or add on ARM), if immediates are too costly
+    #
+    # and also vectorize the code if used for arrays
+
+    # Swap bytes - allow vectorization by using raw Nim impl instead of compiler builtin
+    var n = (n shl 32) or (n shr 32)
+    var m = 0x0000ffff0000ffff'u64
+    n = ((n shr 16) and m) or ((n shl 16) and not m)
+    m = m xor (m shl 8); n = ((n shr 8) and m) or ((n shl 8) and not m)
+
+    # Reverse bits
+    m = m xor (m shl 4); n = ((n shr 4) and m) or ((n shl 4) and not m)
+    m = m xor (m shl 2); n = ((n shr 2) and m) or ((n shl 2) and not m)
+    m = m xor (m shl 1); n = ((n shr 1) and m) or ((n shl 1) and not m)
+
+    return n
\ No newline at end of file
diff --git a/constantine/platforms/intrinsics/bitops.nim b/constantine/platforms/intrinsics/bitops.nim
index 218a8fad..00a7796e 100644
--- a/constantine/platforms/intrinsics/bitops.nim
+++ b/constantine/platforms/intrinsics/bitops.nim
@@ -49,6 +49,12 @@ when GCC_Compatible:
       else:
         builtin_ctz(n.uint32)
 
+  func builtin_swapBytes(n: uint32): uint32 {.importc: "__builtin_bswap32", nodecl.}
+  func builtin_swapBytes(n: uint64): uint64 {.importc: "__builtin_bswap64", nodecl.}
+
+  func swapBytes_c_compiler*(n: SomeUnsignedInt): SomeUnsignedInt {.inline.} =
+    builtin_swapBytes(n)
+
 elif defined(icc):
   func bitScanReverse(r: var uint32, n: uint32): uint8 {.importc: "_BitScanReverse", header: "<immintrin.h>".}
     ## Returns 0 if n is zero and non-zero otherwise
@@ -93,6 +99,12 @@ elif defined(icc):
     else:
       bitscan(bitScanForward, c.uint32, default = 0)
 
+  func builtin_swapBytes(n: uint32): uint32 {.importc: "_bswap", nodecl.}
+  func builtin_swapBytes(n: uint64): uint64 {.importc: "_bswap64", nodecl.}
+
+  func swapBytes_c_compiler*(n: SomeUnsignedInt): SomeUnsignedInt {.inline.} =
+    builtin_swapBytes(n)
+
 elif defined(vcc):
   func bitScanReverse(p: ptr uint32, b: uint32): uint8 {.importc: "_BitScanReverse", header: "<intrin.h>".}
     ## Returns 0 if n s no set bit and non-zero otherwise
@@ -137,5 +149,11 @@ elif defined(vcc):
     else:
       bitscan(bitScanForward, c.uint32, default = sizeof(n) * 8)
 
+  func builtin_swapBytes(n: uint32): uint32 {.importc: "_byteswap_ulong", cdecl, header: "<intrin.h>".}
+  func builtin_swapBytes(n: uint64): uint64 {.importc: "_byteswap_uint64", cdecl, header: "<intrin.h>".}
+
+  func swapBytes_c_compiler*(n: SomeUnsignedInt): SomeUnsignedInt {.inline.} =
+    builtin_swapBytes(n)
+
 else:
   {. error: "Unsupported compiler".}
\ No newline at end of file
diff --git a/constantine/platforms/isa/macro_assembler_x86.nim b/constantine/platforms/isa/macro_assembler_x86.nim
index 47b4fb79..85d2bbb9 100644
--- a/constantine/platforms/isa/macro_assembler_x86.nim
+++ b/constantine/platforms/isa/macro_assembler_x86.nim
@@ -147,10 +147,7 @@ func toString*(nimSymbol: NimNode): string =
 
 func hash(od: OperandDesc): Hash =
   {.noSideEffect.}:
-    try: # Why does this raise a generic exception?
-      hash(od.nimSymbol.toString())
-    except:
-      raise newException(Defect, "Broke Nim")
+    hash(od.nimSymbol.toString())
 
 func len*(opArray: OperandArray): int =
   opArray.buf.len
@@ -212,10 +209,7 @@ func genMemClobber(nimSymbol: NimNode, len: int, memIndirect: MemIndirectAccess)
 
 func asmValue*(nimSymbol: NimNode, rm: RM, constraint: Constraint): Operand =
   {.noSideEffect.}:
-    let symStr = try: # Why does this raise a generic exception?
-      $nimSymbol
-    except:
-      raise newException(Defect, "Broke Nim!")
+    let symStr = $nimSymbol
 
   let desc = OperandDesc(
         asmId: "[" & symStr & "]",
@@ -362,10 +356,7 @@ func setToCarryFlag*(a: var Assembler_x86, carry: NimNode) =
   let nimSymbol = if isHiddenDeref: carry[0]
                   else: carry
   {.noSideEffect.}:
-    let symStr = try: # Why does this raise a generic exception?
-      $nimSymbol
-    except:
-      raise newException(Defect, "Broke Nim!")
+    let symStr = $nimSymbol
 
   let desc = OperandDesc(
     asmId: "",
diff --git a/constantine/platforms/views.nim b/constantine/platforms/views.nim
index 0e47928c..e4887520 100644
--- a/constantine/platforms/views.nim
+++ b/constantine/platforms/views.nim
@@ -47,6 +47,173 @@ func `[]`*[T](v: MutableView[T], idx: int): var T {.inline.} =
 func `[]=`*[T](v: MutableView[T], idx: int, val: T) {.inline.} =
   v.data[idx] = val
 
+# StridedView type
+# ---------------------------------------------------------
+# using the borrow checker with `lent` requires a recent Nim
+# https://github.com/nim-lang/Nim/issues/21674
+
+type
+  StridedView*[T] = object
+    ## A strided view over an (unowned) data buffer
+    len*: int
+    stride: int
+    offset: int
+    data: ptr UncheckedArray[T]
+
+func `[]`*[T](v: StridedView[T], idx: int): lent T {.inline.} =
+  v.data[v.offset + idx*v.stride]
+
+func `[]`*[T](v: var StridedView[T], idx: int): var T {.inline.} =
+  v.data[v.offset + idx*v.stride]
+
+func `[]=`*[T](v: var StridedView[T], idx: int, val: T) {.inline.} =
+  v.data[v.offset + idx*v.stride] = val
+
+func toStridedView*[T](oa: openArray[T]): StridedView[T] {.inline.} =
+  result.len = oa.len
+  result.stride = 1
+  result.offset = 0
+  result.data = cast[ptr UncheckedArray[T]](oa[0].unsafeAddr)
+
+func toStridedView*[T](p: ptr UncheckedArray[T], len: int): StridedView[T] {.inline.} =
+  result.len = len
+  result.stride = 1
+  result.offset = 0
+  result.data = p
+
+iterator items*[T](v: StridedView[T]): lent T =
+  var cur = v.offset
+  for _ in 0 ..< v.len:
+    yield v.data[cur]
+    cur += v.stride
+
+func `$`*(v: StridedView): string =
+  result = "StridedView["
+  var first = true
+  for elem in v:
+    if not first:
+      result &= ", "
+    else:
+      first = false
+    result &= $elem
+  result &= ']'
+
+func toHex*(v: StridedView): string =
+  mixin toHex
+
+  result = "StridedView["
+  var first = true
+  for elem in v:
+    if not first:
+      result &= ", "
+    else:
+      first = false
+    result &= elem.toHex()
+  result &= ']'
+
+# FFT-specific splitting
+# -------------------------------------------------------------------------------
+
+func splitAlternate*(t: StridedView): tuple[even, odd: StridedView] {.inline.} =
+  ## Split the tensor into 2
+  ## partitioning the input every other index
+  ## even: indices [0, 2, 4, ...]
+  ## odd: indices [ 1, 3, 5, ...]
+  assert (t.len and 1) == 0, "The tensor must contain an even number of elements"
+
+  let half = t.len shr 1
+  let skipHalf = t.stride shl 1
+
+  result.even.len = half
+  result.even.stride = skipHalf
+  result.even.offset = t.offset
+  result.even.data = t.data
+
+  result.odd.len = half
+  result.odd.stride = skipHalf
+  result.odd.offset = t.offset + t.stride
+  result.odd.data = t.data
+
+func splitMiddle*(t: StridedView): tuple[left, right: StridedView] {.inline.} =
+  ## Split the tensor into 2
+  ## partitioning into left and right halves.
+  ## left:  indices [0, 1, 2, 3]
+  ## right: indices  [4, 5, 6, 7]
+  assert (t.len and 1) == 0, "The tensor must contain an even number of elements"
+
+  let half = t.len shr 1
+
+  result.left.len = half
+  result.left.stride = t.stride
+  result.left.offset = t.offset
+  result.left.data = t.data
+
+  result.right.len = half
+  result.right.stride = t.stride
+  result.right.offset = t.offset + half
+  result.right.data = t.data
+
+func skipHalf*(t: StridedView): StridedView {.inline.} =
+  ## Pick one every other indices
+  ## output: [0, 2, 4, ...]
+  assert (t.len and 1) == 0, "The tensor must contain an even number of elements"
+
+  result.len = t.len shr 1
+  result.stride = t.stride shl 1
+  result.offset = t.offset
+  result.data = t.data
+
+func slice*(v: StridedView, start, stop, step: int): StridedView {.inline.} =
+  ## Slice a view
+  ## stop is inclusive
+  # General tensor slicing algorithm is
+  # https://github.com/mratsim/Arraymancer/blob/71cf616/src/arraymancer/tensor/private/p_accessors_macros_read.nim#L26-L56
+  #
+  # for i, slice in slices:
+  #   # Check if we start from the end
+  #   let a = if slice.a_from_end: result.shape[i] - slice.a
+  #           else: slice.a
+  #
+  #   let b = if slice.b_from_end: result.shape[i] - slice.b
+  #           else: slice.b
+  #
+  #   # Compute offset:
+  #   result.offset += a * result.strides[i]
+  #   # Now change shape and strides
+  #   result.strides[i] *= slice.step
+  #   result.shape[i] = abs((b-a) div slice.step) + 1
+  #
+  # with slices being of size 1, as we have a monodimensional Tensor
+  # and the slice being a..<b with the reverse case: len-1 -> 0
+  #
+  # result is preinitialized with a copy of v (shape, stride, offset, data)
+  result.offset = v.offset + start * v.stride
+  result.stride = v.stride * step
+  result.len = abs((stop-start) div step) + 1
+  result.data = v.data
+
+func reversed*(v: StridedView): StridedView {.inline.} =
+  # Hopefully the compiler optimizes div by -1
+  v.slice(v.len-1, 0, -1)
+
+# Debugging helpers
+# ---------------------------------------------------------
+
+when defined(debugConstantine):
+  import std/[strformat, strutils]
+
+  func display*[F](name: string, indent: int, oa: openArray[F]) =
+    debugEcho strutils.indent(name & ", openarray of " & $F & " of length " & $oa.len, indent)
+    for i in 0 ..< oa.len:
+      debugEcho strutils.indent(&"    {i:>2}: {oa[i].toHex()}", indent)
+    debugEcho strutils.indent(name & "  " & $F & " -- FIN\n", indent)
+
+  func display*[F](name: string, indent: int, v: StridedView[F]) =
+    debugEcho strutils.indent(name & ", view of " & $F & " of length " & $v.len, indent)
+    for i in 0 ..< v.len:
+      debugEcho strutils.indent(&"    {i:>2}: {v[i].toHex()}", indent)
+    debugEcho strutils.indent(name & "  " & $F & " -- FIN\n", indent)
+
 # Binary blob API
 # ---------------------------------------------------------
 #
diff --git a/constantine/trusted_setups/gen_eth_kzg_testing_setups.nim b/constantine/trusted_setups/gen_eth_kzg_testing_setups.nim
index 7e90efa6..466731bf 100644
--- a/constantine/trusted_setups/gen_eth_kzg_testing_setups.nim
+++ b/constantine/trusted_setups/gen_eth_kzg_testing_setups.nim
@@ -6,6 +6,16 @@
 #   * Apache v2 license (license terms in the root directory or at http://www.apache.org/licenses/LICENSE-2.0).
 # at your option. This file may not be copied, modified, or distributed except according to those terms.
 
+import
+  ../math/polynomials/fft,
+  ../math/[arithmetic, extension_fields, ec_shortweierstrass],
+  ../math/config/curves,
+  ../math/elliptic/[ec_scalar_mul_vartime, ec_shortweierstrass_batch_ops],
+  ../math/io/io_fields,
+  ../math/constants/zoo_generators,
+  ../platforms/abstractions,
+  std/streams
+
 # This tool generates the same testing setups that are used in Ethereum consensus-spec
 # in a Constantine-specific format specified in README.md
 
@@ -24,7 +34,6 @@
 #   python3 ./gen_kzg_trusted_setups.py --secret=1337 --g1-length=4096 --g2-length=65
 #   https://github.com/ethereum/consensus-specs/blob/v1.3.0/Makefile#L209-L210
 
-
 # Roots of unity
 # ------------------------------------------------------------
 #
@@ -56,39 +65,95 @@
 #
 #   BLS12-381 was chosen for its high 2-adicity, as 2^32 is a factor of its order-1
 
-const ctt_eth_kzg_fr_pow2_roots_of_unity = (
+const ctt_eth_kzg_fr_pow2_roots_of_unity = [
   # primitive_root⁽ᵐᵒᵈᵘˡᵘˢ⁻¹⁾/⁽²^ⁱ⁾ for i in [0, 32)
   # The primitive root chosen is 7
-  BigInt[1].fromHex"0x1",
-  BigInt[255].fromHex"0x73eda753299d7d483339d80809a1d80553bda402fffe5bfeffffffff00000000",
-  BigInt[192].fromHex"0x8d51ccce760304d0ec030002760300000001000000000000",
-  BigInt[254].fromHex"0x345766f603fa66e78c0625cd70d77ce2b38b21c28713b7007228fd3397743f7a",
-  BigInt[254].fromHex"0x20b1ce9140267af9dd1c0af834cec32c17beb312f20b6f7653ea61d87742bcce",
-  BigInt[255].fromHex"0x50e0903a157988bab4bcd40e22f55448bf6e88fb4c38fb8a360c60997369df4e",
-  BigInt[255].fromHex"0x45af6345ec055e4d14a1e27164d8fdbd2d967f4be2f951558140d032f0a9ee53",
-  BigInt[255].fromHex"0x6898111413588742b7c68b4d7fdd60d098d0caac87f5713c5130c2c1660125be",
-  BigInt[255].fromHex"0x4f9b4098e2e9f12e6b368121ac0cf4ad0a0865a899e8deff4935bd2f817f694b",
-  BigInt[252].fromHex"0x95166525526a65439feec240d80689fd697168a3a6000fe4541b8ff2ee0434e",
-  BigInt[254].fromHex"0x325db5c3debf77a18f4de02c0f776af3ea437f9626fc085e3c28d666a5c2d854",
-  BigInt[255].fromHex"0x6d031f1b5c49c83409f1ca610a08f16655ea6811be9c622d4a838b5d59cd79e5",
-  BigInt[255].fromHex"0x564c0a11a0f704f4fc3e8acfe0f8245f0ad1347b378fbf96e206da11a5d36306",
-  BigInt[255].fromHex"0x485d512737b1da3d2ccddea2972e89ed146b58bc434906ac6fdd00bfc78c8967",
-  BigInt[255].fromHex"0x56624634b500a166dc86b01c0d477fa6ae4622f6a9152435034d2ff22a5ad9e1",
-  BigInt[254].fromHex"0x3291357ee558b50d483405417a0cbe39c8d5f51db3f32699fbd047e11279bb6e",
-  BigInt[254].fromHex"0x2155379d12180caa88f39a78f1aeb57867a665ae1fcadc91d7118f85cd96b8ad",
-  BigInt[254].fromHex"0x224262332d8acbf4473a2eef772c33d6cd7f2bd6d0711b7d08692405f3b70f10",
-  BigInt[254].fromHex"0x2d3056a530794f01652f717ae1c34bb0bb97a3bf30ce40fd6f421a7d8ef674fb",
-  BigInt[255].fromHex"0x520e587a724a6955df625e80d0adef90ad8e16e84419c750194e8c62ecb38d9d",
-  BigInt[250].fromHex"0x3e1c54bcb947035a57a6e07cb98de4a2f69e02d265e09d9fece7e0e39898d4b",
-  BigInt[255].fromHex"0x47c8b5817018af4fc70d0874b0691d4e46b3105f04db5844cd3979122d3ea03a",
-  BigInt[252].fromHex"0xabe6a5e5abcaa32f2d38f10fbb8d1bbe08fec7c86389beec6e7a6ffb08e3363",
-  BigInt[255].fromHex"0x73560252aa0655b25121af06a3b51e3cc631ffb2585a72db5616c57de0ec9eae",
-  BigInt[254].fromHex"0x291cf6d68823e6876e0bcd91ee76273072cf6a8029b7d7bc92cf4deb77bd779c",
-  BigInt[249].fromHex"0x19fe632fd3287390454dc1edc61a1a3c0ba12bb3da64ca5ce32ef844e11a51e",
-  BigInt[252].fromHex"0xa0a77a3b1980c0d116168bffbedc11d02c8118402867ddc531a11a0d2d75182",
-  BigInt[254].fromHex"0x23397a9300f8f98bece8ea224f31d25db94f1101b1d7a628e2d0a7869f0319ed",
-  BigInt[255].fromHex"0x52dd465e2f09425699e276b571905a7d6558e9e3f6ac7b41d7b688830a4f2089",
-  BigInt[252].fromHex"0xc83ea7744bf1bee8da40c1ef2bb459884d37b826214abc6474650359d8e211b",
-  BigInt[254].fromHex"0x2c6d4e4511657e1e1339a815da8b398fed3a181fabb30adc694341f608c9dd56",
-  BigInt[255].fromHex"0x4b5371495990693fad1715b02e5713b5f070bb00e28a193d63e7cb4906ffc93f"
-)
\ No newline at end of file
+  Fr[BLS12_381].fromHex"0x1",
+  Fr[BLS12_381].fromHex"0x73eda753299d7d483339d80809a1d80553bda402fffe5bfeffffffff00000000",
+  Fr[BLS12_381].fromHex"0x8d51ccce760304d0ec030002760300000001000000000000",
+  Fr[BLS12_381].fromHex"0x345766f603fa66e78c0625cd70d77ce2b38b21c28713b7007228fd3397743f7a",
+  Fr[BLS12_381].fromHex"0x20b1ce9140267af9dd1c0af834cec32c17beb312f20b6f7653ea61d87742bcce",
+  Fr[BLS12_381].fromHex"0x50e0903a157988bab4bcd40e22f55448bf6e88fb4c38fb8a360c60997369df4e",
+  Fr[BLS12_381].fromHex"0x45af6345ec055e4d14a1e27164d8fdbd2d967f4be2f951558140d032f0a9ee53",
+  Fr[BLS12_381].fromHex"0x6898111413588742b7c68b4d7fdd60d098d0caac87f5713c5130c2c1660125be",
+  Fr[BLS12_381].fromHex"0x4f9b4098e2e9f12e6b368121ac0cf4ad0a0865a899e8deff4935bd2f817f694b",
+  Fr[BLS12_381].fromHex"0x95166525526a65439feec240d80689fd697168a3a6000fe4541b8ff2ee0434e",
+  Fr[BLS12_381].fromHex"0x325db5c3debf77a18f4de02c0f776af3ea437f9626fc085e3c28d666a5c2d854",
+  Fr[BLS12_381].fromHex"0x6d031f1b5c49c83409f1ca610a08f16655ea6811be9c622d4a838b5d59cd79e5",
+  Fr[BLS12_381].fromHex"0x564c0a11a0f704f4fc3e8acfe0f8245f0ad1347b378fbf96e206da11a5d36306",
+  Fr[BLS12_381].fromHex"0x485d512737b1da3d2ccddea2972e89ed146b58bc434906ac6fdd00bfc78c8967",
+  Fr[BLS12_381].fromHex"0x56624634b500a166dc86b01c0d477fa6ae4622f6a9152435034d2ff22a5ad9e1",
+  Fr[BLS12_381].fromHex"0x3291357ee558b50d483405417a0cbe39c8d5f51db3f32699fbd047e11279bb6e",
+  Fr[BLS12_381].fromHex"0x2155379d12180caa88f39a78f1aeb57867a665ae1fcadc91d7118f85cd96b8ad",
+  Fr[BLS12_381].fromHex"0x224262332d8acbf4473a2eef772c33d6cd7f2bd6d0711b7d08692405f3b70f10",
+  Fr[BLS12_381].fromHex"0x2d3056a530794f01652f717ae1c34bb0bb97a3bf30ce40fd6f421a7d8ef674fb",
+  Fr[BLS12_381].fromHex"0x520e587a724a6955df625e80d0adef90ad8e16e84419c750194e8c62ecb38d9d",
+  Fr[BLS12_381].fromHex"0x3e1c54bcb947035a57a6e07cb98de4a2f69e02d265e09d9fece7e0e39898d4b",
+  Fr[BLS12_381].fromHex"0x47c8b5817018af4fc70d0874b0691d4e46b3105f04db5844cd3979122d3ea03a",
+  Fr[BLS12_381].fromHex"0xabe6a5e5abcaa32f2d38f10fbb8d1bbe08fec7c86389beec6e7a6ffb08e3363",
+  Fr[BLS12_381].fromHex"0x73560252aa0655b25121af06a3b51e3cc631ffb2585a72db5616c57de0ec9eae",
+  Fr[BLS12_381].fromHex"0x291cf6d68823e6876e0bcd91ee76273072cf6a8029b7d7bc92cf4deb77bd779c",
+  Fr[BLS12_381].fromHex"0x19fe632fd3287390454dc1edc61a1a3c0ba12bb3da64ca5ce32ef844e11a51e",
+  Fr[BLS12_381].fromHex"0xa0a77a3b1980c0d116168bffbedc11d02c8118402867ddc531a11a0d2d75182",
+  Fr[BLS12_381].fromHex"0x23397a9300f8f98bece8ea224f31d25db94f1101b1d7a628e2d0a7869f0319ed",
+  Fr[BLS12_381].fromHex"0x52dd465e2f09425699e276b571905a7d6558e9e3f6ac7b41d7b688830a4f2089",
+  Fr[BLS12_381].fromHex"0xc83ea7744bf1bee8da40c1ef2bb459884d37b826214abc6474650359d8e211b",
+  Fr[BLS12_381].fromHex"0x2c6d4e4511657e1e1339a815da8b398fed3a181fabb30adc694341f608c9dd56",
+  Fr[BLS12_381].fromHex"0x4b5371495990693fad1715b02e5713b5f070bb00e28a193d63e7cb4906ffc93f"
+]
+
+func newTrustedSetup(
+       EC: typedesc, secret: auto, length: int): seq[EC] =
+
+  var ts = newSeq[jacobian(EC)](length)
+  result.newSeq(length)
+
+  var P {.noInit.}: jacobian(EC)
+  P.fromAffine(EC.F.C.getGenerator($EC.G))
+  ts[0] = P
+  for i in 1 ..< length:
+    P.scalarMul_minHammingWeight_windowed_vartime(secret, window = 5)
+    ts[i] = P
+
+  batchAffine(result.asUnchecked(), ts.asUnchecked(), length)
+
+proc genEthereumKzgTestingTrustedSetup(filepath: string, secret: auto, length: int) =
+  ## Generate an Ethereum KZG testing trusted setup
+  ## in the Trusted Setup Interchange Format
+
+  let f = openFileStream(filepath, fmWrite)
+  f.write"\u2203\u222A\u2208\u220E" # ∃⋃∈∎ in UTF-16. (magic bytes)
+
+  # v1.0
+  f.write 'v'
+  f.write uint8 1
+  f.write '.'
+  f.write uint8 0
+
+  # Protocol
+  f.write"ethereum_deneb_kzg"
+
+  # Curve
+  const curve = "bls12_381"
+  f.write curve
+  const padCurve = array[15 - curve.len, byte] # zero-nit padding
+  f.write padCurve
+
+  # Number of fields
+  f.write uint8 3
+
+
+
+when isMainModule:
+  import ../math/io/[io_bigints, io_ec]
+
+  let secret = BigInt[11].fromUint(1337'u64)
+  let ts1 = newTrustedSetup(ECP_ShortW_Aff[Fp[BLS12_381], G1], secret, 4)
+
+  for i in 0 ..< ts1.len:
+    echo "ts1[", i, "]: ", ts1[i].toHex()
+
+  let ts2 = newTrustedSetup(ECP_ShortW_Aff[Fp2[BLS12_381], G2], secret, 65)
+
+  for i in 0 ..< ts2.len:
+    echo "ts2[", i, "]: ", ts2[i].toHex()
\ No newline at end of file
diff --git a/research/kzg/fft_fr.nim b/research/kzg/fft_fr.nim
index b6a11b96..4a77dab6 100644
--- a/research/kzg/fft_fr.nim
+++ b/research/kzg/fft_fr.nim
@@ -192,8 +192,6 @@ proc init*(T: type FFTDescriptor, maxScale: uint8): T =
 #
 # ############################################################
 
-{.experimental: "views".}
-
 when isMainModule:
   import
     std/[times, monotimes, strformat],
diff --git a/research/kzg/fft_g1.nim b/research/kzg/fft_g1.nim
index 70ffce57..4a4d041f 100644
--- a/research/kzg/fft_g1.nim
+++ b/research/kzg/fft_g1.nim
@@ -11,7 +11,7 @@ import
   ../../constantine/math/config/curves,
   ../../constantine/math/arithmetic,
   ../../constantine/math/ec_shortweierstrass,
-  ../../constantine/math/io/[io_fields, io_ec],
+  ../../constantine/math/io/[io_fields, io_ec, io_bigints],
   # Research
   ./strided_views,
   ./fft_lut
@@ -31,12 +31,12 @@ import
 # - https://github.com/zkcrypto/bellman/blob/10c5010/src/domain.rs#L272-L315
 # - Modern Computer Arithmetic, Brent and Zimmermann, p53 algorithm 2.2
 #   https://members.loria.fr/PZimmermann/mca/mca-cup-0.5.9.pdf
-
 # ############################################################
 #
 #            Finite-Field Fast Fourier Transform
 #
 # ############################################################
+
 #
 # This is a research, unoptimized implementation of
 # Finite Field Fast Fourier Transform
@@ -176,7 +176,7 @@ func ifft*[EC](
 
   var invLen {.noInit.}: Fr[EC.F.C]
   invLen.fromUint(vals.len.uint64)
-  invLen.inv()
+  invLen.inv_vartime()
   let inv = invLen.toBig()
 
   for i in 0..< output.len:
@@ -201,8 +201,6 @@ proc init*(T: type FFTDescriptor, maxScale: uint8): T =
 #
 # ############################################################
 
-{.experimental: "views".}
-
 when isMainModule:
   import
     std/[times, monotimes, strformat],
@@ -230,7 +228,7 @@ when isMainModule:
     var res = newSeq[EC_G1](data.len)
     let ifftOk = ifft(fftDesc, res, coefs)
     doAssert ifftOk == FFTS_Success
-    # display("res", 0, coefs)
+    # display("res", 0, res)
 
     for i in 0 ..< res.len:
       if bool(res[i] != data[i]):
diff --git a/research/kzg/fft_lut.nim b/research/kzg/fft_lut.nim
index f3af7530..6120d972 100644
--- a/research/kzg/fft_lut.nim
+++ b/research/kzg/fft_lut.nim
@@ -17,7 +17,7 @@ import
 # we can precompute everything in Sage
 # and auto-generate the file.
 
-const BLS12_381_Fr_primitive_root = 5
+const BLS12_381_Fr_primitive_root = 7
 
 func buildRootLUT(F: type Fr): array[32, F] =
   ## [pow(PRIMITIVE_ROOT, (MODULUS - 1) // (2**i), MODULUS) for i in range(32)]
diff --git a/research/kzg/strided_views.nim b/research/kzg/strided_views.nim
index 979cd197..20794e5b 100644
--- a/research/kzg/strided_views.nim
+++ b/research/kzg/strided_views.nim
@@ -18,15 +18,13 @@
 # Or the minimal tensor implementation challenge:
 # https://github.com/SimonDanisch/julia-challenge/blob/b8ed3b6/nim/nim_sol_mratsim.nim#L4-L26
 
-{.experimental: "views".}
-
 type
   View*[T] = object
     ## A strided view over an (unowned) data buffer
     len*: int
     stride: int
     offset: int
-    data: lent UncheckedArray[T]
+    data: ptr UncheckedArray[T]
 
 func `[]`*[T](v: View[T], idx: int): lent T {.inline.} =
   v.data[v.offset + idx*v.stride]
@@ -43,7 +41,7 @@ func toView*[T](oa: openArray[T]): View[T] {.inline.} =
   result.len = oa.len
   result.stride = 1
   result.offset = 0
-  result.data = cast[lent UncheckedArray[T]](oa[0].unsafeAddr)
+  result.data = cast[ptr UncheckedArray[T]](oa[0].unsafeAddr)
 
 iterator items*[T](v: View[T]): lent T =
   var cur = v.offset
@@ -168,16 +166,16 @@ func reversed*(v: View): View {.inline.} =
 import strformat, strutils
 
 func display*[F](name: string, indent: int, oa: openArray[F]) =
-  debugEcho indent(name & ", openarray of " & $F & " of length " & $oa.len, indent)
+  debugEcho strutils.indent(name & ", openarray of " & $F & " of length " & $oa.len, indent)
   for i in 0 ..< oa.len:
-    debugEcho indent(&"    {i:>2}: {oa[i].toHex()}", indent)
-  debugEcho indent(name & "  " & $F & " -- FIN\n", indent)
+    debugEcho strutils.indent(&"    {i:>2}: {oa[i].toHex()}", indent)
+  debugEcho strutils.indent(name & "  " & $F & " -- FIN\n", indent)
 
 func display*[F](name: string, indent: int, v: View[F]) =
-  debugEcho indent(name & ", view of " & $F & " of length " & $v.len, indent)
+  debugEcho strutils.indent(name & ", view of " & $F & " of length " & $v.len, indent)
   for i in 0 ..< v.len:
-    debugEcho indent(&"    {i:>2}: {v[i].toHex()}", indent)
-  debugEcho indent(name & "  " & $F & " -- FIN\n", indent)
+    debugEcho strutils.indent(&"    {i:>2}: {v[i].toHex()}", indent)
+  debugEcho strutils.indent(name & "  " & $F & " -- FIN\n", indent)
 
 # ############################################################
 #