Differential D13361 Diff 45286 intern/cycles/kernel/device/optix/kernel.cu

Changeset View

Standalone View

intern/cycles/kernel/device/optix/kernel.cu

Show All 15 Lines
*/		*/

// 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/tables.h"

#include "kernel/integrator/state.h"		#include "kernel/integrator/state.h"
#include "kernel/integrator/state_flow.h"		#include "kernel/integrator/state_flow.h"
#include "kernel/integrator/state_util.h"		#include "kernel/integrator/state_util.h"

#include "kernel/integrator/intersect_closest.h"		#include "kernel/integrator/intersect_closest.h"
#include "kernel/integrator/intersect_shadow.h"		#include "kernel/integrator/intersect_shadow.h"
#include "kernel/integrator/intersect_subsurface.h"		#include "kernel/integrator/intersect_subsurface.h"
#include "kernel/integrator/intersect_volume_stack.h"		#include "kernel/integrator/intersect_volume_stack.h"

// clang-format on		// clang-format on

		#define OPTIX_DEFINE_ABI_VERSION_ONLY
		#include <optix_function_table.h>

template<typename T> ccl_device_forceinline T *get_payload_ptr_0()		template<typename T> ccl_device_forceinline T *get_payload_ptr_0()
{		{
return pointer_unpack_from_uint<T>(optixGetPayload_0(), optixGetPayload_1());		return pointer_unpack_from_uint<T>(optixGetPayload_0(), optixGetPayload_1());
}		}
template<typename T> ccl_device_forceinline T *get_payload_ptr_2()		template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
{		{
return pointer_unpack_from_uint<T>(optixGetPayload_2(), optixGetPayload_3());		return pointer_unpack_from_uint<T>(optixGetPayload_2(), optixGetPayload_3());
}		}

ccl_device_forceinline int get_object_id()		ccl_device_forceinline int get_object_id()
{		{
#ifdef __OBJECT_MOTION__		#ifdef __OBJECT_MOTION__
/* Always get the the instance ID from the TLAS		/* Always get 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. */
return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));		return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
#else		#else
return optixGetInstanceId();		return optixGetInstanceId();
#endif		#endif
}		}

extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()		extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()
▲ Show 20 Lines • Show All 137 Lines • ▼ Show 20 Lines	# 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());

const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);		const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
type = segment.type;		type = segment.type;
prim = segment.prim;		prim = segment.prim;

		# if OPTIX_ABI_VERSION < 55
/* Filter out curve endcaps. */		/* Filter out curve endcaps. */
if (u == 0.0f \|\| u == 1.0f) {		if (u == 0.0f \|\| u == 1.0f) {
return optixIgnoreIntersection();		return optixIgnoreIntersection();
}		}
		# endif
}		}
# endif		# endif

# ifndef __TRANSPARENT_SHADOWS__		# ifndef __TRANSPARENT_SHADOWS__
/* No transparent shadows support compiled in, make opaque. */		/* No transparent shadows support compiled in, make opaque. */
optixSetPayload_5(true);		optixSetPayload_5(true);
return optixTerminateRay();		return optixTerminateRay();
# else		# else
▲ Show 20 Lines • Show All 90 Lines • ▼ Show 20 Lines	#endif
if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {		if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
return optixIgnoreIntersection();		return optixIgnoreIntersection();
}		}
}		}

extern "C" __global__ void __anyhit__kernel_optix_visibility_test()		extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
{		{
#ifdef __HAIR__		#ifdef __HAIR__
		# if OPTIX_ABI_VERSION < 55
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
		#endif

#ifdef __VISIBILITY_FLAG__		#ifdef __VISIBILITY_FLAG__
const uint object = get_object_id();		const uint object = get_object_id();
const uint visibility = optixGetPayload_4();		const uint visibility = optixGetPayload_4();
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {		if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
return optixIgnoreIntersection();		return optixIgnoreIntersection();
}		}

▲ Show 20 Lines • Show All 80 Lines • Show Last 20 Lines