Let the caller of cuda_dl set the DP frequency

cudadl.h

@@ -25,12 +25,7 @@
 
 NTL_CLIENT
 
-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);
+		const ZZ &modulus, unsigned int dpfreq, void *data);
 
 #endif

dlrho.cc

@@ -154,8 +154,12 @@ static int p_dl(const ZZ_p &target, const ZZ_p &base, ZZ &exp,
 	ZZ md = ZZ_p::modulus();
 	CBData cbdata(subgroup_base, subgroup_target, fvec[i]);
 	struct timeval st, et;
+	unsigned int dpfreq = 4294967; // 2^32/1000
+	if (fvec[i] < 1000000) {
+	    dpfreq = 4294967295; // 2^32-1 : every point is a DP
+	}
 	gettimeofday(&st, NULL);
-	cuda_dl(subgroup_base, subgroup_target, fvec[i], md, &cbdata);
+	cuda_dl(subgroup_base, subgroup_target, fvec[i], md, dpfreq, &cbdata);
 	ZZ subgroup_dl = cbdata.expon;
 	gettimeofday(&et, NULL);
 	unsigned long us_elapsed = (et.tv_sec-st.tv_sec)*1000000 +

parrhoasm.cu

@@ -188,7 +188,8 @@ __device__ inline bool _gt(const unsigned int * x)
 #define threadDimx 32
 
 __global__ void cudaMulmod(GlobalThreadState *global_ts,
-	unsigned int order_0, unsigned int order_1, unsigned int order_2)
+	unsigned int order_0, unsigned int order_1, unsigned int order_2,
+	unsigned int dpfreq)
 // x <- x * y mod rho
 {
     CIOS_REG_DECLARE;
@@ -282,7 +283,7 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
 
 	// Check for a distinguished point
 	asm("mov.u32 %0, $xr0;" : "=r" (xlow));
-	if ((xlow & 0x000003ff) == 0 || !(order_2|order_1) && order_0 < (1<<20)) {
+	if (xlow <= dpfreq) {
 	    unsigned int *ourbuffer = DPstreamAlloc();
 	    if (ourbuffer) {
 		CIOS_WRITE_DP(ourbuffer, a_0, a_1, a_2, b_0, b_1, b_2);
@@ -307,7 +308,7 @@ int nthreads = 25600;
 int nblocks = 50;
 
 void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
-		const ZZ &modulus, void *cbdata)
+		const ZZ &modulus, unsigned int dpfreq, void *cbdata)
 {
     unsigned long long totmicros = 0;
     ZZ_pBak pbak;
@@ -468,7 +469,8 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 #endif
     {
 	cerr << getpid() << " Launch " << ++launchcount << "...\n";
-	cudaMulmod<<< nblocks, nthreads/nblocks /*tpb*/, 0 >>>(d_ts, order_0, order_1, order_2);
+	cudaMulmod<<< nblocks, nthreads/nblocks /*tpb*/, 0 >>>(d_ts,
+	    order_0, order_1, order_2, dpfreq);
 
 	cudaThreadSynchronize();
 	checkCUDAError("kernel launch");

worker.cc

@@ -85,7 +85,8 @@ static void *worker_thread_start(void *data)
     cuda_dl(to_ZZ_p(wrkctrlstate.current_problem->base),
 	    to_ZZ_p(wrkctrlstate.current_problem->target),
 	    wrkctrlstate.current_problem->order,
-	    wrkctrlstate.current_problem->modulus, NULL);
+	    wrkctrlstate.current_problem->modulus,
+	    wrkctrlstate.current_problem->dpfreq, NULL);
 
     return NULL;
 }