Can save distinguished points to files.

By setting the SAVE_DPS macro, you can save the distinguished points. The points discovered during each subproblem are saved in individual files (or multiple files in the case of multiple dpnodes). The file name format is 'dplist_{subproblemid}.out' for the dlrho version, or 'dplist_{dpnodeid}_{subproblemid}.out' for the controller/worker version. This can generate a large amount of data, so it's best to use it on small problems (B < 2^60).

When run with the DERANDOMIZE macro, it should produce the same points each time, except in the case of small subproblems where dpfreq is large and the device dp buffer is filled. In this case the points will not be the same, and this can be detected by the line "Warning: The device dp buffer was filled, so some points were not recorded." in the worker's stderr.

Makefile

@@ -24,7 +24,7 @@ MPI ?= /usr/local
 CXX ?= g++
 
 WORDS = 24
-MACROS=-DDERANDOMIZE -UVERBOSE
+MACROS=-DDERANDOMIZE -UVERBOSE -DSAVE_DPS
 
 CXXFLAGS=-g -Wall -O2
 CPPFLAGS=-DWORDS=$(WORDS) -I$(LIBEVENT)/include -I$(GMP)/include -std=c++11 $(MACROS)

controller.cc

@@ -485,7 +485,7 @@ static vector<SubproblemProgress> decomp(const ZZ_p &base, const ZZ_p &target,
 	}
 
 	++cursubproblemid;
-	ret.push_back(SubproblemProgress(cursubproblemid, rep(subgroup_base),
+	ret.push_back(SubproblemProgress(cursubproblemid-1, rep(subgroup_base),
 				    rep(subgroup_target), f.factor, order));
     }
 

controller_main.cc

@@ -44,6 +44,13 @@ static void boundcb(const char *boundaddr, unsigned short boundport)
 
 int main(int argc, char **argv)
 {
+#ifdef SAVE_DPS
+    std::cerr << "Note: Saving the distinguished points. Do not run huge problems or else it will use up all of your disk space.\n";
+#ifndef DERANDOMIZE
+    std::cerr << "Saving DPs without derandomization!\n";
+#endif
+#endif
+
     unsigned short bindport = 0;
     Worklist worklist;
 

cudadl.h

@@ -26,6 +26,7 @@
 NTL_CLIENT
 
 void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
-		const ZZ &modulus, unsigned int dpfreq, void *data, unsigned int* launch_count);
+		const ZZ &modulus, unsigned int dpfreq, void *data,
+		unsigned int* launch_count, bool* filled_dp_buffer);
 
 #endif

dlrho.cc

@@ -24,6 +24,7 @@
 #include <NTL/ZZ_p.h>
 
 #include <map>
+#include <sstream>
 
 #include <sys/types.h>
 #include <sys/wait.h>
@@ -38,6 +39,10 @@ NTL_CLIENT
 
 string output_prefix;
 
