Merge branch 'master' of github.com:igfuw/libcloudphxx into kida-1d

pdziekan · pdziekan · commit 8cf24fa4d4af · 2020-09-09T15:03:02.000+02:00
diff --git a/.travis.yml b/.travis.yml
@@ -68,7 +68,7 @@ aliases:
 jobs:
     include:
       - <<: *precache_tests_mpi
-        env: TEST_SUITE=MPI_tests MPI=mvapich2
+        env: TEST_SUITE=tests MPI=mvapich2
         compiler: gcc
 #     clang MPI_tests disabled due to issues with manually built boost python libstdc++ / libc++ issues
 #      - <<: *precache_default
@@ -88,7 +88,7 @@ jobs:
           env: TEST_SUITE=cuda MPI=none
 
         - compiler: clang
-          env: TEST_SUITE=MPI_tests MPI=none
+          env: TEST_SUITE=tests MPI=mvapich2
 
         - compiler: clang
           env: TEST_SUITE=UWLCM_MT MPI=mvapich2
@@ -180,7 +180,7 @@ install:
     - if [[ $MPI == 'none' && ($TEST_SUITE != 'tests' || $TEST_SUITE != 'KidA') ]]; then sudo ln -s /usr/lib/x86_64-linux-gnu/libboost_python-py35.so /usr/lib/x86_64-linux-gnu/libboost_python3.so; fi # different naming conventions for boost python with python 3
 
 # thrust
-    - git clone --depth=1 git://github.com/thrust/thrust.git;
+    - git clone --depth=1 git://github.com/thrust/thrust.git --branch 1.9.10-1;
     - sudo ln -s `pwd`/thrust/thrust /usr/local/include/thrust;
     - if [[ $TEST_SUITE == 'cuda' ]];  then  sudo ln -s `pwd`/thrust/thrust /usr/local/cuda/include/thrust; fi
  
@@ -202,6 +202,4 @@ script:
     - if [[ $TEST_SUITE == 'UWLCM_unit_smg' ]]; then . ./.travis_scripts/UWLCM.sh unit_smg; fi
     # KiD-A 1D
     - if [[ $TEST_SUITE == 'KidA' ]]; then . ./.travis_scripts/KidA.sh; fi
-    # unit tests and drops.py with MPI (drop drops here?)
-    - if [[ $TEST_SUITE == 'MPI_tests' ]]; then . ./.travis_scripts/tests.sh; fi
 
diff --git a/.travis_scripts/deps_install/boost.sh b/.travis_scripts/deps_install/boost.sh
@@ -19,7 +19,7 @@ DEPS_DIR="${TRAVIS_BUILD_DIR}/deps"
       cat $HOME/user-config.jam
       if [[ $COMPILER == 'g++' ]]; then
         ./bootstrap.sh --prefix=${DEPS_DIR}/boost/ --with-libraries=chrono,regex,serialization,mpi,atomic,thread,date_time,system,iostreams,timer,filesystem,program_options,python
-        ./b2 -d0 install
+        travis_wait 30 ./b2 -d0 install
       fi
       if [[ $COMPILER == 'clang++' ]]; then 
         #clang installation taken from https://gist.github.com/jimporter/10442880
diff --git a/include/libcloudph++/common/output.hpp b/include/libcloudph++/common/output.hpp
@@ -16,7 +16,8 @@ namespace libcloudphxx
       outS_VI = chem::S_VI,                        
       outH = chem::H,                          
       outliq_vol,
-      outdry_vol
+      outdry_vol,
+      outprtcl_num
     };
 
     const std::map<output_t, std::string> output_names
@@ -30,7 +31,8 @@ namespace libcloudphxx
       {outS_VI   , "S_VI"},                        
       {outH      , "H"},                          
       {outliq_vol, "liquid_volume"},
-      {outdry_vol, "dry_volume"}
+      {outdry_vol, "dry_volume"},
+      {outprtcl_num, "particle_number"}
     };
   };
 };
