From b862904600eb5d3e59a3e539a1bd544fd2c8c8f0 Mon Sep 17 00:00:00 2001 From: Xaver Riedel Date: Fri, 24 Oct 2025 19:42:07 +0200 Subject: [PATCH] remove line engine neighbor detection & old neighbor buffer --- gpu/include/raygBoundary.hpp | 69 +----------------------- gpu/include/raygTrace.hpp | 2 - gpu/pipelines/GeneralPipelineLine.cu | 57 +++----------------- gpu/pipelines/GeneralPipelineTriangle.cu | 10 ++-- 4 files changed, 12 insertions(+), 126 deletions(-) diff --git a/gpu/include/raygBoundary.hpp b/gpu/include/raygBoundary.hpp index d72d35a..5f423b6 100644 --- a/gpu/include/raygBoundary.hpp +++ b/gpu/include/raygBoundary.hpp @@ -35,8 +35,7 @@ __device__ __inline__ void reflectFromBoundary(viennaray::gpu::PerRayData *prd, prd->numBoundaryHits++; } else if constexpr (std::is_same::value) { - prd->pos = - prd->pos + prd->dir * (optixGetRayTmax() - launchParams.tThreshold); + prd->pos = prd->pos + prd->dir * (optixGetRayTmax()); if (primID == 0 || primID == 1) // x boundary prd->dir[0] -= 2 * prd->dir[0]; } @@ -70,76 +69,12 @@ applyPeriodicBoundary(viennaray::gpu::PerRayData *prd, const SBTData *hsd, prd->numBoundaryHits++; } else if constexpr (std::is_same::value) { - prd->pos = - prd->pos + prd->dir * (optixGetRayTmax() - launchParams.tThreshold); + prd->pos = prd->pos + prd->dir * (optixGetRayTmax()); if (primID == 0) { // xmin prd->pos[0] = hsd->nodes[1][0]; } else if (primID == 1) { // xmax prd->pos[0] = hsd->nodes[0][0]; } } - - // prd->pos = prd->pos + optixGetRayTmax() * prd->dir; - - // // mappedId = primID ^ 2; - // // IMPORTANT: THIS ONLY WORKS WITH CURRENT BOUNDARY SETUP - // // 0, 1 -> -x - // // 2, 3 -> +x - // // 4, 5 -> -y - // // 6, 7 -> +y - // unsigned dim = primID / 4; - // prd->pos[dim] = hsd->vertex[hsd->index[primID ^ 2][0]][dim]; - // prd->numBoundaryHits++; - - // // if (primID == 0 || primID == 1) // x min - // // { - // // prd->pos[dim] = hsd->vertex[hsd->index[2][0]][dim]; // set to x max - // // } else if (primID == 2 || primID == 3) // x max - // // { - // // prd->pos[dim] = hsd->vertex[hsd->index[0][0]][dim]; // set to x min - // // } else if (primID == 4 || primID == 5) // y min - // // { - // // prd->pos[dim] = hsd->vertex[hsd->index[6][0]][dim]; // set to y max - // // } else if (primID == 6 || primID == 7) // y max - // // { - // // prd->pos[dim] = hsd->vertex[hsd->index[4][0]][dim]; // set to y min - // // } } - -// __device__ __inline__ void -// reflectFromBoundaryDisk(viennaray::gpu::PerRayData *prd, -// const viennaray::gpu::HitSBTDataDisk *hsd, -// const int D) { -// using namespace viennacore; -// const unsigned int primID = optixGetPrimitiveIndex(); -// prd->pos = prd->pos + prd->dir * (optixGetRayTmax() - -// launchParams.tThreshold); - -// if (primID == 0 || primID == 1) { -// prd->dir[0] -= 2 * prd->dir[0]; // x boundary -// } else if ((primID == 2 || primID == 3) && D == 3) { -// prd->dir[1] -= 2 * prd->dir[1]; // y boundary -// } -// } - -// __device__ __inline__ void -// applyPeriodicBoundaryDisk(viennaray::gpu::PerRayData *prd, -// const viennaray::gpu::HitSBTDataDisk *hsd, -// const int D) { -// using namespace viennacore; -// const unsigned int primID = optixGetPrimitiveIndex(); -// // prd->pos = prd->pos + prd->dir * optixGetRayTmax(); -// prd->pos = prd->pos + prd->dir * (optixGetRayTmax() - -// launchParams.tThreshold); - -// if (primID == 0) { // xmin -// prd->pos[0] = hsd->point[1][0]; -// } else if (primID == 1) { // xmax -// prd->pos[0] = hsd->point[0][0]; -// } else if (D == 3 && primID == 2) { // ymin -// prd->pos[1] = hsd->point[3][1]; -// } else if (D == 3 && primID == 3) { // ymax -// prd->pos[1] = hsd->point[2][1]; -// } -// } #endif diff --git a/gpu/include/raygTrace.hpp b/gpu/include/raygTrace.hpp index 5b66a8b..445e36a 100644 --- a/gpu/include/raygTrace.hpp +++ b/gpu/include/raygTrace.hpp @@ -340,7 +340,6 @@ template class Trace { for (auto &buffer : materialStickingBuffer_) { buffer.free(); } - neighborsBuffer_.free(); areaBuffer_.free(); } @@ -663,7 +662,6 @@ template class Trace { std::set uniqueMaterialIds_; CudaBuffer materialIdsBuffer_; - CudaBuffer neighborsBuffer_; float gridDelta_ = 0.0f; CudaBuffer areaBuffer_; diff --git a/gpu/pipelines/GeneralPipelineLine.cu b/gpu/pipelines/GeneralPipelineLine.cu index ce9995b..ec456af 100644 --- a/gpu/pipelines/GeneralPipelineLine.cu +++ b/gpu/pipelines/GeneralPipelineLine.cu @@ -41,20 +41,12 @@ extern "C" __global__ void __intersection__() { bool valid = true; float t = CrossProduct((p0 - prd->pos), lineDir)[2] / d[2]; - valid &= t > 1e-4f; + valid &= t > 1e-5f; float s = CrossProduct((p0 - prd->pos), prd->dir)[2] / d[2]; - valid &= s > 1e-4f && s < 1.0f - 1e-4f; + valid &= s > 1e-5f && s < 1.0f - 1e-5f; if (valid) { - // Collect all intersections and filter out overlapping lines in CH shader - if (!sbtData->base.isBoundary && prd->tempCount < MAX_NEIGHBORS) { - prd->tValues[prd->tempCount] = t; - prd->primIDs[prd->tempCount] = primID; - prd->tempCount++; - } - - // Has to pass a dummy t value so later intersections are not ignored - optixReportIntersection(t + launchParams.tThreshold, 0); + optixReportIntersection(t, 0); } } @@ -64,56 +56,19 @@ extern "C" __global__ void __closesthit__() { PerRayData *prd = (PerRayData *)getPRD(); const unsigned int primID = optixGetPrimitiveIndex(); - prd->tMin = optixGetRayTmax() - launchParams.tThreshold; + prd->tMin = optixGetRayTmax(); prd->primID = primID; - const Vec3Df normal = computeNormal(sbtData, primID); - - // If closest hit was on backside, let it through once - if (DotProduct(prd->dir, normal) > 0.0f) { - // If back was hit a second time, kill the ray - if (prd->hitFromBack) { - prd->rayWeight = 0.f; - return; - } - prd->hitFromBack = true; - prd->pos = prd->pos + prd->tMin * prd->dir; - return; - } - if (sbtData->base.isBoundary) { prd->numBoundaryHits++; - // This is effectively the miss shader - if (primID == 2 || primID == 3) { // bottom or top - ymin or ymax - prd->rayWeight = 0.0f; - return; - } - if (launchParams.periodicBoundary) { applyPeriodicBoundary(prd, sbtData, launchParams.D); } else { reflectFromBoundary(prd, sbtData, launchParams.D); } } else { - // ------------- NEIGHBOR FILTERING --------------- // - // Keep only hits close to tMin - prd->ISCount = 0; - for (int i = 0; i < prd->tempCount; ++i) { - if (fabsf(prd->tValues[i] - prd->tMin) < launchParams.tThreshold && - prd->ISCount < MAX_NEIGHBORS) { - prd->TIndex[prd->ISCount++] = prd->primIDs[i]; - } - } - - // // CPU like neighbor detection - // prd->ISCount = 0; - // for (int i = 0; i < prd->tempCount; ++i) { - // float distance = viennacore::Distance(sbtData->point[primID], - // sbtData->point[prd->primIDs[i]]); - // if (distance < 2 * sbtData->radius && prd->ISCount < MAX_NEIGHBORS) { - // prd->TIndex[prd->ISCount++] = prd->primIDs[i]; - // } - // } + prd->ISCount = 1; + prd->TIndex[0] = primID; // ------------- SURFACE COLLISION --------------- // unsigned callIdx = diff --git a/gpu/pipelines/GeneralPipelineTriangle.cu b/gpu/pipelines/GeneralPipelineTriangle.cu index 073053a..8011df6 100644 --- a/gpu/pipelines/GeneralPipelineTriangle.cu +++ b/gpu/pipelines/GeneralPipelineTriangle.cu @@ -36,11 +36,9 @@ extern "C" __global__ void __closesthit__() { if (sbtData->base.isBoundary) { prd->numBoundaryHits++; if (launchParams.periodicBoundary) { - applyPeriodicBoundary(prd, sbtData, - launchParams.D); + applyPeriodicBoundary(prd, sbtData, launchParams.D); } else { - reflectFromBoundary(prd, sbtData, - launchParams.D); + reflectFromBoundary(prd, sbtData, launchParams.D); } } else { prd->ISCount = 1; @@ -80,8 +78,8 @@ extern "C" __global__ void __raygen__() { unsigned callIdx = callableIndex(launchParams.particleType, CallableSlot::INIT); - optixDirectCall(callIdx, nullptr, - &prd); + optixDirectCall( + callIdx, nullptr, &prd); // the values we store the PRD pointer in: uint32_t u0, u1;