blackbeard/miner
1
0
Fork 0
Code Issues 2 Pull Requests Actions Packages Projects Releases 2 Wiki Activity

Fix conditional use of n0_inv in qpow Montgomery kernels

Browse Source
This commit is contained in:
rob thijssen
2025-09-12 18:35:01 +03:00
parent 398c9f9bc9
commit 943d35100f
1 changed files with 3 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
crates/engine-gpu-cuda/src/kernels/qpow_kernel.cu
+3 -3
View File
@@ -241,8 +241,7 @@ extern "C" __global__ void qpow_montgomery_g1_kernel(
// Transform to Montgomery domain
uint64_t yhat[8];
const uint64_t n0i = C_CONSTS_READY ? C_N0_INV : n0_inv;
to_mont_512(y0_loc, r2_loc, n_loc, n0i, yhat);
to_mont_512(y0_loc, r2_loc, n_loc, n0_inv, yhat);
// Iterate and emit normalized y per step
// Output stride per thread: iters_per_thread * 8 limbs
@@ -528,8 +527,9 @@ extern "C" __global__ void qpow_montgomery_g2_kernel(
for (int i = 0; i < 8; ++i) {
y0_loc[i] = y0[tid * 8u + i];
}
const uint64_t n0i = C_CONSTS_READY ? C_N0_INV : n0_inv;
uint64_t yhat[8];
to_mont_512(y0_loc, r2_loc, n_loc, n0_inv, yhat);
to_mont_512(y0_loc, r2_loc, n_loc, n0i, yhat);
// Iterate and check threshold
const uint32_t iters = iters_per_thread;