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

Make it compile and run on a machine with cuda-11.5 and two A100 GPUs

I have no idea whether it's tuned optimally for this configuration
Ian Goldberg 2 месяцев назад
Родитель
Сommit
22fe3a92d9
4 измененных файлов с 8 добавлено и 8 удалено
  1. 1 1
      Makefile
  2. 2 2
      mpi.cc
  3. 4 4
      parrhoasm.cu
  4. 1 1
      worker.cc

+ 1 - 1
Makefile

@@ -17,7 +17,7 @@
 
 
 -include Makefile.local
 -include Makefile.local
 
 
-CUDA_ARCH ?= sm_30
+CUDA_ARCH ?= sm_80
 
 
 GMP ?= /usr/local
 GMP ?= /usr/local
 NTL ?=
 NTL ?=

+ 2 - 2
mpi.cc

@@ -22,12 +22,12 @@ static int controllerfds[2];
 
 
 // Close fds 3 and up, except for the one given (pass -1 to close them
 // Close fds 3 and up, except for the one given (pass -1 to close them
 // all)
 // all)
-static void close_highfds_except(int exceptfd)
+static void close_highfds_except(rlim_t exceptfd)
 {
 {
     // Find the max fd number
     // Find the max fd number
     struct rlimit limit;
     struct rlimit limit;
     getrlimit(RLIMIT_NOFILE, &limit);
     getrlimit(RLIMIT_NOFILE, &limit);
-    for (int fd = 3; fd < limit.rlim_cur; ++fd) {
+    for (rlim_t fd = 3; fd < limit.rlim_cur; ++fd) {
 	if (fd != exceptfd) {
 	if (fd != exceptfd) {
 	    // There's no ill effect from closing a non-open fd, so just
 	    // There's no ill effect from closing a non-open fd, so just
 	    // do it
 	    // do it

+ 4 - 4
parrhoasm.cu

@@ -189,7 +189,7 @@ __device__ inline bool _gt(const unsigned int * x)
 }
 }
 #endif
 #endif
 
 
-#define nmult 1000
+#define nmult 750
 #define nlaunch 2
 #define nlaunch 2
 
 
 #define threadDimx 32
 #define threadDimx 32
@@ -488,8 +488,8 @@ void cuda_dl(const ZZ_p &base, const ZZ_p &target, const ZZ &order,
 	    order_0, order_1, order_2, dpfreq);
 	    order_0, order_1, order_2, dpfreq);
 	checkCUDAError("kernel launch");
 	checkCUDAError("kernel launch");
 
 
-	cudaThreadSynchronize();
-	checkCUDAError("thread sync");
+	cudaDeviceSynchronize();
+	checkCUDAError("device sync");
 	bool single_launch_filled_dp_buffer = false;
 	bool single_launch_filled_dp_buffer = false;
 	stop_computing = DPstreamParse(cbdata, &num_dps, &single_launch_filled_dp_buffer);
 	stop_computing = DPstreamParse(cbdata, &num_dps, &single_launch_filled_dp_buffer);
 	*filled_dp_buffer = (*filled_dp_buffer || single_launch_filled_dp_buffer);
 	*filled_dp_buffer = (*filled_dp_buffer || single_launch_filled_dp_buffer);
@@ -581,7 +581,7 @@ cout << i << ": " << l_Z[i] << " != " << l_z[i + t * WORDS] << "\n";
 #endif
 #endif
 
 
     char timestamp[20];
     char timestamp[20];
-    sprintf(timestamp, "%d.%06d", et.tv_sec, et.tv_usec);
+    sprintf(timestamp, "%ld.%06ld", et.tv_sec, et.tv_usec);
 
 
     unsigned long long totnanos = totmicros * 1000;
     unsigned long long totnanos = totmicros * 1000;
     AtomicWriter(cout) << timestamp << ":" << output_prefix << ": " << totmicros << " us / " << nthreads << " = " << totmicros / nthreads << " us / " << (nmult*local_launchcount) << " = " << totnanos / ((unsigned long long)nthreads * nmult * local_launchcount) << " ns\n";
     AtomicWriter(cout) << timestamp << ":" << output_prefix << ": " << totmicros << " us / " << nthreads << " = " << totmicros / nthreads << " us / " << (nmult*local_launchcount) << " = " << totnanos / ((unsigned long long)nthreads * nmult * local_launchcount) << " ns\n";

+ 1 - 1
worker.cc

@@ -191,7 +191,7 @@ static void start_working(void)
 	wrkctrlstate.worker_thread = WT_NOT_RUNNING;
 	wrkctrlstate.worker_thread = WT_NOT_RUNNING;
 	cerr << "Could not start worker thread\n";
 	cerr << "Could not start worker thread\n";
     }
     }
-    char thread_name[16];
+    char thread_name[35];
     snprintf(thread_name, sizeof(thread_name), "prob:%2d wrkr:%2d", wrkctrlstate.current_problem->problemid, wrkctrlstate.worker_id);
     snprintf(thread_name, sizeof(thread_name), "prob:%2d wrkr:%2d", wrkctrlstate.current_problem->problemid, wrkctrlstate.worker_id);
     pthread_setname_np(wrkctrlstate.worker_thread, thread_name);
     pthread_setname_np(wrkctrlstate.worker_thread, thread_name);
 }
 }