diff --git a/gpu/pipelines/GeneralPipelineDisk.cu b/gpu/pipelines/GeneralPipelineDisk.cu index ee394fd..d66a4d5 100644 --- a/gpu/pipelines/GeneralPipelineDisk.cu +++ b/gpu/pipelines/GeneralPipelineDisk.cu @@ -70,7 +70,7 @@ extern "C" __global__ void __closesthit__() { const Vec3Df &normal = sbtData->base.normal[primID]; - // If closest hit was on backside, let it through + // If closest hit was on backside, let it through if (DotProduct(prd->traceDir, normal) > 0.0f) { if (prd->numBackfaceHits++ > launchParams.maxBackfaceHits) { prd->rayWeight = 0.f; @@ -186,7 +186,7 @@ extern "C" __global__ void __raygen__() { make_float3(prd.pos[0], prd.pos[1], prd.pos[2]), // origin make_float3(prd.traceDir[0], prd.traceDir[1], prd.traceDir[2]), // direction - 1e-4f, // tmin + launchParams.tnear, // tmin 1e20f, // tmax 0.0f, // rayTime OptixVisibilityMask(255), diff --git a/gpu/pipelines/GeneralPipelineLine.cu b/gpu/pipelines/GeneralPipelineLine.cu index d43a63e..f7e33c1 100644 --- a/gpu/pipelines/GeneralPipelineLine.cu +++ b/gpu/pipelines/GeneralPipelineLine.cu @@ -131,7 +131,7 @@ extern "C" __global__ void __raygen__() { make_float3(prd.pos[0], prd.pos[1], prd.pos[2]), // origin make_float3(prd.traceDir[0], prd.traceDir[1], prd.traceDir[2]), // direction - 1e-4f, // tmin + launchParams.tnear, // tmin 1e20f, // tmax 0.0f, // rayTime OptixVisibilityMask(255), diff --git a/gpu/pipelines/GeneralPipelineTriangle.cu b/gpu/pipelines/GeneralPipelineTriangle.cu index b9e7803..f622120 100644 --- a/gpu/pipelines/GeneralPipelineTriangle.cu +++ b/gpu/pipelines/GeneralPipelineTriangle.cu @@ -122,7 +122,7 @@ extern "C" __global__ void __raygen__() { make_float3(prd.pos[0], prd.pos[1], prd.pos[2]), // origin make_float3(prd.traceDir[0], prd.traceDir[1], prd.traceDir[2]), // direction - 1e-4f, // tmin + launchParams.tnear, // tmin 1e20f, // tmax 0.0f, // rayTime OptixVisibilityMask(255), OPTIX_RAY_FLAG_DISABLE_ANYHIT, diff --git a/include/viennaray/gpu/raygLaunchParams.hpp b/include/viennaray/gpu/raygLaunchParams.hpp index 0dfdfa1..80be5ed 100644 --- a/include/viennaray/gpu/raygLaunchParams.hpp +++ b/include/viennaray/gpu/raygLaunchParams.hpp @@ -29,6 +29,7 @@ struct LaunchParams { float rayWeightThreshold = 0.1f; float tThreshold = 0.5f; + float tnear = 1e-4f; unsigned int seed = 0; bool periodicBoundary = false; diff --git a/include/viennaray/gpu/raygTrace.hpp b/include/viennaray/gpu/raygTrace.hpp index d5c9f61..f1512bb 100644 --- a/include/viennaray/gpu/raygTrace.hpp +++ b/include/viennaray/gpu/raygTrace.hpp @@ -82,6 +82,8 @@ template class Trace { launchParams_.rayWeightThreshold = threshold; } + void setTnear(float tnear) { config_.tnear = tnear; } + void apply() { if (particles_.empty()) { VIENNACORE_LOG_ERROR( @@ -126,6 +128,7 @@ template class Trace { // Threshold value for neighbor detection in disk-based geometries assert(gridDelta_ > 0.0f); launchParams_.tThreshold = 1.1 * gridDelta_; // TODO: find the best value + launchParams_.tnear = config_.tnear; launchParams_.maxReflections = config_.maxReflections; launchParams_.maxBoundaryHits = config_.maxBoundaryHits; diff --git a/include/viennaray/rayBoundary.hpp b/include/viennaray/rayBoundary.hpp index 22c632c..f59a164 100644 --- a/include/viennaray/rayBoundary.hpp +++ b/include/viennaray/rayBoundary.hpp @@ -19,10 +19,11 @@ template class Boundary { public: Boundary(RTCDevice &device, boundingBoxType const &boundingBox, BoundaryCondition boundaryConds[D], - const std::array &traceSettings) + const std::array &traceSettings, const float tnear = 1e-4f) : bdBox_(boundingBox), firstDir_(traceSettings[1]), secondDir_(traceSettings[2]), - boundaryConds_({boundaryConds[firstDir_], boundaryConds[secondDir_]}) { + boundaryConds_({boundaryConds[firstDir_], boundaryConds[secondDir_]}), + tnear_(tnear) { initBoundary(device); } @@ -39,7 +40,7 @@ template class Boundary { // let ray pass through reflect = true; const auto impactCoords = getNewOrigin(rayHit.ray); - rayInternal::fillRayPosition(rayHit.ray, impactCoords); + rayInternal::fillRayPosition(rayHit.ray, impactCoords, tnear_); return; } @@ -60,7 +61,7 @@ template class Boundary { // hit at x/y max boundary -> move to min x/y impactCoords[firstDir_] = bdBox_[0][firstDir_]; } - rayInternal::fillRayPosition(rayHit.ray, impactCoords); + rayInternal::fillRayPosition(rayHit.ray, impactCoords, tnear_); reflect = true; return; } else { @@ -88,7 +89,7 @@ template class Boundary { // hit at firstDir max boundary -> move to min firstDir impactCoords[firstDir_] = bdBox_[0][firstDir_]; } - rayInternal::fillRayPosition(rayHit.ray, impactCoords); + rayInternal::fillRayPosition(rayHit.ray, impactCoords, tnear_); reflect = true; return; } else { @@ -112,7 +113,7 @@ template class Boundary { // hit at secondDir max boundary -> move to min secondDir impactCoords[secondDir_] = bdBox_[0][secondDir_]; } - rayInternal::fillRayPosition(rayHit.ray, impactCoords); + rayInternal::fillRayPosition(rayHit.ray, impactCoords, tnear_); reflect = true; return; } else { @@ -258,7 +259,7 @@ template class Boundary { (NumericType)pVertexBuffer_[tt.v2].zz}; } - static void reflectRay(RTCRayHit &rayHit, Vec3D &rayDirection) { + void reflectRay(RTCRayHit &rayHit, Vec3D &rayDirection) const { auto n = *reinterpret_cast *>(&rayHit.hit.Ng_x); Vec3D normal{static_cast(n[0]), static_cast(n[1]), @@ -267,7 +268,7 @@ template class Boundary { rayDirection = ReflectionSpecular(rayDirection, normal); auto const origin = getNewOrigin(rayHit.ray); rayInternal::fillRayDirection(rayHit.ray, rayDirection); - rayInternal::fillRayPosition(rayHit.ray, origin); + rayInternal::fillRayPosition(rayHit.ray, origin, tnear_); } struct vertex_f3_t { @@ -291,6 +292,7 @@ template class Boundary { const int firstDir_ = 0; const int secondDir_ = 1; const std::array boundaryConds_; + const float tnear_ = 1e-4f; static constexpr size_t numTriangles_ = 8; static constexpr size_t numVertices_ = 8; }; diff --git a/include/viennaray/rayTrace.hpp b/include/viennaray/rayTrace.hpp index 427def2..c7c2ced 100644 --- a/include/viennaray/rayTrace.hpp +++ b/include/viennaray/rayTrace.hpp @@ -102,16 +102,21 @@ template class Trace { config_.maxReflections = maxReflections; } + /// Set the maximum number of boundary hits a ray is allowed to perform. void setMaxBoundaryHits(const unsigned maxBoundaryHits) { config_.maxBoundaryHits = maxBoundaryHits; } - // Set the maximum number of backface hits a ray is allowed to perform. Has no - // effect when using triangle geometry. + /// Set the maximum number of backface hits a ray is allowed to perform. Has + /// no effect when using triangle geometry. void setMaxBackfaceHits(const unsigned maxBackfaceHits) { config_.maxBackfaceHits = maxBackfaceHits; } + /// Set the minimum distance a ray is allowed to travel before it can hit a + /// geometry point. This is used to avoid self-intersections. + void setTnear(const float tnear) { config_.tnear = tnear; } + /// Set the source direction, where the rays should be traced from. void setSourceDirection(const TraceDirection direction) { if (sourceDirection_ == direction) { @@ -188,7 +193,7 @@ template class Trace { pBoundary_ = std::make_unique>( device_, committedBoundingBox_, boundaryConditions_, - committedTraceSettings_); + committedTraceSettings_, config_.tnear); scene_.rtcScene = rtcNewScene(device_); rtcSetSceneFlags(scene_.rtcScene, RTC_SCENE_FLAG_NONE); diff --git a/include/viennaray/rayTraceKernel.hpp b/include/viennaray/rayTraceKernel.hpp index c0ee831..6819acb 100644 --- a/include/viennaray/rayTraceKernel.hpp +++ b/include/viennaray/rayTraceKernel.hpp @@ -129,7 +129,7 @@ template class TraceKernel { fillRayDirection(rayHit.ray, rayDirection); auto origin = pSource_->getOrigin(idx, rngState); - fillRayPosition(rayHit.ray, origin); + fillRayPosition(rayHit.ray, origin, config_.tnear); particle->logData(myDataLog); } @@ -179,7 +179,7 @@ template class TraceKernel { // A scattering event starts a new free flight. freeFlightDistance = sampleFreeFlightDistance(lambda, rngState); - fillRayPosition(rayHit.ray, origin); + fillRayPosition(rayHit.ray, origin, config_.tnear); fillRayDirection(rayHit.ray, rayDirection); ++particleHits; @@ -219,7 +219,7 @@ template class TraceKernel { } // Let ray through, i.e., continue. reflect = true; - fillRayPosition(rayHit.ray, hitPoint); + fillRayPosition(rayHit.ray, hitPoint, config_.tnear); // keep ray direction as it is continue; } else { @@ -322,7 +322,7 @@ template class TraceKernel { // Update ray direction and origin rayDirection = std::move(stickingDirection.second); - fillRayPosition(rayHit.ray, hitPoint); + fillRayPosition(rayHit.ray, hitPoint, config_.tnear); fillRayDirection(rayHit.ray, rayDirection); if (lambda > 0.) freeFlightDistance = sampleFreeFlightDistance(lambda, rngState); diff --git a/include/viennaray/rayUtil.hpp b/include/viennaray/rayUtil.hpp index ad4dd11..96d643b 100644 --- a/include/viennaray/rayUtil.hpp +++ b/include/viennaray/rayUtil.hpp @@ -98,14 +98,17 @@ using namespace viennacore; struct KernelConfig { size_t numRaysPerPoint = 1000; size_t numRaysFixed = 0; + unsigned maxReflections = std::numeric_limits::max(); unsigned maxBoundaryHits = 1000; unsigned maxBackfaceHits = 1; - unsigned rngSeed = 0; + float tnear = 1e-4f; + + unsigned rngSeed = 0; bool useRandomSeed = true; - bool printProgress = false; + bool printProgress = false; unsigned runNumber = 1; };