diff options
| author | Stan Seibert <stan@mtrr.org> | 2011-08-13 14:46:10 -0400 |
|---|---|---|
| committer | Stan Seibert <stan@mtrr.org> | 2011-08-13 14:46:10 -0400 |
| commit | a9aed9f7d07048f61e3c5c8214834628e1f65b36 (patch) | |
| tree | bab6c7e18a0adb78da61f6ce80e7e3a76f21e84f /src | |
| parent | c6961d11edec689471d3f268f58bbab68f57a239 (diff) | |
| download | chroma-a9aed9f7d07048f61e3c5c8214834628e1f65b36.tar.gz chroma-a9aed9f7d07048f61e3c5c8214834628e1f65b36.tar.bz2 chroma-a9aed9f7d07048f61e3c5c8214834628e1f65b36.zip | |
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.
Diffstat (limited to 'src')
| -rw-r--r-- | src/kernel.cu | 7 | ||||
| -rw-r--r-- | src/photon.h | 3 |
2 files changed, 9 insertions, 1 deletions
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 |