Skip to content

Cleaner GPU Line Engine #78

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
tobre1 merged 1 commit into from
Oct 27, 2025
Merged

Cleaner GPU Line Engine #78

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
69 changes: 2 additions & 67 deletions gpu/include/raygBoundary.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,8 +35,7 @@ __device__ __inline__ void reflectFromBoundary(viennaray::gpu::PerRayData *prd,
prd->numBoundaryHits++;
} else if constexpr (std::is_same<SBTData,
viennaray::gpu::HitSBTDataLine>::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];
}
Expand Down Expand Up @@ -70,76 +69,12 @@ applyPeriodicBoundary(viennaray::gpu::PerRayData *prd, const SBTData *hsd,
prd->numBoundaryHits++;
} else if constexpr (std::is_same<SBTData,
viennaray::gpu::HitSBTDataLine>::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
2 changes: 0 additions & 2 deletions gpu/include/raygTrace.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -340,7 +340,6 @@ template <class T, int D> class Trace {
for (auto &buffer : materialStickingBuffer_) {
buffer.free();
}
neighborsBuffer_.free();
areaBuffer_.free();
}

Expand Down Expand Up @@ -663,7 +662,6 @@ template <class T, int D> class Trace {
std::set<int> uniqueMaterialIds_;
CudaBuffer materialIdsBuffer_;

CudaBuffer neighborsBuffer_;
float gridDelta_ = 0.0f;

CudaBuffer areaBuffer_;
Expand Down
57 changes: 6 additions & 51 deletions gpu/pipelines/GeneralPipelineLine.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
}

Expand All @@ -64,56 +56,19 @@ extern "C" __global__ void __closesthit__() {
PerRayData *prd = (PerRayData *)getPRD<PerRayData>();

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 =
Expand Down
10 changes: 4 additions & 6 deletions gpu/pipelines/GeneralPipelineTriangle.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,11 +36,9 @@ extern "C" __global__ void __closesthit__() {
if (sbtData->base.isBoundary) {
prd->numBoundaryHits++;
if (launchParams.periodicBoundary) {
applyPeriodicBoundary<viennaray::gpu::HitSBTDataTriangle>(prd, sbtData,
launchParams.D);
applyPeriodicBoundary(prd, sbtData, launchParams.D);
} else {
reflectFromBoundary<viennaray::gpu::HitSBTDataTriangle>(prd, sbtData,
launchParams.D);
reflectFromBoundary(prd, sbtData, launchParams.D);
}
} else {
prd->ISCount = 1;
Expand Down Expand Up @@ -80,8 +78,8 @@ extern "C" __global__ void __raygen__() {

unsigned callIdx =
callableIndex(launchParams.particleType, CallableSlot::INIT);
optixDirectCall<void, const HitSBTDataDisk *, PerRayData *>(callIdx, nullptr,
&prd);
optixDirectCall<void, const HitSBTDataTriangle *, PerRayData *>(
callIdx, nullptr, &prd);

// the values we store the PRD pointer in:
uint32_t u0, u1;
Expand Down