Просмотр исходного кода

Added derandomization feature.

Setting the DERANDOMIZE macro will choose a deterministic problem (base and target), and Pollard's rho parameters (a, b, astep, and bstep). The distinguished points will also be arranged deterministically after each kernel launch.
Steven Engler 8 лет назад
Родитель
Сommit
f627a02297
5 измененных файлов с 80 добавлено и 11 удалено
  1. 2 1
      Makefile
  2. 15 2
      controller.cc
  3. 21 3
      dlrho.cc
  4. 25 0
      dpstream.cu
  5. 17 5
      worker.cc

+ 2 - 1
Makefile

@@ -24,9 +24,10 @@ MPI ?= /usr/local
 CXX ?= g++
 
 WORDS = 24
+MACROS=-DDERANDOMIZE -UVERBOSE
 
 CXXFLAGS=-g -Wall -O2
-CPPFLAGS=-DWORDS=$(WORDS) -I$(LIBEVENT)/include -I$(GMP)/include -UVERBOSE -std=c++11
+CPPFLAGS=-DWORDS=$(WORDS) -I$(LIBEVENT)/include -I$(GMP)/include -std=c++11 $(MACROS)
 NVCCOPTS=-g -arch sm_30 --ptxas-options=-v -O2
 NVCC=nvcc $(NVCCOPTS)
 

+ 15 - 2
controller.cc

@@ -539,9 +539,20 @@ static int generate_problem()
 
     // Generate a DLP mod rho (in the large odd-order subgroup)
     ZZ_p::init(ctrlstate.rho);
+    unsigned int attempt = 0;
     do {
-	ZZ_p base = power(random_ZZ_p(), 2);
-	ZZ_p target = power(random_ZZ_p(), 2);
+	ZZ_p base;
+	ZZ_p target;
+	{
+#ifdef DERANDOMIZE
+	    RandomStreamPush push_seed;
+	    // the seed will be reset to its original value
+	    // once we exit this scope
+	    SetSeed(ctrlstate.rho*ctrlstate.p.factor*ctrlstate.q.factor+attempt);
+#endif
+	    base = power(random_ZZ_p(), 2);
+	    target = power(random_ZZ_p(), 2);
+	}
 	gettimeofday(&ctrlstate.started_working, NULL);
 	ctrlstate.base = rep(base);
 	ctrlstate.target = rep(target);
@@ -551,6 +562,8 @@ static int generate_problem()
 	ctrlstate.subproblems_q = decomp(base, target, ctrlstate.q);
 	num_subproblems_p = ctrlstate.subproblems_p.size();
 	num_subproblems_q = ctrlstate.subproblems_q.size();
+
+	attempt++;
     } while (num_subproblems_p == 0 || num_subproblems_q == 0);
 
     ctrlstate.num_unsolved_subproblems =

+ 21 - 3
dlrho.cc

@@ -169,7 +169,15 @@ static int p_dl(const ZZ_p &target, const ZZ_p &base, ZZ &exp,
 	}
 	gettimeofday(&st, NULL);
 	unsigned int launch_count = 0;
-	cuda_dl(subgroup_base, subgroup_target, fvec[i], md, dpfreq, &cbdata, &launch_count);
+	{
+#ifdef DERANDOMIZE
+	    RandomStreamPush push_seed;
+	    // the seed will be reset to its original value
+	    // once we exit this scope
+	    SetSeed(rep(subgroup_base)*rep(subgroup_target)*fvec[i]*md);
+#endif
+	    cuda_dl(subgroup_base, subgroup_target, fvec[i], md, dpfreq, &cbdata, &launch_count);
+	}
 	ZZ subgroup_dl = cbdata.expon;
 	gettimeofday(&et, NULL);
 	unsigned long us_elapsed = (et.tv_sec-st.tv_sec)*1000000 +
