Factor the dpcallback out of cuda_dl and into the calling code

This will allow for the multinode driver to have its own dpcallback.
cuda_dl calls the function named "dpcallback" directly now, expecting it
to be defined externally.  It would be cleaner if this were passed as
a function pointer, but this causes nvcc 3.1 to segfault.  :-p

cudadl.h

@@ -25,7 +25,12 @@
 
 NTL_CLIENT
 
-ZZ cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
-		const ZZ &modulus);
+typedef 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 cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
+		const ZZ &modulus, void *data);
 
 #endif

dlrho.cc

@@ -21,15 +21,82 @@
 #include <sys/time.h>
 #include <NTL/vec_ZZ.h>
 #include <NTL/ZZ_p.h>
-#include "cudadl.h"
+
+#include <map>
 
 #include <sys/types.h>
 #include <sys/wait.h>
 #include <sys/socket.h>
 #include <cuda_runtime.h>
 
+#include "cudadl.h"
+
 NTL_CLIENT
 
+typedef map<std::string, pair<ZZ,ZZ> > DTable;
+
+struct CBData {
+    const ZZ_p &base;
+    const ZZ_p &target;
+    const ZZ &order;
+    unsigned long long numdp;
+    DTable dtable;
+    bool found_collision;
+    ZZ expon;
+
+    CBData(const ZZ_p &_base, const ZZ_p &_target, const ZZ &_order) :
+	base(_base), target(_target), order(_order), numdp(0),
+	found_collision(false) {}
+};
+
+// 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, 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
+    //    sizeof(unsigned long) == sizeof(unsigned long long) !
+    ZZ zz_a = to_ZZ(a_2);
+    zz_a <<= 32;
+    zz_a += a_1;
+    zz_a <<= 32;
+    zz_a += a_0;
+    ZZ zz_b = to_ZZ(b_2);
+    zz_b <<= 32;
+    zz_b += b_1;
+    zz_b <<= 32;
+    zz_b += b_0;
+
+    //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;
+    res = d->dtable.insert(DTable::value_type(x, ab));
+    if (!res.second) {
+	// Collision!
+	ZZ adiff = to_ZZ(res.first->second.first) - zz_a;
+	ZZ bdiff = zz_b - to_ZZ(res.first->second.second);
+	if (bdiff < 0) bdiff += d->order;
+	ZZ binv;
+	if (InvModStatus(binv, bdiff, d->order) == 0) {
+	    d->expon = MulMod(binv, adiff, d->order);
+	    d->found_collision = true;
+	} else {
+	    if (d->order > (1<<20)) {
+		cerr << "Unhelpful collision\n";
+	    }
+	}
+    }
+
+    return d->found_collision;
+}
+
 // Compute the discrete log of target mod p, to the given base.
 // p must be the current ZZ_p modulus.
 // Place the result in exp.  fvec is a vector of the factors of
@@ -81,10 +148,11 @@ static int p_dl(const ZZ_p &target, const ZZ_p &base, ZZ &exp,
 
 	cout << "\n";
 	ZZ md = ZZ_p::modulus();
+	CBData cbdata(subgroup_base, subgroup_target, fvec[i]);
 	struct timeval st, et;
 	gettimeofday(&st, NULL);
-	ZZ subgroup_dl = cuda_dl(subgroup_base, subgroup_target, fvec[i],
-					md);
+	cuda_dl(subgroup_base, subgroup_target, fvec[i], md, &cbdata);
+	ZZ subgroup_dl = cbdata.expon;
 	gettimeofday(&et, NULL);
 	unsigned long us_elapsed = (et.tv_sec-st.tv_sec)*1000000 +
 	    (et.tv_usec-st.tv_usec);

dpstream.cu

@@ -82,7 +82,7 @@ __device__ inline unsigned int *DPstreamAlloc()
     return ourbuffer;
 }
 
-static bool dpcallback(void *data, unsigned short threadId,
+extern 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);

parrhoasm.cu

@@ -306,68 +306,8 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
 int nthreads = 25600;
 int nblocks = 50;
 
-typedef map<std::string, pair<ZZ,ZZ> > DTable;
-
-struct CBData {
-    const ZZ_p &base;
-    const ZZ_p &target;
-    const ZZ &order;
-    unsigned long long numdp;
-    DTable dtable;
-    bool found_collision;
-    ZZ expon;
-
-    CBData(const ZZ_p &_base, const ZZ_p &_target, const ZZ &_order) :
-	base(_base), target(_target), order(_order), numdp(0),
-	found_collision(false) {}
-};
-
-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
-    //    sizeof(unsigned long) == sizeof(unsigned long long) !
-    ZZ zz_a = to_ZZ(a_2);
-    zz_a <<= 32;
-    zz_a += a_1;
-    zz_a <<= 32;
-    zz_a += a_0;
-    ZZ zz_b = to_ZZ(b_2);
-    zz_b <<= 32;
-    zz_b += b_1;
-    zz_b <<= 32;
-    zz_b += b_0;
-
-    //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;
-    res = d->dtable.insert(DTable::value_type(x, ab));
-    if (!res.second) {
-	// Collision!
-	ZZ adiff = to_ZZ(res.first->second.first) - zz_a;
-	ZZ bdiff = zz_b - to_ZZ(res.first->second.second);
-	if (bdiff < 0) bdiff += d->order;
-	ZZ binv;
-	if (InvModStatus(binv, bdiff, d->order) == 0) {
-	    d->expon = MulMod(binv, adiff, d->order);
-	    d->found_collision = true;
-	} else {
-	    if (d->order > (1<<20)) {
-		cerr << "Unhelpful collision\n";
-	    }
-	}
-    }
-
-    return d->found_collision;
-}
-
-ZZ cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
-		const ZZ &modulus)
+void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
+		const ZZ &modulus, void *cbdata)
 {
     unsigned long long totmicros = 0;
     ZZ_pBak pbak;
@@ -406,7 +346,6 @@ ZZ cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
     zz_p_r = to_ZZ_p(zz_r);
     zz_p_r_inv = to_ZZ_p(1) / zz_p_r;
 
-    CBData cbdata(base, target, order);
 //cout << "r = " << zz_r << "\n";
 //cout << "r mod rho = " << rep(zz_p_r) << "\n";
 //cout << "r_inv = " << rep(zz_p_r_inv) << "\n";
@@ -533,8 +472,8 @@ ZZ cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 
 	cudaThreadSynchronize();
 	checkCUDAError("kernel launch");
-	stop_computing = DPstreamParse(&cbdata);
-	//cudaPrintfExtractDPE(dpcallback, &cbdata);
+	stop_computing = DPstreamParse(cbdata);
+	//cudaPrintfExtractDPE(dpcallback, cbdata);
 	//cudaPrintfDisplay(stdout, true);
 	cerr << getpid() << "\n";
     }
@@ -590,8 +529,6 @@ cout << i << ": " << l_Z[i] << " != " << l_z[i + t * WORDS] << "\n";
     free(l_z);
 
     pbak.restore();
-
-    return cbdata.expon;
 }
 
 #ifdef TEST_CUDA