diff --git a/models/kinematic_2D/tests/paper_GMD_2015/fig_a/CMakeLists.txt b/models/kinematic_2D/tests/paper_GMD_2015/fig_a/CMakeLists.txt
@@ -59,7 +59,7 @@ add_test(travis_2D_kin_cloud_dims_blk bash -c "
 add_test(travis_2D_kin_cloud_diff_blk_1m bash -c "
   for dir in travis_out_blk_1m; do
     echo 'comparing const.h5'                                                                                                &&
-    h5diff --delta=2e-7  -v2 $dir/const.h5           ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/const.h5                       &&
+    h5diff --delta=3e-7  -v2 $dir/const.h5           ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/const.h5                       &&
     echo 'comparing timestep0000000000.h5'                                                                                   &&
     h5diff --delta=2e-5  -v2 $dir/timestep0000000000.h5 ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/timestep0000000000.h5  /rv  &&
     h5diff --delta=2e-5  -v2 $dir/timestep0000000000.h5 ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/timestep0000000000.h5  /rr  &&
@@ -76,7 +76,7 @@ add_test(travis_2D_kin_cloud_diff_blk_1m bash -c "
 add_test(travis_2D_kin_cloud_diff_blk_2m bash -c "
   for dir in travis_out_blk_2m; do
     echo 'comparing const.h5'
-    h5diff --delta=2e-7     -v2 $dir/const.h5              ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/const.h5                    &&
+    h5diff --delta=3e-7     -v2 $dir/const.h5              ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/const.h5                    &&
     echo 'comparing timestep0000000000.h5'                                                                                     &&
     h5diff --relative=1e-9  -v2 $dir/timestep0000000000.h5 ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/timestep0000000000.h5  /rv  &&
     h5diff --relative=1e-9  -v2 $dir/timestep0000000000.h5 ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/timestep0000000000.h5  /rr  &&
@@ -106,14 +106,14 @@ add_test(travis_2D_kin_cloud_dims_lgrngn bash -c "
 add_test(travis_2D_kin_cloud_diff_lgrngn bash -c "
   for dir in travis_out_lgrngn; do
     echo 'comparing const.h5'
-    h5diff --delta=1e-6     -v2 $dir/const.h5                ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_const.h5                       &&
+    h5diff --delta=1e-6      -v2 $dir/const.h5                ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_const.h5                       &&
     echo 'comparing timestep0000000000.h5'                                                                                                  &&
-    h5diff --relative=1e-9  -v2 $dir/timestep0000000000.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000000.h5 /th      &&
-    h5diff --relative=1e-9  -v2 $dir/timestep0000000000.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000000.h5 /rv      &&
-    h5diff --relative=1e-9  -v2 $dir/timestep0000000000.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000000.h5 /sd_conc &&
+    h5diff --relative=1e-9   -v2 $dir/timestep0000000000.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000000.h5 /th      &&
+    h5diff --relative=1e-9   -v2 $dir/timestep0000000000.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000000.h5 /rv      &&
+    h5diff --relative=1e-9   -v2 $dir/timestep0000000000.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000000.h5 /sd_conc &&
     echo 'comparing timestep0000000020.h5'                                                                                                  &&
-    h5diff --relative=0.001 -v2 $dir/timestep0000000020.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000020.h5 /th      &&
-    h5diff --relative=0.002 -v2 $dir/timestep0000000020.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000020.h5 /rv      &&
-    h5diff --delta=1        -v2 $dir/timestep0000000020.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000020.h5 /sd_conc || exit 1;
+    h5diff --relative=0.001  -v2 $dir/timestep0000000020.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000020.h5 /th      &&
+    h5diff --relative=0.0021 -v2 $dir/timestep0000000020.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000020.h5 /rv      &&
+    h5diff --delta=1         -v2 $dir/timestep0000000020.h5   ${CMAKE_CURRENT_SOURCE_DIR}/refdata/$dir/travis_timestep0000000020.h5 /sd_conc || exit 1;
   done  
 ")
diff --git a/src/detail/gpu_assert.hpp b/src/detail/gpu_assert.hpp
@@ -16,7 +16,7 @@ namespace libcloudphxx
   {
     namespace detail
     {   
-      void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
+      inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
       {   
          if (code != cudaSuccess) 
          {
@@ -26,10 +26,10 @@ namespace libcloudphxx
       }   
 
       // max(1, n)
-      int m1(int n) { return n == 0 ? 1 : n; }
+      inline int m1(int n) { return n == 0 ? 1 : n; }
 
       // run a function on a specific gpu
-      void set_device_and_run(int id, std::function<void()> fun)
+      inline void set_device_and_run(int id, std::function<void()> fun)
       {
         gpuErrchk(cudaSetDevice(id));
         fun();
diff --git a/src/detail/urand.hpp b/src/detail/urand.hpp
@@ -120,8 +120,9 @@ namespace libcloudphxx
 	)
 	{
 	  int status = curandGenerateUniform(gen, thrust::raw_pointer_cast(v.data()), n);
-	  assert(status == CURAND_STATUS_SUCCESS /* && "curandGenerateUniform failed"*/);
+          assert(status == CURAND_STATUS_SUCCESS /* && "curandGenerateUniform failed"*/);
           _unused(status);
+
 	}
 
 	void generate_n(
diff --git a/src/impl/particles_impl.ipp b/src/impl/particles_impl.ipp
@@ -354,7 +354,7 @@ namespace libcloudphxx
 
         // if using nvcc, put increase_sstp_coal flag in host memory, but with direct access from device code
 #if defined(__NVCC__)
-        cudaMallocHost(&increase_sstp_coal, sizeof(bool));
+        gpuErrchk(cudaMallocHost(&increase_sstp_coal, sizeof(bool)));
 #else
         increase_sstp_coal = new bool();
 #endif
diff --git a/src/impl/particles_impl_bcnd.ipp b/src/impl/particles_impl_bcnd.ipp
@@ -68,7 +68,7 @@ namespace libcloudphxx
         BOOST_GPU_ENABLED
         real_t operator()(real_t x)
         {
-          return a + fmod((x-a) + (b-a), b-a); // this should call CUDA's fmod!
+          return a + fmod((x-a) + 10 * (b-a), b-a); // assuming that particles dont move more than 10 * domain size; this should call CUDA's fmod!
         }
       };
     };
@@ -242,6 +242,13 @@ namespace libcloudphxx
                     thrust::plus<real_t>()
                   );
 
+                // add total number of particles that fell out in this step
+                output_puddle[common::outprtcl_num] += 
+                  thrust::reduce(
+                    n_filtered.begin(),            // input start
+                    n_filtered.begin() + n_part    // input end
+                  );
+
                 if(opts_init.chem_switch)
                 {
                   for (int i = 0; i < chem_all; ++i)
diff --git a/src/impl/particles_impl_chem_henry.ipp b/src/impl/particles_impl_chem_henry.ipp
@@ -352,7 +352,7 @@ namespace libcloudphxx
         thrust::pair<
           typename thrust_device::vector<thrust_size_t>::iterator,
           typename thrust_device::vector<real_t>::iterator
-        > np =
+        > it_pair =
         thrust::reduce_by_key(
           sorted_ijk.begin(), sorted_ijk.end(),
           thrust::transform_iterator<             // input - values
@@ -369,7 +369,7 @@ namespace libcloudphxx
           count_ijk.begin(),
           mass_new.begin()
         );
-        count_n = np.first - count_ijk.begin();
+        count_n = it_pair.first - count_ijk.begin();
         assert(count_n > 0 && count_n <= n_cell);
 
         // apply the change to the mixing ratios of trace gases
diff --git a/src/impl/particles_impl_hskpng_count.ipp b/src/impl/particles_impl_hskpng_count.ipp
@@ -21,13 +21,13 @@ namespace libcloudphxx
       thrust::pair<
         thrust_device::vector<thrust_size_t>::iterator,
         thrust_device::vector<n_t>::iterator
-      > n = thrust::reduce_by_key(
+      > it_pair = thrust::reduce_by_key(
         sorted_ijk.begin(), sorted_ijk.end(),   // input - keys
         thrust::make_constant_iterator(n_t(1)), // input - values
         count_ijk.begin(),                      // output - keys
         count_num.begin()                       // output - values
       );
-      count_n = n.first - count_ijk.begin();
+      count_n = it_pair.first - count_ijk.begin();
       assert(count_n <= n_cell);
     }   
   };  
diff --git a/src/impl/particles_impl_mass_dens.ipp b/src/impl/particles_impl_mass_dens.ipp
@@ -62,7 +62,7 @@ namespace libcloudphxx
       thrust::pair<
         thrust_device::vector<thrust_size_t>::iterator,
         typename thrust_device::vector<real_t>::iterator
-      > n = thrust::reduce_by_key(
+      > it_pair = thrust::reduce_by_key(
         // input - keys
         sorted_ijk.begin(), sorted_ijk.end(),
         // input - values
@@ -79,6 +79,10 @@ namespace libcloudphxx
         // output - values
         count_mom.begin()
       );
+
+      count_n = it_pair.first - count_ijk.begin();
+      assert(count_n > 0 && count_n <= n_cell);
+
 #if !defined(__NVCC__)
       using std::sqrt;
 #endif
@@ -89,12 +93,8 @@ namespace libcloudphxx
 #else
         sqrt(CUDART_PI / 2.);
 #endif
-
       namespace arg = thrust::placeholders;
 
-      count_n = n.first - count_ijk.begin();
-      assert(count_n > 0 && count_n <= n_cell);
-
       //multiply by prefactor and divide by dv
       thrust::transform(
         count_mom.begin(), count_mom.begin() + count_n,     // input - first arg
diff --git a/src/impl/particles_impl_moms.ipp b/src/impl/particles_impl_moms.ipp
@@ -189,7 +189,7 @@ namespace libcloudphxx
       thrust::pair<
         thrust_device::vector<thrust_size_t>::iterator,
         typename thrust_device::vector<real_t>::iterator
-      > n = thrust::reduce_by_key(
+      > it_pair = thrust::reduce_by_key(
         // input - keys
         sorted_ijk.begin(), sorted_ijk.end(),  
         // input - values
@@ -206,7 +206,7 @@ namespace libcloudphxx
         count_mom.begin()
       );  
 
-      count_n = n.first - count_ijk.begin();
+      count_n = it_pair.first - count_ijk.begin();
 #if !defined(NDEBUG)
       {
         int nan_count = thrust::transform_reduce(count_mom.begin(), count_mom.begin() + count_n, isnaninf(), 0, thrust::plus<bool>());
diff --git a/src/impl/particles_impl_mpi_exchange.ipp b/src/impl/particles_impl_mpi_exchange.ipp
@@ -66,7 +66,7 @@ namespace libcloudphxx
         // At least pre Thrust 1.9.4. 1.9.4 made all calls (that are not explicitly asynchronous) synchronous, see https://github.com/thrust/thrust/blob/master/doc/changelog.md#thrust-194-cuda-101
         // TODO: Test with Thurst >= 1.9.4 if this sync is still needed
 #if defined(__NVCC__)
-        cudaDeviceSynchronize();
+        gpuErrchk(cudaDeviceSynchronize());
 #endif
 
 //        std::cerr << "mpi exchange: sending n left, lft_count = " << lft_count << " sum of out_n_bfr = " << thrust::reduce(out_n_bfr.begin(), out_n_bfr.begin() + lft_count) << std::endl;
@@ -131,7 +131,7 @@ namespace libcloudphxx
 
 //      check comment for the previous device sync
 #if defined(__NVCC__)
-        cudaDeviceSynchronize();
+        gpuErrchk(cudaDeviceSynchronize());
 #endif
 
 //        std::cerr << "mpi exchange: sending real left, lft_count = " << lft_count << " sum of out_real_bfr = " << thrust::reduce(out_real_bfr.begin(), out_real_bfr.begin() + lft_count * distmem_real_vctrs.size()) << std::endl;
diff --git a/src/impl/particles_impl_src_dry_distros.ipp b/src/impl/particles_impl_src_dry_distros.ipp
@@ -232,13 +232,13 @@ namespace libcloudphxx
             thrust::pair<
               thrust_device::vector<thrust_size_t>::iterator,
               typename thrust_device::vector<thrust_size_t>::iterator
-            > np =  thrust::reduce_by_key(
+            > it_pair =  thrust::reduce_by_key(
               sorted_ijk.begin(), sorted_ijk.begin() + count_bins, 
               thrust::make_constant_iterator<thrust_size_t>(1), 
               sorted_ijk.begin(), 
               tmp_bin_no.begin()
             );
-            count_bins = np.first - sorted_ijk.begin(); // now it counts no of cells that have any bins matched
+            count_bins = it_pair.first - sorted_ijk.begin(); // now it counts no of cells that have any bins matched
           }
 
           // set count_num to the number of SDs matched per cell
diff --git a/src/impl/particles_impl_unpack.ipp b/src/impl/particles_impl_unpack.ipp
@@ -31,6 +31,20 @@ namespace libcloudphxx
         thrust::copy( in_real_bfr.begin() + distance * n_copied, in_real_bfr.begin() + (distance+1) * n_copied, (*it)->begin() + n_part_old);
         it++;
       }
+#if !defined(NDEBUG)
+      auto min_it = thrust::min_element(x.begin() + n_part_old, x.end());
+      if(*min_it < opts_init.x0)
+      {
+        std::cerr << "x (" << *min_it << ")  < opts_init.x0 (" << opts_init.x0 << ") after unpacking, potentially SD moved by more than one process/GPU domain size" << std::endl;
+        assert(0);
+      }
+      auto max_it = thrust::max_element(x.begin() + n_part_old, x.end());
+      if(*max_it >= opts_init.x1)
+      {
+        std::cerr << "x (" << *max_it << ")  >= opts_init.x1 (" << opts_init.x1 << ") after unpacking, potentially SD moved by more than one process/GPU domain size" << std::endl;
+        assert(0);
+      }
+#endif
     }
 
     template <typename real_t, backend_t device>
diff --git a/src/impl/particles_impl_update_th_rv.ipp b/src/impl/particles_impl_update_th_rv.ipp
@@ -109,13 +109,13 @@ namespace libcloudphxx
       thrust::pair<
         thrust_device::vector<thrust_size_t>::iterator,
         typename thrust_device::vector<real_t>::iterator
-      > n = thrust::reduce_by_key(
+      > it_pair = thrust::reduce_by_key(
         sorted_ijk.begin(), sorted_ijk.end(),
         thrust::make_permutation_iterator(pdstate.begin(), sorted_id.begin()),
         count_ijk.begin(),
         count_mom.begin()
       );
-      count_n = n.first - count_ijk.begin();
+      count_n = it_pair.first - count_ijk.begin();
 
       // add this sum to dstate
       thrust::transform(
diff --git a/src/impl_multi_gpu/particles_multi_gpu_impl.ipp b/src/impl_multi_gpu/particles_multi_gpu_impl.ipp
@@ -69,6 +69,12 @@ namespace libcloudphxx
         // copy actual dev_count to glob_opts_init
         glob_opts_init.dev_count = dev_count;
      
+        if(glob_opts_init.dev_id >= 0)
+        {
+          std::cout << "Libcloudph++ warning: opts_init.dev_id is not compatible with the multi_CUDA backend, ignoring it's value." << std::endl;
+          glob_opts_init.dev_id = -1;
+        }
+     
         // check if all GPUs support UVA
         // TODO: other checks?, see CUDA samples 0_simple/simpleP2P
         for (int i = 0; i < dev_count; ++i)
@@ -170,6 +176,7 @@ namespace libcloudphxx
           }
           // store dev_count in the thread; regular ctor zeroes it
           particles[dev_id]->pimpl->opts_init.dev_count = dev_count;
+          gpuErrchk(cudaDeviceSynchronize());
         }
       }
 
diff --git a/src/impl_multi_gpu/particles_multi_gpu_impl_step_async_and_copy.ipp b/src/impl_multi_gpu/particles_multi_gpu_impl_step_async_and_copy.ipp
@@ -39,7 +39,7 @@ namespace libcloudphxx
       particles[dev_id]->step_async(opts);
 
       // --- copy advected SDs to other devices on the same node ---
-      if(opts.adve && glob_opts_init.dev_count>1)
+      if((opts.adve || opts.turb_adve) && glob_opts_init.dev_count>1)
       {
         namespace arg = thrust::placeholders;
         typedef unsigned long long n_t; // TODO: same typedef is in impl struct !! particles::impl::n_t ? 
diff --git a/src/lib_cuda.cu b/src/lib_cuda.cu
@@ -3,6 +3,7 @@
 #include <thrust/system/cuda/vector.h>
 namespace thrust_device = ::thrust::cuda;
 
+#include "detail/gpu_assert.hpp"
 #include "particles.tpp"
 
 namespace libcloudphxx
diff --git a/src/lib_multicuda.cu b/src/lib_multicuda.cu
@@ -3,6 +3,7 @@
 #include <thrust/system/cuda/vector.h>
 namespace thrust_device = ::thrust::cuda;
 
+#include "detail/gpu_assert.hpp"
 #include "particles.tpp"
 #include "particles_multi_gpu.tpp"
 
diff --git a/src/particles_ctor.ipp b/src/particles_ctor.ipp
@@ -64,7 +64,7 @@ namespace libcloudphxx
       // use the desired GPU card, TODO: remove it? can be done using CUDA_VISIBLE_DEVICES
 #if defined(__NVCC__)
       if(opts_init.dev_id >= 0)
-        cudaSetDevice(opts_init.dev_id);
+        gpuErrchk(cudaSetDevice(opts_init.dev_id));
 #endif
       if(opts_init.dev_count < 2) // no distmem
         n_x_tot = opts_init.nx;
@@ -75,7 +75,7 @@ namespace libcloudphxx
       pimpl->sanity_checks();
 
       // init output map to 0
-      for(int i=0; i < chem_all+2; ++i)
+      for(int i=0; i < common::output_names.size(); ++i)
         pimpl->output_puddle[static_cast<common::output_t>(i)] = 0.;
     }
 
diff --git a/src/particles_multi_gpu.tpp b/src/particles_multi_gpu.tpp
@@ -6,7 +6,6 @@
   * @brief Thrust-based CPU/GPU particle-tracking logic for Lagrangian microphysics
   */
 
-#include "detail/gpu_assert.hpp"
 #include "detail/distmem_opts.hpp"
 #include "impl_multi_gpu/particles_multi_gpu_impl.ipp"
 #include "impl_multi_gpu/particles_multi_gpu_impl_step_async_and_copy.ipp"
diff --git a/src/particles_multi_gpu_diag.ipp b/src/particles_multi_gpu_diag.ipp
diff --git a/src/particles_step.ipp b/src/particles_step.ipp
diff --git a/tests/python/physics/puddle.py b/tests/python/physics/puddle.py

Original file line number	Diff line number	Diff line change
`@@ -16,7 +16,7 @@ namespace libcloudphxx`
`16`	`16`	`{`
`17`	`17`	`namespace detail`
`18`	`18`	`{`
`19`		`- void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)`
	`19`	`+ inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)`
`20`	`20`	`{`
`21`	`21`	`if (code != cudaSuccess)`
`22`	`22`	`{`
`@@ -26,10 +26,10 @@ namespace libcloudphxx`
`26`	`26`	`}`
`27`	`27`
`28`	`28`	`// max(1, n)`
`29`		`- int m1(int n) { return n == 0 ? 1 : n; }`
	`29`	`+ inline int m1(int n) { return n == 0 ? 1 : n; }`
`30`	`30`
`31`	`31`	`// run a function on a specific gpu`
`32`		`- void set_device_and_run(int id, std::function<void()> fun)`
	`32`	`+ inline void set_device_and_run(int id, std::function<void()> fun)`
`33`	`33`	`{`
`34`	`34`	`gpuErrchk(cudaSetDevice(id));`
`35`	`35`	`fun();`
Original file line number	Diff line number	Diff line change
`@@ -120,8 +120,9 @@ namespace libcloudphxx`
`120`	`120`	`)`
`121`	`121`	`{`
`122`	`122`	`int status = curandGenerateUniform(gen, thrust::raw_pointer_cast(v.data()), n);`
`123`		`- assert(status == CURAND_STATUS_SUCCESS /* && "curandGenerateUniform failed"*/);`
	`123`	`+ assert(status == CURAND_STATUS_SUCCESS /* && "curandGenerateUniform failed"*/);`
`124`	`124`	`_unused(status);`
	`125`	`+`
`125`	`126`	`}`
`126`	`127`
`127`	`128`	`void generate_n(`