Changeset View
Changeset View
Standalone View
Standalone View
intern/cycles/kernel/device/optix/kernel.cu
| Show All 13 Lines | |||||
| * 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 | // clang-format off | ||||
| #include "kernel/device/optix/compat.h" | #include "kernel/device/optix/compat.h" | ||||
| #include "kernel/device/optix/globals.h" | #include "kernel/device/optix/globals.h" | ||||
| #include "kernel/device/gpu/image.h" // Texture lookup uses normal CUDA intrinsics | #include "kernel/device/gpu/image.h" /* Texture lookup uses normal CUDA intrinsics. */ | ||||
| #include "kernel/integrator/integrator_state.h" | #include "kernel/integrator/integrator_state.h" | ||||
| #include "kernel/integrator/integrator_state_flow.h" | #include "kernel/integrator/integrator_state_flow.h" | ||||
| #include "kernel/integrator/integrator_state_util.h" | #include "kernel/integrator/integrator_state_util.h" | ||||
| #include "kernel/integrator/integrator_intersect_closest.h" | #include "kernel/integrator/integrator_intersect_closest.h" | ||||
| #include "kernel/integrator/integrator_intersect_shadow.h" | #include "kernel/integrator/integrator_intersect_shadow.h" | ||||
| #include "kernel/integrator/integrator_intersect_subsurface.h" | #include "kernel/integrator/integrator_intersect_subsurface.h" | ||||
| #include "kernel/integrator/integrator_intersect_volume_stack.h" | #include "kernel/integrator/integrator_intersect_volume_stack.h" | ||||
| // clang-format on | // 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()); | ||||
| } | } | ||||
| template<bool always = false> ccl_device_forceinline uint get_object_id() | template<bool always = false> ccl_device_forceinline uint get_object_id() | ||||
| { | { | ||||
| #ifdef __OBJECT_MOTION__ | #ifdef __OBJECT_MOTION__ | ||||
| // Always get the the instance ID from the TLAS | /* Always get the the instance ID from the TLAS. | ||||
| // There might be a motion transform node between TLAS and BLAS which does not have one | * There might be a motion transform node between TLAS and BLAS which does not have one. */ | ||||
| uint object = optixGetInstanceIdFromHandle(optixGetTransformListHandle(0)); | uint object = optixGetInstanceIdFromHandle(optixGetTransformListHandle(0)); | ||||
| #else | #else | ||||
| uint object = optixGetInstanceId(); | uint object = optixGetInstanceId(); | ||||
| #endif | #endif | ||||
| // Choose between always returning object ID or only for instances | /* Choose between always returning object ID or only for instances. */ | ||||
| if (always || (object & 1) == 0) | if (always || (object & 1) == 0) | ||||
| // Can just remove the low bit since instance always contains object ID | /* Can just remove the low bit since instance always contains object ID. */ | ||||
| return object >> 1; | return object >> 1; | ||||
| else | else | ||||
| // Set to OBJECT_NONE if this is not an instanced object | /* Set to OBJECT_NONE if this is not an instanced object. */ | ||||
| return OBJECT_NONE; | return OBJECT_NONE; | ||||
| } | } | ||||
| extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest() | extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest() | ||||
| { | { | ||||
| const int global_index = optixGetLaunchIndex().x; | const int global_index = optixGetLaunchIndex().x; | ||||
| const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] : | const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] : | ||||
| global_index; | global_index; | ||||
| Show All 21 Lines | extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_volume_stack() | ||||
| const int global_index = optixGetLaunchIndex().x; | const int global_index = optixGetLaunchIndex().x; | ||||
| const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] : | const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] : | ||||
| global_index; | global_index; | ||||
| integrator_intersect_volume_stack(nullptr, path_index); | integrator_intersect_volume_stack(nullptr, path_index); | ||||
| } | } | ||||
| extern "C" __global__ void __miss__kernel_optix_miss() | extern "C" __global__ void __miss__kernel_optix_miss() | ||||
| { | { | ||||
| // 'kernel_path_lamp_emission' checks intersection distance, so need to set it even on a miss | /* 'kernel_path_lamp_emission' checks intersection distance, so need to set it even on a miss. */ | ||||
| optixSetPayload_0(__float_as_uint(optixGetRayTmax())); | optixSetPayload_0(__float_as_uint(optixGetRayTmax())); | ||||
| optixSetPayload_5(PRIMITIVE_NONE); | optixSetPayload_5(PRIMITIVE_NONE); | ||||
| } | } | ||||
| extern "C" __global__ void __anyhit__kernel_optix_local_hit() | extern "C" __global__ void __anyhit__kernel_optix_local_hit() | ||||
| { | { | ||||
| #ifdef __HAIR__ | |||||
| if (!optixIsTriangleHit()) { | |||||
| /* Ignore curves. */ | |||||
| return optixIgnoreIntersection(); | |||||
| } | |||||
| #endif | |||||
| #ifdef __BVH_LOCAL__ | #ifdef __BVH_LOCAL__ | ||||
| const uint object = get_object_id<true>(); | const uint object = get_object_id<true>(); | ||||
| if (object != optixGetPayload_4() /* local_object */) { | if (object != optixGetPayload_4() /* local_object */) { | ||||
| // Only intersect with matching object | /* Only intersect with matching object. */ | ||||
| return optixIgnoreIntersection(); | return optixIgnoreIntersection(); | ||||
| } | } | ||||
| const uint max_hits = optixGetPayload_5(); | const uint max_hits = optixGetPayload_5(); | ||||
| if (max_hits == 0) { | if (max_hits == 0) { | ||||
| // Special case for when no hit information is requested, just report that something was hit | /* Special case for when no hit information is requested, just report that something was hit */ | ||||
| optixSetPayload_5(true); | optixSetPayload_5(true); | ||||
| return optixTerminateRay(); | return optixTerminateRay(); | ||||
| } | } | ||||
| int hit = 0; | int hit = 0; | ||||
| uint *const lcg_state = get_payload_ptr_0<uint>(); | uint *const lcg_state = get_payload_ptr_0<uint>(); | ||||
| LocalIntersection *const local_isect = get_payload_ptr_2<LocalIntersection>(); | LocalIntersection *const local_isect = get_payload_ptr_2<LocalIntersection>(); | ||||
| Show All 10 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 | /* Record closest intersection only. | ||||
| // Do not terminate ray here, since there is no guarantee about distance ordering in any-hit | * 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); | ||||
| 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; | ||||
| // 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)); | ||||
| // Continue tracing (without this the trace call would return after the first hit) | /* Continue tracing (without this the trace call would return after the first hit). */ | ||||
| optixIgnoreIntersection(); | optixIgnoreIntersection(); | ||||
| #endif | #endif | ||||
| } | } | ||||
| extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() | extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() | ||||
| { | { | ||||
| #ifdef __SHADOW_RECORD_ALL__ | #ifdef __SHADOW_RECORD_ALL__ | ||||
| bool ignore_intersection = false; | bool ignore_intersection = false; | ||||
| Show All 12 Lines | if (optixIsTriangleHit()) { | ||||
| u = 1.0f - barycentrics.y - barycentrics.x; | u = 1.0f - barycentrics.y - barycentrics.x; | ||||
| v = barycentrics.x; | v = barycentrics.x; | ||||
| } | } | ||||
| # ifdef __HAIR__ | # ifdef __HAIR__ | ||||
| else { | else { | ||||
| u = __uint_as_float(optixGetAttribute_0()); | u = __uint_as_float(optixGetAttribute_0()); | ||||
| v = __uint_as_float(optixGetAttribute_1()); | v = __uint_as_float(optixGetAttribute_1()); | ||||
| // Filter out curve endcaps | /* Filter out curve endcaps. */ | ||||
| if (u == 0.0f || u == 1.0f) { | if (u == 0.0f || u == 1.0f) { | ||||
| ignore_intersection = true; | ignore_intersection = true; | ||||
| } | } | ||||
| } | } | ||||
| # endif | # endif | ||||
| int num_hits = optixGetPayload_2(); | int num_hits = optixGetPayload_2(); | ||||
| int record_index = num_hits; | int record_index = num_hits; | ||||
| Show All 34 Lines | if (!ignore_intersection) { | ||||
| isect->u = u; | isect->u = u; | ||||
| isect->v = v; | isect->v = v; | ||||
| isect->t = optixGetRayTmax(); | isect->t = optixGetRayTmax(); | ||||
| isect->prim = prim; | isect->prim = prim; | ||||
| 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); | ||||
| # 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) || max_hits == 0) { | if (!shader_transparent_shadow(NULL, isect) || max_hits == 0) { | ||||
| # endif | # endif | ||||
| // If no transparent shadows, all light is blocked and we can stop immediately | /* If no transparent shadows, all light is blocked and we can stop immediately. */ | ||||
| optixSetPayload_5(true); | optixSetPayload_5(true); | ||||
| return optixTerminateRay(); | return optixTerminateRay(); | ||||
| # ifdef __TRANSPARENT_SHADOWS__ | # ifdef __TRANSPARENT_SHADOWS__ | ||||
| } | } | ||||
| # endif | # endif | ||||
| } | } | ||||
| // Continue tracing | /* Continue tracing. */ | ||||
| optixIgnoreIntersection(); | optixIgnoreIntersection(); | ||||
| #endif | #endif | ||||
| } | } | ||||
| extern "C" __global__ void __anyhit__kernel_optix_visibility_test() | extern "C" __global__ void __anyhit__kernel_optix_volume_test() | ||||
| { | { | ||||
| uint visibility = optixGetPayload_4(); | #ifdef __HAIR__ | ||||
| if (!optixIsTriangleHit()) { | |||||
| /* Ignore curves. */ | |||||
| return optixIgnoreIntersection(); | |||||
| } | |||||
| #endif | |||||
| #ifdef __VISIBILITY_FLAG__ | #ifdef __VISIBILITY_FLAG__ | ||||
| const uint prim = optixGetPrimitiveIndex(); | const uint prim = optixGetPrimitiveIndex(); | ||||
| const uint visibility = optixGetPayload_4(); | |||||
| if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) { | if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) { | ||||
| return optixIgnoreIntersection(); | return optixIgnoreIntersection(); | ||||
| } | } | ||||
| #endif | #endif | ||||
| const uint object = get_object_id<true>(); | |||||
| if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) { | |||||
| return optixIgnoreIntersection(); | |||||
| } | |||||
| } | |||||
| extern "C" __global__ void __anyhit__kernel_optix_visibility_test() | |||||
| { | |||||
| #ifdef __HAIR__ | #ifdef __HAIR__ | ||||
| if (!optixIsTriangleHit()) { | if (!optixIsTriangleHit()) { | ||||
| // Filter out curve endcaps | /* Filter out curve endcaps. */ | ||||
| const float u = __uint_as_float(optixGetAttribute_0()); | const float u = __uint_as_float(optixGetAttribute_0()); | ||||
| if (u == 0.0f || u == 1.0f) { | if (u == 0.0f || u == 1.0f) { | ||||
| return optixIgnoreIntersection(); | return optixIgnoreIntersection(); | ||||
| } | } | ||||
| } | } | ||||
| #endif | #endif | ||||
| // Shadow ray early termination | #ifdef __VISIBILITY_FLAG__ | ||||
| const uint prim = optixGetPrimitiveIndex(); | |||||
| const uint visibility = optixGetPayload_4(); | |||||
| if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) { | |||||
| return optixIgnoreIntersection(); | |||||
| } | |||||
| /* Shadow ray early termination. */ | |||||
| if (visibility & PATH_RAY_SHADOW_OPAQUE) { | if (visibility & PATH_RAY_SHADOW_OPAQUE) { | ||||
| return optixTerminateRay(); | return optixTerminateRay(); | ||||
| } | } | ||||
| #endif | |||||
| } | } | ||||
| 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()); // Same as 'optixGetCurveParameter()' | 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>(); | ||||
| const uint visibility = optixGetPayload_4(); | const uint visibility = optixGetPayload_4(); | ||||
| float3 P = optixGetObjectRayOrigin(); | float3 P = optixGetObjectRayOrigin(); | ||||
| float3 dir = optixGetObjectRayDirection(); | float3 dir = optixGetObjectRayDirection(); | ||||
| // The direction is not normalized by default, but the curve intersection routine expects that | /* The direction is not normalized by default, but the curve intersection routine expects that */ | ||||
| float len; | float len; | ||||
| dir = normalize_len(dir, &len); | dir = normalize_len(dir, &len); | ||||
| # ifdef __OBJECT_MOTION__ | # ifdef __OBJECT_MOTION__ | ||||
| const float time = optixGetRayTime(); | const float time = optixGetRayTime(); | ||||
| # else | # else | ||||
| const float time = 0.0f; | const float time = 0.0f; | ||||
| # endif | # endif | ||||
| Intersection isect; | Intersection isect; | ||||
| isect.t = optixGetRayTmax(); | isect.t = optixGetRayTmax(); | ||||
| // Transform maximum distance into object space | /* Transform maximum distance into object space. */ | ||||
| if (isect.t != FLT_MAX) | if (isect.t != FLT_MAX) | ||||
| isect.t *= len; | isect.t *= len; | ||||
| if (curve_intersect(NULL, &isect, P, dir, isect.t, visibility, object, prim, time, type)) { | if (curve_intersect(NULL, &isect, P, dir, isect.t, 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)) { | ||||
| optix_intersection_curve(prim, type); | optix_intersection_curve(prim, type); | ||||
| } | } | ||||
| } | } | ||||
| #endif | #endif | ||||