Have the dpcallback consume an unparsed array of unsigned ints

Don't have the CUDA driver parse the array of unsigned ints logged
by the CUDA kernel for each DP it found.  Rather, just pass a pointer
to that array directly to the dpcallback.

The WORDS+7 unsigned ints are:
 - 1 word of threadID/blockID
 - WORDS words of the dp value
 - 3 words of a
 - 3 words of b

dlrho.cc

@@ -53,31 +53,33 @@ struct CBData {
 // It calls the function named "dpcallback" directly.  It would be
 // cleaner if this were passed as a function pointer to cuda_dl, but
 // that makes nvcc 3.1 segfault. :-p
-bool dpcallback(void *cbdata, unsigned short threadId,
-    unsigned short blockId, unsigned int demux, unsigned int *dpwords,
-    unsigned int a_0, unsigned int a_1, unsigned int a_2,
-    unsigned int b_0, unsigned int b_1, unsigned int b_2)
+// dp points to an array of WORDS+7 unsigned ints:
+// - 1 word of threadID/blockID
+// - WORDS words of the dp value
+// - 3 words of a
+// - 3 words of b
+bool dpcallback(void *cbdata, unsigned int *dpwords)
 {
     CBData *d = (CBData*)cbdata;
     // WARNING: this assumes
     //    sizeof(unsigned long) == sizeof(unsigned long long) !
-    ZZ zz_a = to_ZZ(a_2);
+    ZZ zz_a = to_ZZ(dpwords[WORDS+3]);
     zz_a <<= 32;
-    zz_a += a_1;
+    zz_a += dpwords[WORDS+2];
     zz_a <<= 32;
-    zz_a += a_0;
-    ZZ zz_b = to_ZZ(b_2);
+    zz_a += dpwords[WORDS+1];
+    ZZ zz_b = to_ZZ(dpwords[WORDS+6]);
     zz_b <<= 32;
-    zz_b += b_1;
+    zz_b += dpwords[WORDS+5];
     zz_b <<= 32;
-    zz_b += b_0;
+    zz_b += dpwords[WORDS+4];
 
     //ZZ_p dp = power(d->base, zz_a) * power(d->target, zz_b);
 
     // cerr << "DP " << ++(d->numdp) << "\r";
     pair<ZZ,ZZ> ab(zz_a,zz_b);
     pair<DTable::iterator, bool> res;
-    string x((const char *)(dpwords), WORDS*sizeof(unsigned int));
+    string x((const char *)(dpwords+1), WORDS*sizeof(unsigned int));
 
     res = d->dtable.insert(DTable::value_type(x, ab));
     if (!res.second) {

dpstream.cu

@@ -82,16 +82,12 @@ __device__ inline unsigned int *DPstreamAlloc()
     return ourbuffer;
 }
 
-extern bool dpcallback(void *data, unsigned short threadId,
-    unsigned short blockId, unsigned int demux, unsigned int *dpwords,
-    unsigned int a_0, unsigned int a_1, unsigned int a_2,
-    unsigned int b_0, unsigned int b_1, unsigned int b_2);
-
-#if WORDS > 1
-#define DEMUXWORD 2
-#else
-#define DEMUXWORD 1
-#endif
+// dp points to an array of WORDS+7 unsigned ints:
+// - 1 word of threadID/blockID
+// - WORDS words of the dp value
+// - 3 words of a
+// - 3 words of b
+extern bool dpcallback(void *data, unsigned int *dpwords);
 
 // Return true if we should stop computing
 bool DPstreamParse(void *data)
@@ -107,17 +103,7 @@ bool DPstreamParse(void *data)
 	    cudaMemcpyDeviceToHost);
 
 	for (unsigned int *dp=dpbuf; dp < dpbuf+bufsize; dp += DPrecordsize) {
-	    unsigned short threadId = (unsigned short)(dp[0]);
-	    unsigned short blockId = dp[0]>>16;
-	    unsigned int a_0 = dp[WORDS+1];
-	    unsigned int a_1 = dp[WORDS+2];
-	    unsigned int a_2 = dp[WORDS+3];
-	    unsigned int b_0 = dp[WORDS+4];
-	    unsigned int b_1 = dp[WORDS+5];
-	    unsigned int b_2 = dp[WORDS+6];
-
-	    if (dpcallback(data, threadId, blockId, dp[DEMUXWORD], dp+1,
-			    a_0, a_1, a_2, b_0, b_1, b_2)) {
+	    if (dpcallback(data, dp)) {
 		stop_computing = true;
 	    }
 	}

worker.cc

@@ -49,29 +49,29 @@ static struct WrkControllerState {
 
 // ----- Below this line are the functions running in the worker thread.
 
+#if WORDS > 1
+#define DEMUXWORD 8
+#else
+#define DEMUXWORD 7
+#endif
 
 // This function is called from inside cuda_dl for each DP it encounters.
 // It calls the function named "dpcallback" directly.  It would be
 // cleaner if this were passed as a function pointer to cuda_dl, but
 // that makes nvcc 3.1 segfault. :-p
-bool dpcallback(void *cbdata, unsigned short threadId,
-    unsigned short blockId, unsigned int demux, unsigned int *dpwords,
-    unsigned int a_0, unsigned int a_1, unsigned int a_2,
-    unsigned int b_0, unsigned int b_1, unsigned int b_2)
+// dp points to an array of WORDS+7 unsigned ints:
+// - 1 word of threadID/blockID
+// - WORDS words of the dp value
+// - 3 words of a
+// - 3 words of b
+bool dpcallback(void *cbdata, unsigned int *dpwords)
 {
+    unsigned int demux = dpwords[DEMUXWORD];
+
     struct bufferevent *bev =
 	wrkctrlstate.dpnodes[demux % wrkctrlstate.num_connected_dpnodes];
 
-    unsigned char DPbuf[(WORDS+6)*sizeof(unsigned int)];
-    memmove(DPbuf, dpwords, WORDS*sizeof(unsigned int));
-    memmove(DPbuf+WORDS*sizeof(unsigned int), &a_0, sizeof(unsigned int));
-    memmove(DPbuf+(WORDS+1)*sizeof(unsigned int), &a_1, sizeof(unsigned int));
-    memmove(DPbuf+(WORDS+2)*sizeof(unsigned int), &a_2, sizeof(unsigned int));
-    memmove(DPbuf+(WORDS+3)*sizeof(unsigned int), &b_0, sizeof(unsigned int));
-    memmove(DPbuf+(WORDS+4)*sizeof(unsigned int), &b_1, sizeof(unsigned int));
-    memmove(DPbuf+(WORDS+5)*sizeof(unsigned int), &b_2, sizeof(unsigned int));
-
-    bufferevent_write(bev, DPbuf, (WORDS+6)*sizeof(unsigned int));
+    bufferevent_write(bev, dpwords+1, (WORDS+6)*sizeof(unsigned int));
 
     // If worker_thread_state changes to WT_SHOULD_STOP, then signal to
     // stop computation by returning true.  If for some reason, it