Prepare to move the dpcallback out of parrhoasm.cu

So that we will eventually be able to replace the local distinguished
points table with a distributed one, touch up dpcallback so that it will
be able to be lifted out of parrhoasm.cu entirely.

dpstream.cu

@@ -83,17 +83,26 @@ __device__ inline unsigned int *DPstreamWrite(unsigned int a_0, unsigned int a_1
     return ourbuffer;
 }
 
-static void dpcallback(void *data, unsigned short threadId,
-    unsigned short blockId, string x, unsigned int a_0, unsigned int a_1,
-    unsigned int a_2, unsigned int b_0, unsigned int b_1, unsigned int b_2);
+static bool dpcallback(void *data, unsigned short threadId,
+    unsigned short blockId, unsigned int demux, string x,
+    unsigned int a_0, unsigned int a_1, unsigned int a_2,
+    unsigned int b_0, unsigned int b_1, unsigned int b_2);
 
-void DPstreamParse(void *data)
+#if WORDS > 1
+#define DEMUXWORD 1
+#else
+#define DEMUXWORD 0
+#endif
+
+// Return true if we should stop computing
+bool DPstreamParse(void *data)
 {
     // Get the number of words in the buffer
     unsigned int bufsize;
     cudaMemcpyFromSymbol(&bufsize, DPbuffertail, sizeof(unsigned int));
     cerr << getpid() << " " << bufsize / DPrecordsize << " DPs\n";
     unsigned int *dpbuf = (unsigned int*)calloc(bufsize, sizeof(unsigned int));
+    bool stop_computing = false;
     if (dpbuf) {
 	cudaMemcpy(dpbuf, DPbuffer_device, bufsize * sizeof(unsigned int),
 	    cudaMemcpyDeviceToHost);
@@ -109,12 +118,16 @@ void DPstreamParse(void *data)
 	    unsigned int b_2 = dp[WORDS+6];
 	    string x((const char *)(dp+1), WORDS*sizeof(unsigned int));
 
-	    dpcallback(data, threadId, blockId, x, a_0, a_1, a_2, b_0, b_1, b_2);
+	    if (dpcallback(data, threadId, blockId, dp[DEMUXWORD], x, a_0, a_1, a_2, b_0, b_1, b_2)) {
+		stop_computing = true;
+	    }
 	}
     }
     free(dpbuf);
     // Reset the buffer
     unsigned int zero = 0;
     cudaMemcpyToSymbol(DPbuffertail, &zero, sizeof(unsigned int));
+
+    return stop_computing;
 }
 #endif

parrhoasm.cu

@@ -322,9 +322,10 @@ struct CBData {
 	found_collision(false) {}
 };
 
-static void dpcallback(void *cbdata, unsigned short threadId,
-    unsigned short blockId, string x, unsigned int a_0, unsigned int a_1,
-    unsigned int a_2, unsigned int b_0, unsigned int b_1, unsigned int b_2)
+static bool dpcallback(void *cbdata, unsigned short threadId,
+    unsigned short blockId, unsigned int demux, string x,
+    unsigned int a_0, unsigned int a_1, unsigned int a_2,
+    unsigned int b_0, unsigned int b_1, unsigned int b_2)
 {
     CBData *d = (CBData*)cbdata;
     // WARNING: this assumes
@@ -361,6 +362,8 @@ static void dpcallback(void *cbdata, unsigned short threadId,
 	    }
 	}
     }
+
+    return d->found_collision;
 }
 
 ZZ cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
@@ -521,7 +524,8 @@ ZZ cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 #ifdef CHECK_RESULTS
     for (int ln=0; ln<nlaunch; ++ln)
 #else
-    while(cbdata.found_collision == false)
+    bool stop_computing = false;
+    while(stop_computing == false)
 #endif
     {
 	cerr << getpid() << " Launch " << ++launchcount << "...\n";
@@ -529,7 +533,7 @@ ZZ cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 
 	cudaThreadSynchronize();
 	checkCUDAError("kernel launch");
-	DPstreamParse(&cbdata);
+	stop_computing = DPstreamParse(&cbdata);
 	//cudaPrintfExtractDPE(dpcallback, &cbdata);
 	//cudaPrintfDisplay(stdout, true);
 	cerr << getpid() << "\n";