Fixed the CPU validation.

The previous version of the CPU validation only worked if a distinguished point was not found before 'nlaunch' kernel launches. This occured most commonly when 'nlaunch' was high and 'nmult' was low, or 'nthreads' was high.

parrhoasm.cu

@@ -20,6 +20,7 @@
 #include "cudadl.h"
 #include <fstream>
 #include <NTL/vec_ZZ.h>
+#include <NTL/vec_ZZ_p.h>
 #include <string.h>
 #include <limits.h>
 #include <sys/time.h>
@@ -407,12 +408,13 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
     //dump("y", l_y, WORDS);
 
 #ifdef CHECK_RESULTS
-    int totmult = 0;
-    bool fail = false;
-    vec_ZZ aexp, bexp;  // The expected values of a and b at completion
-    aexp.SetLength(nthreads);
-    bexp.SetLength(nthreads);
+    vec_ZZ astart, bstart;  // the starting values of a and b
+    vec_ZZ_p zz_p_xr_start;
+    astart.SetLength(nthreads);
+    bstart.SetLength(nthreads);
+    zz_p_xr_start.SetLength(nthreads);
 #endif
+
     ZZ a = RandomBits_ZZ(orderbits-1);
     ZZ b = RandomBits_ZZ(orderbits-1);
     ZZ astep = RandomBits_ZZ(orderbits-1);
@@ -431,35 +433,15 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 	// dump("l_x", l_ts[t].x, WORDS);
 	BytesFromZZ((unsigned char *) l_ts[t].a, a, 3 * sizeof(unsigned int));
 	BytesFromZZ((unsigned char *) l_ts[t].b, b, 3 * sizeof(unsigned int));
-	//dump("x", l_ts[t].x, WORDS);
-//cout << "x = " << rep(zz_p_x) << "\n";
-//cout << "xr = " << rep(zz_p_xr) << "\n";
-	//dump("rho", rho, WORDS);
 #ifdef CHECK_RESULTS
-	for (int l = 0; l < nmult * nlaunch; ++l)
-	{
-	    unsigned int w = trunc_long(rep(zz_p_xr), 32);
-	    if (w < TWO_32_DIV_3) {
-		zz_p_xr *= base;
-		a++;
-	    } else if (w < TWO_32_DIV_3_X2) {
-		zz_p_xr *= target;
-		b++;
-	    } else {
-		zz_p_xr *= zz_p_xr * zz_p_r_inv;
-		a += a;
-		b += b;
-	    }
-	    if (a > order) a -= order;
-	    if (b > order) b -= order;
-	    ++totmult;
-	}     
-	BytesFromZZ((unsigned char *) (l_z + t * WORDS), rep(zz_p_xr), WORDS * sizeof(unsigned int));
-	aexp[t] = a;
-	bexp[t] = b;
+	astart[t] = a;
+	bstart[t] = b;
+	zz_p_xr_start[t] = zz_p_xr;
 #endif
-//cout << "zr = " << rep(zz_p_xr) << "\n";
-	//dump("z", l_z + t * WORDS, WORDS);
+	//dump("x", l_ts[t].x, WORDS);
+	//cout << "x = " << rep(zz_p_x) << "\n";
+	//cout << "xr = " << rep(zz_p_xr) << "\n";
+	//dump("rho", rho, WORDS);
     }
 
     cudaMemcpy(d_ts, l_ts, nthreads * sizeof(GlobalThreadState), cudaMemcpyHostToDevice);
@@ -473,16 +455,16 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
     while(stop_computing == false)
 #endif
     {
+	launchcount++;
 #ifdef VERBOSE
-	cerr << getpid() << " Launch " << ++launchcount << "...\n";
-#else
-	++launchcount;
+	cerr << getpid() << " Launch " << launchcount << "...\n";
 #endif
 	cudaMulmod<<< nblocks, nthreads/nblocks /*tpb*/, 0 >>>(d_ts,
 	    order_0, order_1, order_2, dpfreq);
+	checkCUDAError("kernel launch");
 
 	cudaThreadSynchronize();
-	checkCUDAError("kernel launch");
+	checkCUDAError("thread sync");
 	stop_computing = DPstreamParse(cbdata);
 	//cudaPrintfExtractDPE(dpcallback, cbdata);
 	//cudaPrintfDisplay(stdout, true);
@@ -493,7 +475,48 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
     gettimeofday(&et, NULL);
 
     totmicros = (unsigned long long)(et.tv_sec - st.tv_sec) * 1000000 + (et.tv_usec - st.tv_usec);
+
 #ifdef CHECK_RESULTS
+    int totmult = 0;
+    bool fail = false;
+    vec_ZZ aexp, bexp;  // the expected values of a and b at completion
+    aexp.SetLength(nthreads);
+    bexp.SetLength(nthreads);
+    ZZ atemp;
+    ZZ btemp;
+    ZZ_p zz_p_xr_temp;
+
+    for (int t=0; t<nthreads; ++t)
+    {
+	atemp = astart[t];
+	btemp = bstart[t];
+	zz_p_xr_temp = zz_p_xr_start[t];
+
+	// launchcount should be <= nlaunch
+	for (int l = 0; l < nmult * launchcount; ++l)
+	{
+	    unsigned int w = trunc_long(rep(zz_p_xr_temp), 32);
+	    if (w < TWO_32_DIV_3) {
+		zz_p_xr_temp *= base;
+		atemp++;
+	    } else if (w < TWO_32_DIV_3_X2) {
+		zz_p_xr_temp *= target;
+		btemp++;
+	    } else {
+		zz_p_xr_temp *= zz_p_xr_temp * zz_p_r_inv;
+		atemp += atemp;
+		btemp += btemp;
+	    }
+	    if (atemp > order) atemp -= order;
+	    if (btemp > order) btemp -= order;
+	    ++totmult;
+	}
+	BytesFromZZ((unsigned char *) (l_z + t * WORDS), rep(zz_p_xr_temp), WORDS * sizeof(unsigned int));
+	aexp[t] = atemp;
+	bexp[t] = btemp;
+	//cout << "zr = " << rep(zz_p_xr_temp) << "\n";
+	//dump("z", l_z + t * WORDS, WORDS);
+    }
     cudaMemcpy(l_ts, d_ts, nthreads * sizeof(GlobalThreadState), cudaMemcpyDeviceToHost);
     checkCUDAError("memcpy");