Merge pull request #1218 from HannoSpreeuw/Fix_1208

Fix 1208

Author: HannoSpreeuw

lib/sapporo_light/sapporo.cpp

@@ -270,3 +270,8 @@ int sapporo::get_ngb_list(int cluster_id,
     sort(nbl, nbl + min(nblen, maxlength));
     return overflow;
 }
+
+int sapporo::get_nj_max() const
+{
+    return nj_max;
+}

lib/sapporo_light/sapporo.h

@@ -103,6 +103,7 @@ extern "C"
 {
   int  get_device_count();
   cudaError_t host_evaluate_gravity(sapporo_multi_struct);
+  int g6_get_nj_max_();
 }
 
 class sapporo {
@@ -169,7 +170,7 @@ class sapporo {
 
     predict = false;
     nj_modified = 0;
-    
+
     device.address_j = NULL;
 
     device.t_j   = NULL;
@@ -182,13 +183,15 @@ class sapporo {
     device.acc_i = NULL;
     device.vel_i = NULL;
     device.jrk_i = NULL;
+
 
   };
   ~sapporo() {};
 
   int open(int cluster_id);
   int close(int cluster_id);
   int get_n_pipes();
+  int get_nj_max() const;
   int set_ti(int cluster_id, double ti);
 
   int set_j_particle(int cluster_id,

lib/sapporo_light/sapporoG6lib.cpp

@@ -49,6 +49,9 @@ extern "C" {
     return grav.calc_lasthalf2(*cluster_id, *nj, *ni,
 			       index, xi, vi, *eps2, h2, acc, jerk, pot, inn);
   }
+  int g6_get_nj_max_() {
+    return grav.get_nj_max();
+  }
 
   int g6_initialize_jp_buffer_(int* cluster_id, int* buf_size) {return 0;}
   int g6_flush_jp_buffer_(int* cluster_id) {return 0;}
@@ -80,4 +83,3 @@ extern "C" {
 
 
 }
-

lib/sapporo_light/send_fetch_data.cpp

@@ -65,6 +65,11 @@ void sapporo::send_j_particles_to_device(int ignore) {
     dev_struct &dev = device;
 
     int nj = address_j.size();
+
+    if (nj > nj_max) {
+      fprintf(stdout, "ERROR: nj = %d exceeds GPU capacity nj_max = %d\n", nj, nj_max);
+      abort;
+    }
 
     CUDA_SAFE_CALL(cudaMemcpy( dev.address_j,      &address_j[0], nj * sizeof(int),    cudaMemcpyHostToDevice));
     CUDA_SAFE_CALL(cudaMemcpy( &dev.t_j  [nj_max], &t_j[0],       nj * sizeof(DS2),    cudaMemcpyHostToDevice));

src/amuse_ph4/interface.cc

@@ -936,3 +936,9 @@ int get_id_of_updated_particle(int index, int * index_of_particle, int * status)
     *status = jd->UpdatedParticles[index].status;
     return 0;
 }
+
+int get_nj_max(int * value)
+{
+  *value = jd->get_nj_max();
+  return 0;
+}

src/amuse_ph4/interface.py

@@ -464,6 +464,13 @@ def recompute_timesteps():
         function.result_type = 'int32'
         return function
 
+    @legacy_function
+    def get_nj_max():
+        function = LegacyFunctionSpecification()
+        function.addParameter('nj_max', dtype='int32',
+                              direction=function.OUT)
+        function.result_type = 'int32'
+        return function
 
 class ph4(GravitationalDynamics,GravityFieldCode):
 

src/amuse_ph4/src/jdata.cc

@@ -34,6 +34,7 @@
 
 #ifdef GPU
 #include "grape.h"
+#include "sapporo.h"
 #endif
 
 // AMUSE STOPPING CONDITIONS SUPPORT
@@ -1153,3 +1154,12 @@ real jdata::get_tnext()
     int ndum;
     return sched->get_list(NULL, ndum);
 }
+
+int jdata::get_nj_max()
+{
+#ifdef GPU
+    return g6_get_nj_max_();
+#else
+    return std::numeric_limits<int>::max();
+#endif
+}

src/amuse_ph4/src/jdata.h

@@ -212,6 +212,7 @@ class jdata {
     void spec_output(const char *s = NULL);
     void to_com();
     real get_tnext();
+    int get_nj_max();
 
     // In gpu.cc:
 

src/tests/compile_tests/async_implementation/test_async.py

@@ -8,6 +8,7 @@
 import os
 from pathlib import Path
 import time
+import sys
 
 
 class ForTesting(InCodeComponentImplementation):
@@ -120,7 +121,9 @@ def test7(self):
             self.assertEqual(x[0].result(), x[1])
         instance1.stop()
         instance2.stop()
-        self.assertTrue(t2-t1 < 2.)
+        # Do not run this test on MacOS runners, since they seem very slow.
+        if sys.platform != "darwin":
+            self.assertTrue(t2-t1 < 2.)
 
     def test8(self):
         from threading import Thread