Changeset View
Changeset View
Standalone View
Standalone View
intern/cycles/kernel/kernels/optix/kernel_optix.cu
| Show All 9 Lines | |||||
| * | * | ||||
| * Unless required by applicable law or agreed to in writing, software | * Unless required by applicable law or agreed to in writing, software | ||||
| * distributed under the License is distributed on an "AS IS" BASIS, | * distributed under the License is distributed on an "AS IS" BASIS, | ||||
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||||
| * See the License for the specific language governing permissions and | * See the License for the specific language governing permissions and | ||||
| * limitations under the License. | * limitations under the License. | ||||
| */ | */ | ||||
| // clang-format off | |||||
| #include "kernel/kernel_compat_optix.h" | #include "kernel/kernel_compat_optix.h" | ||||
| #include "util/util_atomic.h" | #include "util/util_atomic.h" | ||||
| #include "kernel/kernel_types.h" | #include "kernel/kernel_types.h" | ||||
| #include "kernel/kernel_globals.h" | #include "kernel/kernel_globals.h" | ||||
| #include "../cuda/kernel_cuda_image.h" // Texture lookup uses normal CUDA intrinsics | #include "../cuda/kernel_cuda_image.h" // Texture lookup uses normal CUDA intrinsics | ||||
| #include "kernel/kernel_path.h" | #include "kernel/kernel_path.h" | ||||
| #include "kernel/kernel_bake.h" | #include "kernel/kernel_bake.h" | ||||
| // clang-format on | |||||
| template<typename T> ccl_device_forceinline T *get_payload_ptr_0() | template<typename T> ccl_device_forceinline T *get_payload_ptr_0() | ||||
| { | { | ||||
| return (T *)(((uint64_t)optixGetPayload_1() << 32) | optixGetPayload_0()); | return (T *)(((uint64_t)optixGetPayload_1() << 32) | optixGetPayload_0()); | ||||
| } | } | ||||
| template<typename T> ccl_device_forceinline T *get_payload_ptr_2() | template<typename T> ccl_device_forceinline T *get_payload_ptr_2() | ||||
| { | { | ||||
| return (T *)(((uint64_t)optixGetPayload_3() << 32) | optixGetPayload_2()); | return (T *)(((uint64_t)optixGetPayload_3() << 32) | optixGetPayload_2()); | ||||
| ▲ Show 20 Lines • Show All 100 Lines • ▼ Show 20 Lines | if (local_isect->num_hits > max_hits) { | ||||
| hit = lcg_step_uint(lcg_state) % local_isect->num_hits; | hit = lcg_step_uint(lcg_state) % local_isect->num_hits; | ||||
| if (hit >= max_hits) { | if (hit >= max_hits) { | ||||
| return optixIgnoreIntersection(); | return optixIgnoreIntersection(); | ||||
| } | } | ||||
| } | } | ||||
| } | } | ||||
| else { | else { | ||||
| if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) { | if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) { | ||||
| // Record closest intersection only (do not terminate ray here, since there is no guarantee | // Record closest intersection only | ||||
| // about distance ordering in anyhit) | // Do not terminate ray here, since there is no guarantee about distance ordering in any-hit | ||||
| return optixIgnoreIntersection(); | return optixIgnoreIntersection(); | ||||
| } | } | ||||
| local_isect->num_hits = 1; | local_isect->num_hits = 1; | ||||
| } | } | ||||
| Intersection *isect = &local_isect->hits[hit]; | Intersection *isect = &local_isect->hits[hit]; | ||||
| isect->t = optixGetRayTmax(); | isect->t = optixGetRayTmax(); | ||||
| isect->prim = optixGetPrimitiveIndex(); | isect->prim = optixGetPrimitiveIndex(); | ||||
| isect->object = get_object_id(); | isect->object = get_object_id(); | ||||
| isect->type = kernel_tex_fetch(__prim_type, isect->prim); | isect->type = kernel_tex_fetch(__prim_type, isect->prim); | ||||
| if (optixIsTriangleHit()) { | |||||
| const float2 barycentrics = optixGetTriangleBarycentrics(); | const float2 barycentrics = optixGetTriangleBarycentrics(); | ||||
| isect->u = 1.0f - barycentrics.y - barycentrics.x; | isect->u = 1.0f - barycentrics.y - barycentrics.x; | ||||
| isect->v = barycentrics.x; | isect->v = barycentrics.x; | ||||
| } | |||||
| else { | |||||
| isect->u = __uint_as_float(optixGetAttribute_0()); | |||||
| isect->v = __uint_as_float(optixGetAttribute_1()); | |||||
| } | |||||
| // Record geometric normal | // Record geometric normal | ||||
| const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim); | const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim); | ||||
| const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0)); | const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0)); | ||||
| const float3 tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1)); | const float3 tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1)); | ||||
| const float3 tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2)); | const float3 tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2)); | ||||
| local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a)); | local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a)); | ||||
| Show All 20 Lines | # endif | ||||
| isect->object = get_object_id(); | isect->object = get_object_id(); | ||||
| isect->type = kernel_tex_fetch(__prim_type, prim); | isect->type = kernel_tex_fetch(__prim_type, prim); | ||||
| if (optixIsTriangleHit()) { | if (optixIsTriangleHit()) { | ||||
| const float2 barycentrics = optixGetTriangleBarycentrics(); | const float2 barycentrics = optixGetTriangleBarycentrics(); | ||||
| isect->u = 1.0f - barycentrics.y - barycentrics.x; | isect->u = 1.0f - barycentrics.y - barycentrics.x; | ||||
| isect->v = barycentrics.x; | isect->v = barycentrics.x; | ||||
| } | } | ||||
| # ifdef __HAIR__ | |||||
| else { | else { | ||||
| isect->u = __uint_as_float(optixGetAttribute_0()); | const float u = __uint_as_float(optixGetAttribute_0()); | ||||
| isect->u = u; | |||||
| isect->v = __uint_as_float(optixGetAttribute_1()); | isect->v = __uint_as_float(optixGetAttribute_1()); | ||||
| // Filter out curve endcaps | |||||
| if (u == 0.0f || u == 1.0f) { | |||||
| return optixIgnoreIntersection(); | |||||
| } | } | ||||
| } | |||||
| # endif | |||||
| # ifdef __TRANSPARENT_SHADOWS__ | # ifdef __TRANSPARENT_SHADOWS__ | ||||
| // Detect if this surface has a shader with transparent shadows | // Detect if this surface has a shader with transparent shadows | ||||
| if (!shader_transparent_shadow(NULL, isect) || optixGetPayload_2() >= optixGetPayload_3()) { | if (!shader_transparent_shadow(NULL, isect) || optixGetPayload_2() >= optixGetPayload_3()) { | ||||
| # endif | # endif | ||||
| // This is an opaque hit or the hit limit has been reached, abort traversal | // This is an opaque hit or the hit limit has been reached, abort traversal | ||||
| optixSetPayload_5(true); | optixSetPayload_5(true); | ||||
| return optixTerminateRay(); | return optixTerminateRay(); | ||||
| # ifdef __TRANSPARENT_SHADOWS__ | # ifdef __TRANSPARENT_SHADOWS__ | ||||
| } | } | ||||
| // TODO(pmours): Do we need REQUIRE_UNIQUE_ANYHIT for this to work? | |||||
| optixSetPayload_2(optixGetPayload_2() + 1); // num_hits++ | optixSetPayload_2(optixGetPayload_2() + 1); // num_hits++ | ||||
| // Continue tracing | // Continue tracing | ||||
| optixIgnoreIntersection(); | optixIgnoreIntersection(); | ||||
| # endif | # endif | ||||
| #endif | #endif | ||||
| } | } | ||||
| extern "C" __global__ void __anyhit__kernel_optix_visibility_test() | extern "C" __global__ void __anyhit__kernel_optix_visibility_test() | ||||
| { | { | ||||
| uint visibility = optixGetPayload_4(); | uint visibility = optixGetPayload_4(); | ||||
| #ifdef __VISIBILITY_FLAG__ | #ifdef __VISIBILITY_FLAG__ | ||||
| const uint prim = optixGetPrimitiveIndex(); | const uint prim = optixGetPrimitiveIndex(); | ||||
| if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) | if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) { | ||||
| return optixIgnoreIntersection(); | |||||
| } | |||||
| #endif | |||||
| #ifdef __HAIR__ | |||||
| if (!optixIsTriangleHit()) { | |||||
| // Filter out curve endcaps | |||||
| const float u = __uint_as_float(optixGetAttribute_0()); | |||||
| if (u == 0.0f || u == 1.0f) { | |||||
| return optixIgnoreIntersection(); | return optixIgnoreIntersection(); | ||||
| } | |||||
| } | |||||
| #endif | #endif | ||||
| // Shadow ray early termination | // Shadow ray early termination | ||||
| if (visibility & PATH_RAY_SHADOW_OPAQUE) | if (visibility & PATH_RAY_SHADOW_OPAQUE) { | ||||
| return optixTerminateRay(); | return optixTerminateRay(); | ||||
| } | } | ||||
| } | |||||
| extern "C" __global__ void __closesthit__kernel_optix_hit() | extern "C" __global__ void __closesthit__kernel_optix_hit() | ||||
| { | { | ||||
| optixSetPayload_0(__float_as_uint(optixGetRayTmax())); // Intersection distance | optixSetPayload_0(__float_as_uint(optixGetRayTmax())); // Intersection distance | ||||
| optixSetPayload_3(optixGetPrimitiveIndex()); | optixSetPayload_3(optixGetPrimitiveIndex()); | ||||
| optixSetPayload_4(get_object_id()); | optixSetPayload_4(get_object_id()); | ||||
| // Can be PRIMITIVE_TRIANGLE and PRIMITIVE_MOTION_TRIANGLE or curve type and segment index | // Can be PRIMITIVE_TRIANGLE and PRIMITIVE_MOTION_TRIANGLE or curve type and segment index | ||||
| optixSetPayload_5(kernel_tex_fetch(__prim_type, optixGetPrimitiveIndex())); | optixSetPayload_5(kernel_tex_fetch(__prim_type, optixGetPrimitiveIndex())); | ||||
| if (optixIsTriangleHit()) { | if (optixIsTriangleHit()) { | ||||
| const float2 barycentrics = optixGetTriangleBarycentrics(); | const float2 barycentrics = optixGetTriangleBarycentrics(); | ||||
| optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x)); | optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x)); | ||||
| optixSetPayload_2(__float_as_uint(barycentrics.x)); | optixSetPayload_2(__float_as_uint(barycentrics.x)); | ||||
| } | } | ||||
| else { | else { | ||||
| optixSetPayload_1(optixGetAttribute_0()); | optixSetPayload_1(optixGetAttribute_0()); // Same as 'optixGetCurveParameter()' | ||||
| optixSetPayload_2(optixGetAttribute_1()); | optixSetPayload_2(optixGetAttribute_1()); | ||||
| } | } | ||||
| } | } | ||||
| #ifdef __HAIR__ | #ifdef __HAIR__ | ||||
| ccl_device_inline void optix_intersection_curve(const uint prim, const uint type) | ccl_device_inline void optix_intersection_curve(const uint prim, const uint type) | ||||
| { | { | ||||
| const uint object = get_object_id<true>(); | const uint object = get_object_id<true>(); | ||||
| Show All 19 Lines | if (isect.t != FLT_MAX) | ||||
| isect.t *= len; | isect.t *= len; | ||||
| if (curve_intersect(NULL, &isect, P, dir, visibility, object, prim, time, type)) { | if (curve_intersect(NULL, &isect, P, dir, visibility, object, prim, time, type)) { | ||||
| optixReportIntersection(isect.t / len, | optixReportIntersection(isect.t / len, | ||||
| type & PRIMITIVE_ALL, | type & PRIMITIVE_ALL, | ||||
| __float_as_int(isect.u), // Attribute_0 | __float_as_int(isect.u), // Attribute_0 | ||||
| __float_as_int(isect.v)); // Attribute_1 | __float_as_int(isect.v)); // Attribute_1 | ||||
| } | } | ||||
| } | } | ||||
| extern "C" __global__ void __intersection__curve_ribbon() | extern "C" __global__ void __intersection__curve_ribbon() | ||||
| { | { | ||||
| const uint prim = optixGetPrimitiveIndex(); | const uint prim = optixGetPrimitiveIndex(); | ||||
| const uint type = kernel_tex_fetch(__prim_type, prim); | const uint type = kernel_tex_fetch(__prim_type, prim); | ||||
| if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) { | if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) { | ||||
| Show All 18 Lines | |||||