@@ -307,8 +315,18 @@ int main(int argc, char **argv)
 
     // Generate a DLP mod rho (in the large odd-order subgroup)
     ZZ_p::init(rho);
-    ZZ_p base = power(random_ZZ_p(), 2);
-    ZZ_p target = power(random_ZZ_p(), 2);
+    ZZ_p base;
+    ZZ_p target;
+    {
+#ifdef DERANDOMIZE
+	RandomStreamPush push_seed;
+	// the seed will be reset to its original value
+	// once we exit this scope
+	SetSeed(rho*p*q);
+#endif
+	base = power(random_ZZ_p(), 2);
+	target = power(random_ZZ_p(), 2);
+    }
 
     cout << "base = " << base << "\n";
     cout << "target = " << target << "\n";

+ 25 - 0
dpstream.cu

@@ -22,6 +22,7 @@
 
 #include <sm_20_atomic_functions.h>
 #include <unistd.h>
+#include <stdlib.h>
 
 /* Routines to allow CUDA threads to write distinguished points into
    a buffer that will be read by the host */
@@ -83,6 +84,25 @@ __device__ inline unsigned int *DPstreamAlloc()
     return ourbuffer;
 }
 
+#ifdef DERANDOMIZE
+int DPrecordComparator(const void* p1, const void* p2)
+{
+    const unsigned int* p1_uint = (const unsigned int*)p1;
+    const unsigned int* p2_uint = (const unsigned int*)p2;
+
+    for (int i=0; i<DPrecordsize; i++) {
+	// compare each word
+	if (p1_uint[i] < p2_uint[i]) {
+	    return 1;
+	} else if (p1_uint[i] > p2_uint[i]) {
+	    return -1;
+	}
+    }
+
+    return 0;
+}
+#endif
+
 // dp points to an array of WORDS+7 unsigned ints:
 // - 1 word of threadID/blockID
 // - WORDS words of the dp value
@@ -104,6 +124,11 @@ bool DPstreamParse(void *data, unsigned int *num_dps)
 	cudaMemcpy(dpbuf, DPbuffer_device, bufsize * sizeof(unsigned int),
 	    cudaMemcpyDeviceToHost);
 
+#ifdef DERANDOMIZE
+	qsort(dpbuf, *num_dps, DPrecordsize*sizeof(unsigned int), DPrecordComparator);
+	// make the order of the DPs deterministic after each kernel launch
+#endif
+
 	for (unsigned int *dp=dpbuf; dp < dpbuf+bufsize; dp += DPrecordsize) {
 	    if (dpcallback(data, dp)) {
 		stop_computing = true;

+ 17 - 5
worker.cc

@@ -101,11 +101,23 @@ static void *worker_thread_start(void *data)
 	exit(1);
     }
 
-    cuda_dl(to_ZZ_p(wrkctrlstate.current_problem->base),
-	    to_ZZ_p(wrkctrlstate.current_problem->target),
-	    wrkctrlstate.current_problem->order,
-	    wrkctrlstate.current_problem->modulus,
-	    wrkctrlstate.current_problem->dpfreq, NULL, &wrkctrlstate.kernel_launch_count);
+    {
+#ifdef DERANDOMIZE
+	RandomStreamPush push_seed;
+	// the seed will be reset to its original value
+	// once we exit this scope
+	SetSeed(wrkctrlstate.current_problem->base*
+		wrkctrlstate.current_problem->target*
+		wrkctrlstate.current_problem->order*
+		wrkctrlstate.current_problem->modulus);
+#endif
+	cuda_dl(to_ZZ_p(wrkctrlstate.current_problem->base),
+	        to_ZZ_p(wrkctrlstate.current_problem->target),
+		wrkctrlstate.current_problem->order,
+		wrkctrlstate.current_problem->modulus,
+		wrkctrlstate.current_problem->dpfreq, NULL,
+		&wrkctrlstate.kernel_launch_count);
+    }
 
     return NULL;
 }