From bb08302b6844e6da06aecdd63f39dc4345341395 Mon Sep 17 00:00:00 2001 From: Hanno Spreeuw Date: Thu, 18 Dec 2025 11:00:47 +0100 Subject: [PATCH 1/2] Fixes #1195 Apparently, this line of code: " float inv_ds = rsqrt(ds2 + EPS2) * (__float_as_int(pos_i.w.y) != __float_as_int(pos_j.w.y)); " which should result in "inv_ds==0" for "i==j" results in "inv_ds==NaN" for "i==j" for modern CUDA versions. We suspect that for older CUDA versions this worked well and resulted in "inv_ds==0" for "i==j", i.e. for fictious self-gravitation - which we exclude, of course. Most likely, older CUDA versions did not comply with IEEE 754. --- lib/sapporo_light/dev_evaluate_gravity.cu | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) 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 From b47d5ad87da814ba06e677aa3331f8d9f43f2961 Mon Sep 17 00:00:00 2001 From: Hanno Spreeuw Date: Thu, 18 Dec 2025 11:22:14 +0100 Subject: [PATCH 2/2] "cudaThreadExit" is deprecated See https://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/html/group__CUDART__THREAD__DEPRECATED_gf423ba04af587d42b52799455a7c094d.html: "cudaThreadExit" has been deprecated: " Note that this function is deprecated because its name does not reflect its behavior. Its functionality is identical to the non-deprecated function cudaDeviceReset(), which should be used instead. " --- lib/sapporo_light/send_fetch_data.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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"); }