Improved the logging for dlrho.

Each line of the output from every dlrho subprocess is now written atomically. The timing information (number of kernel launches) is also written to stdout.

atomic_iostream.h

@@ -0,0 +1,21 @@
+#include <sstream>
+#include <iostream>
+
+// copied from StackOverflow: https://stackoverflow.com/a/34403537
+
+class AtomicWriter {
+    std::ostringstream st;
+    std::ostream &stream;
+public:
+    AtomicWriter(std::ostream &s=std::cout):stream(s) { }
+    template <typename T>
+    AtomicWriter& operator<<(T const& t) {
+	st << t;
+	return *this;
+    }
+    AtomicWriter& operator<<( std::ostream&(*f)(std::ostream&) ) {
+	st << f;
+	return *this;
+    }
+    ~AtomicWriter() { stream << st.str(); }
+};

dlrho.cc

@@ -30,6 +30,7 @@
 #include <sys/socket.h>
 #include <cuda_runtime.h>
 
+#include "atomic_iostream.h"
 #include "cudadl.h"
 
 NTL_CLIENT
@@ -112,7 +113,7 @@ bool dpcallback(void *cbdata, unsigned int *dpwords)
 // label is "p" or "q", to be printed to report progress.
 // Return 0 on failure, 1 on success.
 static int p_dl(const ZZ_p &target, const ZZ_p &base, ZZ &exp,
-    const vec_ZZ &fvec, const string &label)
+    const vec_ZZ &fvec, const string &label, unsigned int initial_subproblem_id)
 {
     const int flen = fvec.length();
 
@@ -129,23 +130,29 @@ static int p_dl(const ZZ_p &target, const ZZ_p &base, ZZ &exp,
     curexp = 0;
 
     for (int i = 0; i < flen; ++i) {
-	cout << label << " submodulus " << i+1 << " of " << flen << "... ";
-	cout.flush();
-	// Figure out exp mod fvec[i] by taking each side to the power
-	// of phirho/fvec[i] so that we're working in the
-	// subgroup of order fvec[i].
-	ZZ quotient = phip / fvec[i];
-	ZZ_p subgroup_base = power(base, quotient);
-	ZZ_p subgroup_target = power(target, quotient);
-
-	if (subgroup_base == 1) {
-	    // The original base wasn't a generator of the whole group.
-	    if (subgroup_target == 1) {
-		cout << "Non-unique solution (mod " << fvec[i] <<")\n";
-		continue;
-	    } else {
-		cout << "Target not in subgroup generated by base\n";
-		return 0;
+	ZZ quotient;
+	ZZ_p subgroup_base;
+	ZZ_p subgroup_target;
+	
+	{
+	    AtomicWriter atomic_cout(cout);
+	    atomic_cout << label << " submodulus " << i+1 << " of " << flen << "...\n";
+	    // Figure out exp mod fvec[i] by taking each side to the power
+	    // of phirho/fvec[i] so that we're working in the
+	    // subgroup of order fvec[i].
+	    quotient = phip / fvec[i];
+	    subgroup_base = power(base, quotient);
+	    subgroup_target = power(target, quotient);
+
+	    if (subgroup_base == 1) {
+		// The original base wasn't a generator of the whole group.
+		if (subgroup_target == 1) {
+		    atomic_cout << "Non-unique solution (mod " << fvec[i] <<")\n";
+		    continue;
+		} else {
+		    atomic_cout << "Target not in subgroup generated by base\n";
+		    return 0;
+		}
 	    }
 	}
 
@@ -153,7 +160,6 @@ static int p_dl(const ZZ_p &target, const ZZ_p &base, ZZ &exp,
 	// subgroup_target with base subgroup_base, knowing that it's in
 	// the range [0,fvec[i]).
 
-	cout << "\n";
 	ZZ md = ZZ_p::modulus();
 	CBData cbdata(subgroup_base, subgroup_target, fvec[i]);
 	struct timeval st, et;
@@ -171,6 +177,11 @@ static int p_dl(const ZZ_p &target, const ZZ_p &base, ZZ &exp,
 	printf("%ld.%06ld seconds elapsed\n", us_elapsed/1000000,
 	    us_elapsed % 1000000);
 
+	AtomicWriter(cout) << "Timing (subproblemid, label, launches): "
+	                   << i+initial_subproblem_id << ", "
+	                   << label << ", "
+	                   << launch_count << "\n" << std::flush;
+
 	CRT(curexp, curmodulus, subgroup_dl, fvec[i]);
 	// cout << "CRT\n";
 	// cout << "curexp = " << curexp << "\n";
@@ -192,7 +203,7 @@ typedef struct {
 
 // Behave like p_dl, but do the work in an asynchronous subprocess
 static PDLHandle* p_dl_fork_start(const ZZ &p, const ZZ_p &target,
-    const ZZ_p &base, const vec_ZZ &fvec, const string &label, int deviceid)
+    const ZZ_p &base, const vec_ZZ &fvec, const string &label, int deviceid, unsigned int initial_subproblem_id)
 {
     PDLHandle *handle = new PDLHandle;
     int fds[2];
@@ -203,6 +214,11 @@ static PDLHandle* p_dl_fork_start(const ZZ &p, const ZZ_p &target,
 	return NULL;
     }
     handle->rfd = fds[0];
+
+    cout.flush();
+    cerr.flush();
+    // flush the output before forking
+
     pid_t childpid = fork();
     if (childpid == -1) {
 	perror("fork");
@@ -222,10 +238,10 @@ static PDLHandle* p_dl_fork_start(const ZZ &p, const ZZ_p &target,
 
 	cudaError_t cudares = cudaSetDevice(deviceid);
 	if (cudares != cudaSuccess) {
-	    cerr << "Error setting CUDA device: " << cudaGetErrorString(cudares) << "\n";
+	    AtomicWriter(cerr) << "Error setting CUDA device: " << cudaGetErrorString(cudares) << "\n";
 	    exit(1);
 	}
-	int res = p_dl(target_p, base_p, exp, fvec, label);
+	int res = p_dl(target_p, base_p, exp, fvec, label, initial_subproblem_id);
 
 	if (res) {
 	    // Write the result back to the parent using wfd
@@ -303,8 +319,8 @@ int main(int argc, char **argv)
 
     PDLHandle *handle_p, *handle_q;
 
-    handle_p = p_dl_fork_start(p, target, base, pfvec, "p", 0);
-    handle_q = p_dl_fork_start(q, target, base, qfvec, "q", 1);
+    handle_p = p_dl_fork_start(p, target, base, pfvec, "p", 0, 0);
+    handle_q = p_dl_fork_start(q, target, base, qfvec, "q", 1, pfvec.length());
     res_p = p_dl_fork_join(handle_p, exp_p);
     res_q = p_dl_fork_join(handle_q, exp_q);
 

dpstream.cu

@@ -91,14 +91,13 @@ __device__ inline unsigned int *DPstreamAlloc()
 extern bool dpcallback(void *data, unsigned int *dpwords);
 
 // Return true if we should stop computing
-bool DPstreamParse(void *data)
+bool DPstreamParse(void *data, unsigned int *num_dps)
 {
     // Get the number of words in the buffer
     unsigned int bufsize;
     cudaMemcpyFromSymbol(&bufsize, DPbuffertail, sizeof(unsigned int));
-#ifdef VERBOSE
-    cerr << getpid() << " " << bufsize / DPrecordsize << " DPs\n";
-#endif
+    *num_dps = bufsize/DPrecordsize;
+
     unsigned int *dpbuf = (unsigned int*)calloc(bufsize, sizeof(unsigned int));
     bool stop_computing = false;
     if (dpbuf) {

parrhoasm.cu

@@ -31,6 +31,7 @@
 #include <utility>
 #include <map>
 
+#include "atomic_iostream.h"
 #include "worker.h"
 
 // #define CHECK_RESULTS
@@ -334,13 +335,13 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 
     if (NumBits(modulus) <= ((WORDS-1)*BITS_PER_WORD) ||
 	    NumBits(modulus) > WORDS*BITS_PER_WORD) {
-	cerr << "modulus is not " << WORDS << " words long.\n";
+	AtomicWriter(cerr) << "modulus is not " << WORDS << " words long.\n";
 	exit(1);
     }
 
     long orderbits = NumBits(order);
     if (orderbits > 92) {
-	cerr << "order is larger than 92 bits.\n";
+	AtomicWriter(cerr) << "order is larger than 92 bits.\n";
 	exit(1);
     }
 
@@ -391,7 +392,7 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 */
 
     if (nthreads % nblocks) {
-	cerr << "Error: " << nthreads << " not a multiple of " << nblocks << "\n";
+	AtomicWriter(cerr) << "Error: " << nthreads << " not a multiple of " << nblocks << "\n";
 	exit(1);
     }
     const int threadsPerBlock = nthreads / nblocks;
@@ -461,6 +462,7 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
     gettimeofday(&st, NULL);
     int local_launchcount = 0;
     bool stop_computing = false;
+    unsigned int num_dps = 0;
 #ifdef CHECK_RESULTS
     for (int ln=0; ln<nlaunch && stop_computing == false; ++ln)
 #else
@@ -469,7 +471,7 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
     {
 	local_launchcount++;
 #ifdef VERBOSE
-	cerr << getpid() << " Launch " << local_launchcount << "...\n";
+	AtomicWriter(cerr) << getpid() << " Launch " << local_launchcount << "...\n";
 #endif
 	cudaMulmod<<< nblocks, nthreads/nblocks /*tpb*/, 0 >>>(d_ts,
 	    order_0, order_1, order_2, dpfreq);
@@ -477,12 +479,12 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 
 	cudaThreadSynchronize();
 	checkCUDAError("thread sync");
-	stop_computing = DPstreamParse(cbdata);
-	//cudaPrintfExtractDPE(dpcallback, cbdata);
-	//cudaPrintfDisplay(stdout, true);
+	stop_computing = DPstreamParse(cbdata, &num_dps);
 #ifdef VERBOSE
-	cerr << getpid() << "\n";
+	AtomicWriter(cerr) << getpid() << " Finished " << local_launchcount << " (" << num_dps << " DPs)\n";
 #endif
+	//cudaPrintfExtractDPE(dpcallback, cbdata);
+	//cudaPrintfDisplay(stdout, true);
     }
     gettimeofday(&et, NULL);
 
@@ -552,15 +554,16 @@ cout << i << ": " << l_Z[i] << " != " << l_z[i + t * WORDS] << "\n";
 	ZZ ares = (to_ZZ(l_ts[t].a[2]) << 64) + (to_ZZ(l_ts[t].a[1]) << 32) + to_ZZ(l_ts[t].a[0]);
 	ZZ bres = (to_ZZ(l_ts[t].b[2]) << 64) + (to_ZZ(l_ts[t].b[1]) << 32) + to_ZZ(l_ts[t].b[0]);
 	if (ares != aexp[t] || bres != bexp[t]) {
-	    cerr << "ares = " << ares << "\n";
-	    cerr << "aexp = " << aexp[t] << "\n";
-	    cerr << "bres = " << bres << "\n";
-	    cerr << "bexp = " << bexp[t] << "\n";
+	    AtomicWriter atomic_cerr(cerr);
+	    atomic_cerr << "ares = " << ares << "\n";
+	    atomic_cerr << "aexp = " << aexp[t] << "\n";
+	    atomic_cerr << "bres = " << bres << "\n";
+	    atomic_cerr << "bexp = " << bexp[t] << "\n";
 	}
 	++j;
     }
     if (!fail) {
-	cerr << "Results correct.\n";
+	AtomicWriter(cerr) << "Results correct.\n";
     }
 #endif
 
@@ -568,7 +571,7 @@ cout << i << ": " << l_Z[i] << " != " << l_z[i + t * WORDS] << "\n";
     sprintf(timestamp, "%d.%06d", et.tv_sec, et.tv_usec);
 
     unsigned long long totnanos = totmicros * 1000;
-    cout << timestamp << ":" << output_prefix << ": " << totmicros << " us / " << nthreads << " = " << totmicros / nthreads << " us / " << (nmult*local_launchcount) << " = " << totnanos / ((unsigned long long)nthreads * nmult * local_launchcount) << " ns\n";
+    AtomicWriter(cout) << timestamp << ":" << output_prefix << ": " << totmicros << " us / " << nthreads << " = " << totmicros / nthreads << " us / " << (nmult*local_launchcount) << " = " << totnanos / ((unsigned long long)nthreads * nmult * local_launchcount) << " ns\n";
     cout.flush();
     //cudaPrintfEnd();
     DPstreamEnd();
@@ -698,7 +701,8 @@ int main(int argc, char** argv)
 	// Try to find the DL_g of y
 	cerr << "DL_" << g << "(" << y << ") mod " << p << "\n";
 	CBData cbdata(g, y, pfvec[f]);
-	cuda_dl(g, y, pfvec[f], p, 4294967, &cbdata);
+	unsigned int* launch_count;
+	cuda_dl(g, y, pfvec[f], p, 4294967, &cbdata, &launch_count);
 	cerr << "e = " << cbdata.expon << "\n";
 	cerr << ( (power(g,cbdata.expon) == y) ? "CORRECT!" : "INCORRECT!" ) << "\n";
     }