diff --git a/lib/sapporo_light/dev_evaluate_gravity.cu b/lib/sapporo_light/dev_evaluate_gravity.cu index 2d5cff41bc..d4c967105c 100644 --- a/lib/sapporo_light/dev_evaluate_gravity.cu +++ b/lib/sapporo_light/dev_evaluate_gravity.cu @@ -73,7 +73,11 @@ __device__ void body_body_interaction(float &ds_min, } - float inv_ds = rsqrt(ds2 + EPS2) * (__float_as_int(pos_i.w.y) != __float_as_int(pos_j.w.y)); + float inv_ds = 0.0f; + if (__float_as_int(pos_i.w.y) != __float_as_int(pos_j.w.y)) { + inv_ds = rsqrt(ds2 + EPS2); + } + float mass = pos_j.w.x; float inv_ds2 = inv_ds*inv_ds; // 1 FLOP float inv_ds3 = mass * inv_ds*inv_ds2; // 2 FLOP diff --git a/lib/sapporo_light/send_fetch_data.cpp b/lib/sapporo_light/send_fetch_data.cpp index 7bcc2f5684..f74c6c0fc5 100644 --- a/lib/sapporo_light/send_fetch_data.cpp +++ b/lib/sapporo_light/send_fetch_data.cpp @@ -23,7 +23,7 @@ void sapporo::free_cuda_memory(int ignore) { CUDA_SAFE_CALL(cudaFree( (void*)dev.ngb_list_i)); - CUDA_SAFE_CALL(cudaThreadExit()); + CUDA_SAFE_CALL(cudaDeviceReset()); CUT_CHECK_ERROR("Failedn"); }