+#ifdef SAVE_DPS
+static ofstream dp_file_stream;
+#endif
+
 typedef map<std::string, pair<ZZ,ZZ> > DTable;
 
 struct CBData {
@@ -86,6 +91,15 @@ bool dpcallback(void *cbdata, unsigned int *dpwords)
     pair<DTable::iterator, bool> res;
     string x((const char *)(dpwords+1), WORDS*sizeof(unsigned int));
 
+#ifdef SAVE_DPS
+    if (!d->found_collision) {
+	ZZ zz_x;
+	ZZFromBytes(zz_x, (const unsigned char *)(dpwords+1),
+		    WORDS*sizeof(unsigned int));
+	dp_file_stream << zz_x << "\n";
+    }
+#endif
+
     res = d->dtable.insert(DTable::value_type(x, ab));
     if (!res.second) {
 	// Collision!
@@ -165,6 +179,12 @@ static int p_dl(const ZZ_p &target, const ZZ_p &base, ZZ &exp,
 	    }
 	}
 
+#ifdef SAVE_DPS
+	std::ostringstream oss;
+	oss << "dplist_" << i+initial_subproblem_id << ".out";
+	dp_file_stream.open(oss.str().c_str());
+#endif
+
 	// Now use your favourite method to get the DL of
 	// subgroup_target with base subgroup_base, knowing that it's in
 	// the range [0,fvec[i]).
@@ -172,8 +192,13 @@ 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;
+
+#ifdef SAVE_DPS
+	dp_file_stream << "Subproblem " << i+initial_subproblem_id << "\n";
+#endif
 	gettimeofday(&st, NULL);
 	unsigned int launch_count = 0;
+	bool filled_dp_buffer = false;
 	{
 #ifdef DERANDOMIZE
 	    RandomStreamPush push_seed;
@@ -181,7 +206,7 @@ static int p_dl(const ZZ_p &target, const ZZ_p &base, ZZ &exp,
 	    // 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);
+	    cuda_dl(subgroup_base, subgroup_target, fvec[i], md, dpfreq, &cbdata, &launch_count, &filled_dp_buffer);
 	}
 	ZZ subgroup_dl = cbdata.expon;
 	gettimeofday(&et, NULL);
@@ -199,6 +224,15 @@ static int p_dl(const ZZ_p &target, const ZZ_p &base, ZZ &exp,
 	// cout << "CRT\n";
 	// cout << "curexp = " << curexp << "\n";
 	// cout << "curmodulus = " << curmodulus << "\n\n";
+
+#ifdef SAVE_DPS
+	if (filled_dp_buffer) {
+	    AtomicWriter atomic_cout(cout);
+	    atomic_cout << "Warning: The device dp buffer was filled, so some points were not recorded.\n";
+	    atomic_cout << "These points will not be reproducible (subproblem " << i+initial_subproblem_id << ").\n";
+	}
+        dp_file_stream.close();
+#endif
     }
     // We'd like a non-negative answer back
     if (curexp >= 0) {
@@ -308,6 +342,13 @@ static int p_dl_fork_join(PDLHandle *handle, ZZ &exp)
 
 int main(int argc, char **argv)
 {
+#ifdef SAVE_DPS
+    cerr << "Note: Saving the distinguished points. Do not run huge problems or else it will use up all of your disk space.\n";
+#ifndef DERANDOMIZE
+    cerr << "Saving DPs without derandomization!\n";
+#endif
+#endif
+
     if (argc != 2) {
 	cerr << "Usage: " << argv[0] << " total_mem_GB";
 	exit(1);

dpnode.cc

@@ -18,12 +18,19 @@ extern "C" {
 #include <string.h>
 #include <errno.h>
 
+#include <sstream>
+#include <fstream>
+
 #include "evutils.h"
 #include "subproblem.h"
 #include "dpnode.h"
 
 // #undef VERBOSE
 
+#ifdef SAVE_DPS
+static ofstream dp_file_stream;
+#endif
+
 typedef map<std::string, pair<ZZ,ZZ> > DTable;
 
 typedef enum {
@@ -139,6 +146,14 @@ static void dpnode_reader(struct bufferevent *bev, void *ctx)
 		3*sizeof(unsigned int));
 	string x((const char *)(dp), WORDS*sizeof(unsigned int));
 	pair<ZZ,ZZ> ab(zz_a,zz_b);
+
+#ifdef SAVE_DPS
+	ZZ zz_x;
+	ZZFromBytes(zz_x, (const unsigned char *)(dp),
+		    WORDS*sizeof(unsigned int));
+	dp_file_stream << zz_x << "\n";
+#endif
+
 	pair<DTable::iterator, bool> res =
 		dpctrlstate.table.insert(DTable::value_type(x,ab));
 #ifdef LOG_MEM
@@ -232,6 +247,11 @@ static void stop_problem(void)
 #ifdef VERBOSE
     cerr << "Stopping problem\n";
 #endif
+
+#ifdef SAVE_DPS
+    dp_file_stream.close();
+#endif
+
     if (dpctrlstate.current_problem) {
 	delete dpctrlstate.current_problem;
 	dpctrlstate.current_problem = NULL;
@@ -274,6 +294,13 @@ static void start_problem(struct bufferevent *bev,
     dpctrlstate.current_problem->dump(cerr);
 #endif
 
+#ifdef SAVE_DPS
+    std::ostringstream oss;
+    oss << "dplist_" << dpctrlstate.id << "_" << dpctrlstate.current_problem->problemid << ".out";
+    dp_file_stream.open(oss.str().c_str());
+    dp_file_stream << "Subproblem " << dpctrlstate.current_problem->problemid << "\n";
+#endif
+
     // Create the DPNode server socket
     dpctrlstate.listener = dpnode_create(bufferevent_get_base(bev),
 	    &myip, &myport);

dpstream.cu

@@ -111,12 +111,16 @@ int DPrecordComparator(const void* p1, const void* p2)
 extern bool dpcallback(void *data, unsigned int *dpwords);
 
 // Return true if we should stop computing
-bool DPstreamParse(void *data, unsigned int *num_dps)
+bool DPstreamParse(void *data, unsigned int *num_dps, bool *filled_dp_buffer)
 {
     // Get the number of words in the buffer
     unsigned int bufsize;
     cudaMemcpyFromSymbol(&bufsize, DPbuffertail, sizeof(unsigned int));
+    unsigned int max_bufsize;
+    cudaMemcpyFromSymbol(&max_bufsize, DPbuffersize, sizeof(unsigned int));
+
     *num_dps = bufsize/DPrecordsize;
+    *filled_dp_buffer = (bufsize == max_bufsize);
 
     unsigned int *dpbuf = (unsigned int*)calloc(bufsize, sizeof(unsigned int));
     bool stop_computing = false;

parrhoasm.cu

@@ -327,7 +327,7 @@ int nblocks = 50;
 
 void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 		const ZZ &modulus, unsigned int dpfreq, void *cbdata,
-		unsigned int* launch_count)
+		unsigned int* launch_count, bool* filled_dp_buffer)
 {
     unsigned long long totmicros = 0;
     ZZ_pBak pbak;
@@ -479,7 +479,9 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 
 	cudaThreadSynchronize();
 	checkCUDAError("thread sync");
-	stop_computing = DPstreamParse(cbdata, &num_dps);
+	bool single_launch_filled_dp_buffer = false;
+	stop_computing = DPstreamParse(cbdata, &num_dps, &single_launch_filled_dp_buffer);
+	*filled_dp_buffer = (*filled_dp_buffer || single_launch_filled_dp_buffer);
 #ifdef VERBOSE
 	AtomicWriter(cerr) << getpid() << " Finished " << local_launchcount << " (" << num_dps << " DPs)\n";
 #endif

worker.cc

@@ -101,6 +101,8 @@ static void *worker_thread_start(void *data)
 	exit(1);
     }
 
+    bool filled_dp_buffer = false;
+
     {
 #ifdef DERANDOMIZE
 	RandomStreamPush push_seed;
@@ -116,8 +118,15 @@ static void *worker_thread_start(void *data)
 		wrkctrlstate.current_problem->order,
 		wrkctrlstate.current_problem->modulus,
 		wrkctrlstate.current_problem->dpfreq, NULL,
-		&wrkctrlstate.kernel_launch_count);
+		&wrkctrlstate.kernel_launch_count, &filled_dp_buffer);
+    }
+
+#ifdef SAVE_DPS
+    if (filled_dp_buffer) {
+	cerr << "Warning: The device dp buffer was filled, so some points were not recorded.\n";
+	cerr << "These points will not be reproducible.\n";
     }
+#endif
 
     return NULL;
 }