diff --git a/gpu/pipelines/GeneralPipelineDisk.cu b/gpu/pipelines/GeneralPipelineDisk.cu index 8bc70c7..f5a3d1e 100644 --- a/gpu/pipelines/GeneralPipelineDisk.cu +++ b/gpu/pipelines/GeneralPipelineDisk.cu @@ -172,21 +172,33 @@ extern "C" __global__ void __raygen__() { // the values we store the PRD pointer in: uint32_t u0, u1; packPointer((void *)&prd, u0, u1); + unsigned int hintBitLength = 2; while (continueRay(launchParams, prd)) { - optixTrace(launchParams.traversable, // traversable GAS - make_float3(prd.pos[0], prd.pos[1], prd.pos[2]), // origin - make_float3(prd.dir[0], prd.dir[1], prd.dir[2]), // direction - 1e-4f, // tmin - 1e20f, // tmax - 0.0f, // rayTime - OptixVisibilityMask(255), - OPTIX_RAY_FLAG_DISABLE_ANYHIT, // OPTIX_RAY_FLAG_NONE, - 0, // SBT offset - 1, // SBT stride - 0, // missSBTIndex - u0, u1); // Payload - prd.totalCount = 0; // Reset PerRayData + optixTraverse(launchParams.traversable, // traversable GAS + make_float3(prd.pos[0], prd.pos[1], prd.pos[2]), // origin + make_float3(prd.dir[0], prd.dir[1], prd.dir[2]), // direction + 1e-4f, // tmin + 1e20f, // tmax + 0.0f, // rayTime + OptixVisibilityMask(255), + OPTIX_RAY_FLAG_DISABLE_ANYHIT, // OPTIX_RAY_FLAG_NONE, + 0, // SBT offset + 1, // SBT stride + 0, // missSBTIndex + u0, u1); // Payload + unsigned int hint = 0; + if (prd.rayWeight < launchParams.rayWeightThreshold || prd.energy < 0.f) { + hint |= (1 << 0); + } + if (optixHitObjectIsHit()) { + const HitSBTDataDisk *hitData = reinterpret_cast( + optixHitObjectGetSbtDataPointer()); + hint |= hitData->base.isBoundary << 1; + } + optixReorder(hint, hintBitLength); + optixInvoke(u0, u1); + prd.totalCount = 0; // Reset PerRayData prd.numReflections++; } } diff --git a/gpu/pipelines/GeneralPipelineLine.cu b/gpu/pipelines/GeneralPipelineLine.cu index ecc45ea..ddb97db 100644 --- a/gpu/pipelines/GeneralPipelineLine.cu +++ b/gpu/pipelines/GeneralPipelineLine.cu @@ -114,21 +114,33 @@ extern "C" __global__ void __raygen__() { // the values we store the PRD pointer in: uint32_t u0, u1; packPointer((void *)&prd, u0, u1); + unsigned int hintBitLength = 2; while (continueRay(launchParams, prd)) { - optixTrace(launchParams.traversable, // traversable GAS - make_float3(prd.pos[0], prd.pos[1], prd.pos[2]), // origin - make_float3(prd.dir[0], prd.dir[1], prd.dir[2]), // direction - 1e-4f, // tmin - 1e20f, // tmax - 0.0f, // rayTime - OptixVisibilityMask(255), - OPTIX_RAY_FLAG_DISABLE_ANYHIT, // OPTIX_RAY_FLAG_NONE, - 0, // SBT offset - 1, // SBT stride - 0, // missSBTIndex - u0, u1); // Payload - prd.totalCount = 0; // Reset PerRayData + optixTraverse(launchParams.traversable, // traversable GAS + make_float3(prd.pos[0], prd.pos[1], prd.pos[2]), // origin + make_float3(prd.dir[0], prd.dir[1], prd.dir[2]), // direction + 1e-4f, // tmin + 1e20f, // tmax + 0.0f, // rayTime + OptixVisibilityMask(255), + OPTIX_RAY_FLAG_DISABLE_ANYHIT, // OPTIX_RAY_FLAG_NONE, + 0, // SBT offset + 1, // SBT stride + 0, // missSBTIndex + u0, u1); // Payload + unsigned int hint = 0; + if (prd.rayWeight < launchParams.rayWeightThreshold || prd.energy < 0.f) { + hint |= (1 << 0); + } + if (optixHitObjectIsHit()) { + const HitSBTDataLine *hitData = reinterpret_cast( + optixHitObjectGetSbtDataPointer()); + hint |= hitData->base.isBoundary << 1; + } + optixReorder(hint, hintBitLength); + optixInvoke(u0, u1); + prd.totalCount = 0; // Reset PerRayData prd.numReflections++; } } diff --git a/gpu/pipelines/GeneralPipelineTriangle.cu b/gpu/pipelines/GeneralPipelineTriangle.cu index 3fd166a..7e293ff 100644 --- a/gpu/pipelines/GeneralPipelineTriangle.cu +++ b/gpu/pipelines/GeneralPipelineTriangle.cu @@ -84,20 +84,32 @@ extern "C" __global__ void __raygen__() { // the values we store the PRD pointer in: uint32_t u0, u1; packPointer((void *)&prd, u0, u1); + unsigned int hintBitLength = 2; while (continueRay(launchParams, prd)) { - optixTrace(launchParams.traversable, // traversable GAS - make_float3(prd.pos[0], prd.pos[1], prd.pos[2]), // origin - make_float3(prd.dir[0], prd.dir[1], prd.dir[2]), // direction - 1e-4f, // tmin - 1e20f, // tmax - 0.0f, // rayTime - OptixVisibilityMask(255), - OPTIX_RAY_FLAG_DISABLE_ANYHIT, // OPTIX_RAY_FLAG_NONE, - 0, // SBT offset - 1, // SBT stride - 0, // missSBTIndex - u0, u1); // Payload + optixTraverse(launchParams.traversable, // traversable GAS + make_float3(prd.pos[0], prd.pos[1], prd.pos[2]), // origin + make_float3(prd.dir[0], prd.dir[1], prd.dir[2]), // direction + 1e-4f, // tmin + 1e20f, // tmax + 0.0f, // rayTime + OptixVisibilityMask(255), + OPTIX_RAY_FLAG_DISABLE_ANYHIT, // OPTIX_RAY_FLAG_NONE, + 0, // SBT offset + 1, // SBT stride + 0, // missSBTIndex + u0, u1); // Payload + unsigned int hint = 0; + if (prd.rayWeight < launchParams.rayWeightThreshold || prd.energy < 0.f) { + hint |= (1 << 0); + } + if (optixHitObjectIsHit()) { + const HitSBTDataTriangle *hitData = reinterpret_cast( + optixHitObjectGetSbtDataPointer()); + hint |= hitData->base.isBoundary << 1; + } + optixReorder(hint, hintBitLength); + optixInvoke(u0, u1); prd.numReflections++; #ifdef COUNT_RAYS int *counter = reinterpret_cast(launchParams.customData);