From a9aed9f7d07048f61e3c5c8214834628e1f65b36 Mon Sep 17 00:00:00 2001 From: Stan Seibert Date: Sat, 13 Aug 2011 14:46:10 -0400 Subject: A faulty optical process can make the position or direction of the photon into NaN on the GPU. Now we abort these photons rather than let them lock up the intersect_mesh() method. There is a new history bit (#31) that indicates when a NAN_ABORT has occurred, and this bit is checked for by GPU.propagate(). If set for any of the photons, a warning message is printed. While not as good as preventing the NaN problem in the first place, this at least ensures we are aware of the problem. --- src/kernel.cu | 7 +++++++ 1 file changed, 7 insertions(+) (limited to 'src/kernel.cu') diff --git a/src/kernel.cu b/src/kernel.cu index cc567b8..307412e 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -345,6 +345,13 @@ __global__ void propagate(int first_photon, int nthreads, int command; + // check for NaN and fail + if (isnan(p.direction.x*p.direction.y*p.direction.z*p.position.x*p.position.y*p.position.z)) + { + p.history |= NO_HIT | NAN_ABORT; + break; + } + fill_state(s, p); if (p.last_hit_triangle == -1) -- cgit