|
@@ -322,8 +322,14 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
|
|
|
// cuPrintf("d_b = %08X%08X\n", global_b_base[1], global_b_base[0]);
|
|
// cuPrintf("d_b = %08X%08X\n", global_b_base[1], global_b_base[0]);
|
|
|
}
|
|
}
|
|
|
|
|
|
|
|
|
|
+#if defined(NTHREADS) || defined(NBLOCKS)
|
|
|
|
|
+// must define both together
|
|
|
|
|
+int nthreads = NTHREADS;
|
|
|
|
|
+int nblocks = NBLOCKS;
|
|
|
|
|
+#else
|
|
|
int nthreads = 25600;
|
|
int nthreads = 25600;
|
|
|
int nblocks = 50;
|
|
int nblocks = 50;
|
|
|
|
|
+#endif
|
|
|
|
|
|
|
|
void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
|
|
void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
|
|
|
const ZZ &modulus, unsigned int dpfreq, void *cbdata,
|
|
const ZZ &modulus, unsigned int dpfreq, void *cbdata,
|
|
@@ -463,6 +469,11 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
|
|
|
int local_launchcount = 0;
|
|
int local_launchcount = 0;
|
|
|
bool stop_computing = false;
|
|
bool stop_computing = false;
|
|
|
unsigned int num_dps = 0;
|
|
unsigned int num_dps = 0;
|
|
|
|
|
+
|
|
|
|
|
+#ifdef VERBOSE
|
|
|
|
|
+ AtomicWriter(cerr) << getpid() << " Using (nblocks, nthreads): " << nblocks << ", " << nthreads << "\n";
|
|
|
|
|
+#endif
|
|
|
|
|
+
|
|
|
#ifdef CHECK_RESULTS
|
|
#ifdef CHECK_RESULTS
|
|
|
for (int ln=0; ln<nlaunch && stop_computing == false; ++ln)
|
|
for (int ln=0; ln<nlaunch && stop_computing == false; ++ln)
|
|
|
#else
|
|
#else
|