ntt/*: omit radix6_twiddles_12.

supranational · Jan 28, 2025 · cb1bc09 · cb1bc09
1 parent ce5854c
commit cb1bc09
Show file tree

Hide file tree

Showing 4 changed files with 3 additions and 23 deletions.
diff --git a/ntt/kernels/ct_mixed_radix_wide.cu b/ntt/kernels/ct_mixed_radix_wide.cu
@@ -213,15 +213,6 @@ public:
                 _CT_NTT<1><<<NTT_CONFIGURATION>>>(NTT_ARGUMENTS);
             }
             break;
-        case 12:
-            if (iterations <= 6) {
-                intermediate_twiddle_shift = 6;
-                d_intermediate_twiddles = ntt_parameters.radix6_twiddles_12;
-                _CT_NTT<2><<<NTT_CONFIGURATION>>>(NTT_ARGUMENTS);
-            } else {
-                _CT_NTT<1><<<NTT_CONFIGURATION>>>(NTT_ARGUMENTS);
-            }
-            break;
         default:
             _CT_NTT<1><<<NTT_CONFIGURATION>>>(NTT_ARGUMENTS);
             break;

diff --git a/ntt/kernels/gs_mixed_radix_wide.cu b/ntt/kernels/gs_mixed_radix_wide.cu
@@ -209,15 +209,6 @@ public:
                 _GS_NTT<1><<<NTT_CONFIGURATION>>>(NTT_ARGUMENTS);
             }
             break;
-        case 12:
-            if (iterations <= 6) {
-                intermediate_twiddle_shift = 6;
-                d_intermediate_twiddles = ntt_parameters.radix6_twiddles_12;
-                _GS_NTT<2><<<NTT_CONFIGURATION>>>(NTT_ARGUMENTS);
-            } else {
-                _GS_NTT<1><<<NTT_CONFIGURATION>>>(NTT_ARGUMENTS);
-            }
-            break;
         default:
             _GS_NTT<1><<<NTT_CONFIGURATION>>>(NTT_ARGUMENTS);
             break;

diff --git a/ntt/ntt.cuh b/ntt/ntt.cuh
@@ -105,7 +105,7 @@ private:
 
         if (lg_domain_size <= 10) {
             params.step(lg_domain_size);
-        } else if (lg_domain_size <= 17) {
+        } else if (lg_domain_size <= 18) {
             int step = lg_domain_size / 2;
             params.step(step + lg_domain_size % 2);
             params.step(step);
@@ -135,7 +135,7 @@ private:
 
         if (lg_domain_size <= 10) {
             params.step(lg_domain_size);
-        } else if (lg_domain_size <= 17) {
+        } else if (lg_domain_size <= 18) {
             int step = lg_domain_size / 2;
             params.step(step);
             params.step(step + lg_domain_size % 2);

diff --git a/ntt/parameters.cuh b/ntt/parameters.cuh
@@ -186,7 +186,7 @@ public:
     fr_t (*partial_group_gen_powers)[WINDOW_SIZE]; // for LDE
 
 #if !defined(FEATURE_BABY_BEAR) && !defined(FEATURE_GOLDILOCKS)
-    fr_t* radix6_twiddles_6, * radix6_twiddles_12, * radix7_twiddles_7,
+    fr_t* radix6_twiddles_6, * radix7_twiddles_7,
         * radix8_twiddles_8, * radix9_twiddles_9;
 
 private:
@@ -235,7 +235,6 @@ public:
 
 #if !defined(FEATURE_BABY_BEAR) && !defined(FEATURE_GOLDILOCKS)
         radix6_twiddles_6 = twiddles_X(64, 64, roots[12]);
-        radix6_twiddles_12 = twiddles_X(4096, 64, roots[18]);
         radix7_twiddles_7 = twiddles_X(128, 128, roots[14]);
         radix8_twiddles_8 = twiddles_X(256, 256, roots[16]);
         radix9_twiddles_9 = twiddles_X(512, 512, roots[18]);
@@ -270,7 +269,6 @@ public:
             (void)cudaFreeAsync(radix9_twiddles_9, gpu);
             (void)cudaFreeAsync(radix8_twiddles_8, gpu);
             (void)cudaFreeAsync(radix7_twiddles_7, gpu);
-            (void)cudaFreeAsync(radix6_twiddles_12, gpu);
             (void)cudaFreeAsync(radix6_twiddles_6, gpu);
 #endif
             (void)cudaFreeAsync(twiddles[1], gpu);