Differential D12880 Diff 43509 intern/cycles/kernel/device/optix/kernel.cu

Changeset View

Standalone View

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

Show First 20 Lines • Show All 204 Lines • ▼ Show 20 Lines	# ifdef __HAIR__
}		}
# 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_4(true);		optixSetPayload_4(true);
return optixTerminateRay();		return optixTerminateRay();
# else		# else
const int max_hits = optixGetPayload_2();		const uint max_hits = optixGetPayload_2();
		const uint num_hits_packed = optixGetPayload_1();
		const uint num_recorded_hits = uint16_unpack_from_uint_0(num_hits_packed);
		const uint num_hits = uint16_unpack_from_uint_1(num_hits_packed);

/* 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. */
if (max_hits == 0 \|\|		if (num_hits >= max_hits \|\|
!(intersection_get_shader_flags(NULL, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) {		!(intersection_get_shader_flags(NULL, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) {
optixSetPayload_4(true);		optixSetPayload_4(true);
return optixTerminateRay();		return optixTerminateRay();
}		}

		/* Always use baked shadow transparency for curves. */
		if (type & PRIMITIVE_ALL_CURVE) {
		float throughput = __uint_as_float(optixGetPayload_5());
		throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, u);
		optixSetPayload_5(__float_as_uint(throughput));
		optixSetPayload_1(uint16_pack_to_uint(num_recorded_hits, num_hits + 1));

		if (throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) {
		optixSetPayload_4(true);
		return optixTerminateRay();
		}
		else {
		/* Continue tracing. */
		optixIgnoreIntersection();
		return;
		}
		}

/* Record transparent intersection. */		/* Record transparent intersection. */
const int num_hits = optixGetPayload_1();		optixSetPayload_1(uint16_pack_to_uint(num_recorded_hits + 1, num_hits + 1));
int record_index = num_hits;

optixSetPayload_1(num_hits + 1);		uint record_index = num_recorded_hits;

const IntegratorShadowState state = optixGetPayload_0();		const IntegratorShadowState state = optixGetPayload_0();

if (record_index >= max_hits) {		const uint max_record_hits = min(max_hits, INTEGRATOR_SHADOW_ISECT_SIZE);
		if (record_index >= max_record_hits) {
/* If maximum number of hits reached, find a hit to replace. */		/* If maximum number of hits reached, find a hit to replace. */
float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t);		float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t);
int max_recorded_hit = 0;		uint max_recorded_hit = 0;

for (int i = 1; i < max_hits; i++) {		for (int i = 1; i < max_record_hits; i++) {
const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t);		const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t);
if (isect_t > max_recorded_t) {		if (isect_t > max_recorded_t) {
max_recorded_t = isect_t;		max_recorded_t = isect_t;
max_recorded_hit = i;		max_recorded_hit = i;
}		}
}		}

if (optixGetRayTmax() >= max_recorded_t) {		if (optixGetRayTmax() >= max_recorded_t) {
/* Accept hit, so that OptiX won't consider any more hits beyond the distance of the		/* Accept hit, so that OptiX won't consider any more hits beyond the distance of the
* current hit anymore. */		* current hit anymore. */
return;		return;
}		}

record_index = max_recorded_hit;		record_index = max_recorded_hit;
}		}

INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, u) = u;		INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, u) = u;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, v) = v;		INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, v) = v;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, t) = optixGetRayTmax();		INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, t) = optixGetRayTmax();
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, prim) = prim;		INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, prim) = prim;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object;		INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type;		INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type;

		/* Continue tracing. */
optixIgnoreIntersection();		optixIgnoreIntersection();
# endif /* __TRANSPARENT_SHADOWS__ */		# endif /* __TRANSPARENT_SHADOWS__ */
#endif /* __SHADOW_RECORD_ALL__ */		#endif /* __SHADOW_RECORD_ALL__ */
}		}

extern "C" __global__ void __anyhit__kernel_optix_volume_test()		extern "C" __global__ void __anyhit__kernel_optix_volume_test()
{		{
#ifdef __HAIR__		#ifdef __HAIR__
▲ Show 20 Lines • Show All 118 Lines • Show Last 20 Lines