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 +++++++ src/photon.h | 3 ++- 2 files changed, 9 insertions(+), 1 deletion(-) (limited to 'src') 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) diff --git a/src/photon.h b/src/photon.h index f471866..4203a50 100644 --- a/src/photon.h +++ b/src/photon.h @@ -44,7 +44,8 @@ enum SURFACE_ABSORB = 0x1 << 3, RAYLEIGH_SCATTER = 0x1 << 4, REFLECT_DIFFUSE = 0x1 << 5, - REFLECT_SPECULAR = 0x1 << 6 + REFLECT_SPECULAR = 0x1 << 6, + NAN_ABORT = 0x1 << 31 }; // processes enum {BREAK, CONTINUE, PASS}; // return value from propagate_to_boundary -- cgit