Change the dpcallback to take a pointer to the dp

The dpcallback used to take a C++ string containing the bytes of the
numerical value of the DP.  Now it just takes a pointer, and can do what
it likes.

dlrho.cc

@@ -54,7 +54,7 @@ struct CBData {
 // 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, string x,
+    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)
 {
@@ -77,6 +77,8 @@ bool dpcallback(void *cbdata, unsigned short threadId,
     // 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));
+
     res = d->dtable.insert(DTable::value_type(x, ab));
     if (!res.second) {
 	// Collision!

dpstream.cu

@@ -83,7 +83,7 @@ __device__ inline unsigned int *DPstreamAlloc()
 }
 
 extern bool dpcallback(void *data, unsigned short threadId,
-    unsigned short blockId, unsigned int demux, string x,
+    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);
 
@@ -115,9 +115,9 @@ bool DPstreamParse(void *data)
 	    unsigned int b_0 = dp[WORDS+4];
 	    unsigned int b_1 = dp[WORDS+5];
 	    unsigned int b_2 = dp[WORDS+6];
-	    string x((const char *)(dp+1), WORDS*sizeof(unsigned int));
 
-	    if (dpcallback(data, threadId, blockId, dp[DEMUXWORD], x, a_0, a_1, a_2, b_0, b_1, b_2)) {
+	    if (dpcallback(data, threadId, blockId, dp[DEMUXWORD], dp+1,
+			    a_0, a_1, a_2, b_0, b_1, b_2)) {
 		stop_computing = true;
 	    }
 	}