Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions gpu/pipelines/GeneralPipelineDisk.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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),
Expand Down
2 changes: 1 addition & 1 deletion gpu/pipelines/GeneralPipelineLine.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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),
Expand Down
2 changes: 1 addition & 1 deletion gpu/pipelines/GeneralPipelineTriangle.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
1 change: 1 addition & 0 deletions include/viennaray/gpu/raygLaunchParams.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
3 changes: 3 additions & 0 deletions include/viennaray/gpu/raygTrace.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -82,6 +82,8 @@ template <class T, int D> class Trace {
launchParams_.rayWeightThreshold = threshold;
}

void setTnear(float tnear) { config_.tnear = tnear; }

void apply() {
if (particles_.empty()) {
VIENNACORE_LOG_ERROR(
Expand Down Expand Up @@ -126,6 +128,7 @@ template <class T, int D> 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;
Expand Down
18 changes: 10 additions & 8 deletions include/viennaray/rayBoundary.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -19,10 +19,11 @@ template <typename NumericType, int D> class Boundary {
public:
Boundary(RTCDevice &device, boundingBoxType const &boundingBox,
BoundaryCondition boundaryConds[D],
const std::array<int, 5> &traceSettings)
const std::array<int, 5> &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);
}

Expand All @@ -39,7 +40,7 @@ template <typename NumericType, int D> 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;
}

Expand All @@ -60,7 +61,7 @@ template <typename NumericType, int D> 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 {
Expand Down Expand Up @@ -88,7 +89,7 @@ template <typename NumericType, int D> 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 {
Expand All @@ -112,7 +113,7 @@ template <typename NumericType, int D> 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 {
Expand Down Expand Up @@ -258,7 +259,7 @@ template <typename NumericType, int D> class Boundary {
(NumericType)pVertexBuffer_[tt.v2].zz};
}

static void reflectRay(RTCRayHit &rayHit, Vec3D<NumericType> &rayDirection) {
void reflectRay(RTCRayHit &rayHit, Vec3D<NumericType> &rayDirection) const {
auto n = *reinterpret_cast<Vec3D<float> *>(&rayHit.hit.Ng_x);
Vec3D<NumericType> normal{static_cast<NumericType>(n[0]),
static_cast<NumericType>(n[1]),
Expand All @@ -267,7 +268,7 @@ template <typename NumericType, int D> class Boundary {
rayDirection = ReflectionSpecular<NumericType>(rayDirection, normal);
auto const origin = getNewOrigin(rayHit.ray);
rayInternal::fillRayDirection<D>(rayHit.ray, rayDirection);
rayInternal::fillRayPosition(rayHit.ray, origin);
rayInternal::fillRayPosition(rayHit.ray, origin, tnear_);
}

struct vertex_f3_t {
Expand All @@ -291,6 +292,7 @@ template <typename NumericType, int D> class Boundary {
const int firstDir_ = 0;
const int secondDir_ = 1;
const std::array<BoundaryCondition, 2> boundaryConds_;
const float tnear_ = 1e-4f;
static constexpr size_t numTriangles_ = 8;
static constexpr size_t numVertices_ = 8;
};
Expand Down
11 changes: 8 additions & 3 deletions include/viennaray/rayTrace.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -102,16 +102,21 @@ template <class NumericType, int D> 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) {
Expand Down Expand Up @@ -188,7 +193,7 @@ template <class NumericType, int D> class Trace {

pBoundary_ = std::make_unique<Boundary<NumericType, D>>(
device_, committedBoundingBox_, boundaryConditions_,
committedTraceSettings_);
committedTraceSettings_, config_.tnear);

scene_.rtcScene = rtcNewScene(device_);
rtcSetSceneFlags(scene_.rtcScene, RTC_SCENE_FLAG_NONE);
Expand Down
8 changes: 4 additions & 4 deletions include/viennaray/rayTraceKernel.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -129,7 +129,7 @@ template <typename NumericType, int D, typename GeometryT> class TraceKernel {
fillRayDirection<D>(rayHit.ray, rayDirection);

auto origin = pSource_->getOrigin(idx, rngState);
fillRayPosition(rayHit.ray, origin);
fillRayPosition(rayHit.ray, origin, config_.tnear);

particle->logData(myDataLog);
}
Expand Down Expand Up @@ -179,7 +179,7 @@ template <typename NumericType, int D, typename GeometryT> 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<D>(rayHit.ray, rayDirection);

++particleHits;
Expand Down Expand Up @@ -219,7 +219,7 @@ template <typename NumericType, int D, typename GeometryT> 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 {
Expand Down Expand Up @@ -322,7 +322,7 @@ template <typename NumericType, int D, typename GeometryT> class TraceKernel {

// Update ray direction and origin
rayDirection = std::move(stickingDirection.second);
fillRayPosition(rayHit.ray, hitPoint);
fillRayPosition(rayHit.ray, hitPoint, config_.tnear);
fillRayDirection<D>(rayHit.ray, rayDirection);
if (lambda > 0.)
freeFlightDistance = sampleFreeFlightDistance(lambda, rngState);
Expand Down
7 changes: 5 additions & 2 deletions include/viennaray/rayUtil.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -98,14 +98,17 @@ using namespace viennacore;
struct KernelConfig {
size_t numRaysPerPoint = 1000;
size_t numRaysFixed = 0;

unsigned maxReflections = std::numeric_limits<unsigned>::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;
};

Expand Down
Loading