complete correctness for idcct

ingonyama-zk · Sep 26, 2024 · ee59011 · ee59011
1 parent 2fc4614
commit ee59011
Show file tree

Hide file tree

Showing 2 changed files with 48 additions and 48 deletions.
diff --git a/icicle/src/ntt/kernel_ntt.cu b/icicle/src/ntt/kernel_ntt.cu
@@ -524,19 +524,19 @@ namespace mxntt {
     }
 
     // if (s_meta.ntt_block_id > 14 && s_meta.ntt_block_id < 18 && s_meta.ntt_inp_id == 0)
-      printf(
-        "T BEFORE: B: %d, I: %d\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n",
-        s_meta.ntt_block_id,
-        s_meta.ntt_inp_id,
-        engine.X[0].limbs_storage.limbs[0],
-        engine.X[1].limbs_storage.limbs[0],
-        engine.X[2].limbs_storage.limbs[0],
-        engine.X[3].limbs_storage.limbs[0],
-        engine.X[4].limbs_storage.limbs[0],
-        engine.X[5].limbs_storage.limbs[0],
-        engine.X[6].limbs_storage.limbs[0],
-        engine.X[7].limbs_storage.limbs[0]
-      );
+      // printf(
+      //   "T BEFORE: B: %d, I: %d\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n",
+      //   s_meta.ntt_block_id,
+      //   s_meta.ntt_inp_id,
+      //   engine.X[0].limbs_storage.limbs[0],
+      //   engine.X[1].limbs_storage.limbs[0],
+      //   engine.X[2].limbs_storage.limbs[0],
+      //   engine.X[3].limbs_storage.limbs[0],
+      //   engine.X[4].limbs_storage.limbs[0],
+      //   engine.X[5].limbs_storage.limbs[0],
+      //   engine.X[6].limbs_storage.limbs[0],
+      //   engine.X[7].limbs_storage.limbs[0]
+      // );
 
     if (dit) {
       if (inv) {
@@ -563,19 +563,19 @@ namespace mxntt {
     }
 
     // if (s_meta.ntt_block_id > 14 && s_meta.ntt_block_id < 18 && s_meta.ntt_inp_id == 0)
-      printf(
-        "T AFTER: B: %d, I: %d\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n",
-        s_meta.ntt_block_id,
-        s_meta.ntt_inp_id,
-        engine.X[0].limbs_storage.limbs[0],
-        engine.X[1].limbs_storage.limbs[0],
-        engine.X[2].limbs_storage.limbs[0],
-        engine.X[3].limbs_storage.limbs[0],
-        engine.X[4].limbs_storage.limbs[0],
-        engine.X[5].limbs_storage.limbs[0],
-        engine.X[6].limbs_storage.limbs[0],
-        engine.X[7].limbs_storage.limbs[0]
-      );
+      // printf(
+      //   "T AFTER: B: %d, I: %d\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n0x%x\n",
+      //   s_meta.ntt_block_id,
+      //   s_meta.ntt_inp_id,
+      //   engine.X[0].limbs_storage.limbs[0],
+      //   engine.X[1].limbs_storage.limbs[0],
+      //   engine.X[2].limbs_storage.limbs[0],
+      //   engine.X[3].limbs_storage.limbs[0],
+      //   engine.X[4].limbs_storage.limbs[0],
+      //   engine.X[5].limbs_storage.limbs[0],
+      //   engine.X[6].limbs_storage.limbs[0],
+      //   engine.X[7].limbs_storage.limbs[0]
+      // );
 
     engine.loadBasicTwiddlesGeneric(basic_twiddles, twiddle_stride, log_data_stride, s_meta, tw_log_size, twiddles_offset, 5, inv, dit, true);
     if (dit) {
@@ -1515,7 +1515,7 @@ namespace mxntt {
 // #ifndef DCCT
     // reverse_input = eRevType::NaturalToMixedRev;
     reverse_input = eRevType::None;
-    is_normalize = false;
+    // is_normalize = false;
     dit = true;
     if (reverse_input != eRevType::None) {
       const bool is_reverse_in_place = (d_input == d_output);

diff --git a/icicle/src/ntt/thread_ntt.cu b/icicle/src/ntt/thread_ntt.cu
@@ -74,18 +74,18 @@ public:
     uint32_t exp;
 
     // if ((s_meta.ntt_block_id == 15 || s_meta.ntt_block_id == 16) && s_meta.ntt_inp_id == 0) {
-      printf(
-        "T: %d, II: %d, B: %d, block_size: %d, tw_order: %d, tw_log_order: %d, tw_log_size: %d, block_offset: %d, ntt_inp_offset: %d\n",
-        threadIdx.x,
-        s_meta.ntt_inp_id,
-        s_meta.ntt_block_id,
-        s_meta.ntt_block_size,
-        tw_order,
-        tw_log_order,
-        tw_log_size,
-        block_offset,
-        ntt_inp_offset
-      );
+      // printf(
+      //   "T: %d, II: %d, B: %d, block_size: %d, tw_order: %d, tw_log_order: %d, tw_log_size: %d, block_offset: %d, ntt_inp_offset: %d\n",
+      //   threadIdx.x,
+      //   s_meta.ntt_inp_id,
+      //   s_meta.ntt_block_id,
+      //   s_meta.ntt_block_size,
+      //   tw_order,
+      //   tw_log_order,
+      //   tw_log_size,
+      //   block_offset,
+      //   ntt_inp_offset
+      // );
     // }
 
     UNROLL
@@ -98,15 +98,15 @@ public:
 
         // if ((s_meta.ntt_block_id == 255 && s_meta.ntt_inp_id == 3) || (s_meta.ntt_block_id == 0 && s_meta.ntt_inp_id == 0) || (s_meta.ntt_block_id == 511 && s_meta.ntt_inp_id == 1)) {
         // if ((s_meta.ntt_block_id == 15 || s_meta.ntt_block_id == 16) && s_meta.ntt_inp_id == 0) {
-          printf(
-            "T: %d, I: %d, II: %d, B: %d, exp: %d, tw: 0x%x\n",
-            threadIdx.x,
-            stage * 4 + i,
-            s_meta.ntt_inp_id,
-            s_meta.ntt_block_id,
-            exp,
-            basic_twiddles[inv ? (tw_size - 1 - exp) : exp].limbs_storage.limbs[0]
-          );
+          // printf(
+          //   "T: %d, I: %d, II: %d, B: %d, exp: %d, tw: 0x%x\n",
+          //   threadIdx.x,
+          //   stage * 4 + i,
+          //   s_meta.ntt_inp_id,
+          //   s_meta.ntt_block_id,
+          //   exp,
+          //   basic_twiddles[inv ? (tw_size - 1 - exp) : exp].limbs_storage.limbs[0]
+          // );
         // }
 
         WB[stage * 4 + i] = basic_twiddles[inv ? (tw_size - 1 - exp) : exp];