From da08434db3dcfeeb94948346e3293bbbc8c3d1aa Mon Sep 17 00:00:00 2001
From: Andy Polyakov <appro@cryptogams.org>
Date: Wed, 27 Nov 2024 15:34:08 +0100
Subject: [PATCH] polynomial/div_by_x_minus_z.cuh: move a sanity check to the
 host side.

---
 polynomial/div_by_x_minus_z.cuh | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/polynomial/div_by_x_minus_z.cuh b/polynomial/div_by_x_minus_z.cuh
index 75f8143..44c46da 100644
--- a/polynomial/div_by_x_minus_z.cuh
+++ b/polynomial/div_by_x_minus_z.cuh
@@ -55,7 +55,9 @@ void d_div_by_x_minus_z(fr_t d_inout[], size_t len, fr_t z)
         }
     };
 
+#if 0
     assert(blockDim.x%WARP_SZ == 0 && gridDim.x <= blockDim.x);
+#endif
 
     const uint32_t tid    = threadIdx.x + blockDim.x*blockIdx.x;
     const uint32_t laneid = threadIdx.x % WARP_SZ;
@@ -379,6 +381,7 @@ void div_by_x_minus_z(fr_t d_inout[], size_t len, const fr_t& z,
             cudaFuncAttributes attr;
             CUDA_OK(cudaFuncGetAttributes(&attr, d_div_by_x_minus_z<fr_t, rotate, BSZ>));
             saved_blockDim = attr.maxThreadsPerBlock;
+            assert(saved_blockDim%WARP_SZ == 0);
         }
 
         blockDim = saved_blockDim;