bartvdbraak/blender
1
0
Fork 4
Code Issues 4 Pull Requests 1 Actions 1 Packages Projects Releases Wiki Activity

Cycles: Dormant fixes for adaptive feature compilation

Browse Source 
This PR fixes the (currently unused) scene-based selective feature compilation macros. These feature based macros haven't been used for a few years, and enabling them currently results in compilation errors.

The only functional change in this PR is in geom/primitive.h where undef-ing `__HAIR__` had exposed an inconsistency in how pointcloud attributes were being fetched. Using the more general `primitive_surface_attribute_float4` (instead of `curve_attribute_float4`) fixed a compilation error that occurred when rendering pointcloud unit test scenes with adaptive compilation enabled.

Pull Request: https://projects.blender.org/blender/blender/pulls/121216
This commit is contained in:
Michael Jones 2024-04-30 12:56:22 +02:00 committed by Michael Jones (Apple)
parent 33c6e9f92c
commit 99f5433445
19 changed files with 292 additions and 214 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
intern/cycles/kernel/bake/bake.h
View File

@ -95,6 +95,7 @@ ccl_device void kernel_curve_shadow_transparency_evaluate(
ccl_global float *output,
const int offset)
{
#ifdef __HAIR__
/* Setup shader data. */
const KernelShaderEvalInput in = input[offset];
@ -108,6 +109,7 @@ ccl_device void kernel_curve_shadow_transparency_evaluate(
/* Write output. */
output[offset] = clamp(average(surface_shader_transparency(kg, &sd)), 0.0f, 1.0f);
#endif
}
CCL_NAMESPACE_END

4
intern/cycles/kernel/bvh/shadow_all.h
View File

@ -175,7 +175,7 @@ ccl_device_inline
break;
}
#endif
#if BVH_FEATURE(BVH_HAIR)
#if BVH_FEATURE(BVH_HAIR) && defined(__HAIR__)
case PRIMITIVE_CURVE_THICK:
case PRIMITIVE_MOTION_CURVE_THICK:
case PRIMITIVE_CURVE_RIBBON:
@ -195,7 +195,7 @@ ccl_device_inline
break;
}
#endif
#if BVH_FEATURE(BVH_POINTCLOUD)
#if BVH_FEATURE(BVH_POINTCLOUD) && defined(__POINTCLOUD__)
case PRIMITIVE_POINT:
case PRIMITIVE_MOTION_POINT: {
if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {

4
intern/cycles/kernel/bvh/traversal.h
View File

@ -178,7 +178,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
break;
}
#endif /* BVH_FEATURE(BVH_MOTION) */
#if BVH_FEATURE(BVH_HAIR)
#if BVH_FEATURE(BVH_HAIR) && defined(__HAIR__)
case PRIMITIVE_CURVE_THICK:
case PRIMITIVE_MOTION_CURVE_THICK:
case PRIMITIVE_CURVE_RIBBON:
@ -201,7 +201,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
break;
}
#endif /* BVH_FEATURE(BVH_HAIR) */
#if BVH_FEATURE(BVH_POINTCLOUD)
#if BVH_FEATURE(BVH_POINTCLOUD) && defined(__POINTCLOUD__)
case PRIMITIVE_POINT:
case PRIMITIVE_MOTION_POINT: {
if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {

2
intern/cycles/kernel/device/gpu/kernel.h
View File

@ -190,12 +190,14 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
ccl_global const int *path_index_array,
const int work_size)
{
# ifdef __VOLUME__
const int global_index = ccl_gpu_global_id_x();
if (ccl_gpu_kernel_within_bounds(global_index, work_size)) {
const int state = (path_index_array) ? path_index_array[global_index] : global_index;
ccl_gpu_kernel_call(integrator_intersect_volume_stack(NULL, state));
}
# endif
}
ccl_gpu_kernel_postfix

40
intern/cycles/kernel/device/metal/bvh.h
View File

@ -39,6 +39,7 @@ struct MetalRTIntersectionShadowPayload {
bool result;
};
#ifdef __HAIR__
ccl_device_forceinline bool curve_ribbon_accept(
KernelGlobals kg, float u, float t, ccl_private const Ray *ray, int object, int prim, int type)
{
@ -63,11 +64,11 @@ ccl_device_forceinline bool curve_ribbon_accept(
float3 ray_D = ray->D;
if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
float3 idir;
#if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
bvh_instance_motion_push(NULL, object, ray, &ray_P, &ray_D, &idir);
#else
# else
bvh_instance_push(NULL, object, ray, &ray_P, &ray_D, &idir);
#endif
# endif
}
/* ignore self intersections */
@ -78,11 +79,11 @@ ccl_device_forceinline bool curve_ribbon_accept(
ccl_device_forceinline float curve_ribbon_v(
KernelGlobals kg, float u, float t, ccl_private const Ray *ray, int object, int prim, int type)
{
#if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
float time = ray->time;
#else
# else
float time = 0.0f;
#endif
# endif
const bool is_motion = (type & PRIMITIVE_MOTION);
@ -108,11 +109,11 @@ ccl_device_forceinline float curve_ribbon_v(
float3 ray_D = ray->D;
if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
float3 idir;
#if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
bvh_instance_motion_push(NULL, object, ray, &ray_P, &ray_D, &idir);
#else
# else
bvh_instance_push(NULL, object, ray, &ray_P, &ray_D, &idir);
#endif
# endif
}
const float4 P_curve4 = metal::catmull_rom(u, curve[0], curve[1], curve[2], curve[3]);
@ -130,6 +131,7 @@ ccl_device_forceinline float curve_ribbon_v(
float v = dot(P - P_curve, bitangent) / r_curve;
return clamp(v, -1.0, 1.0f);
}
#endif /* __HAIR__ */
/* Scene intersection. */
@ -207,6 +209,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
isect->u = intersection.triangle_barycentric_coord.x;
isect->v = intersection.triangle_barycentric_coord.y;
}
#ifdef __HAIR__
else if (kernel_data.bvh.have_curves && intersection.type == intersection_type::curve) {
int prim = intersection.primitive_id + intersection.user_instance_id;
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
@ -227,6 +230,8 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
isect->v = 0.0f;
}
}
#endif /* __HAIR__ */
#ifdef __POINTCLOUD__
else if (kernel_data.bvh.have_points && intersection.type == intersection_type::bounding_box) {
const int object = intersection.instance_id;
const uint prim = intersection.primitive_id + intersection.user_instance_id;
@ -234,11 +239,11 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
float3 idir;
#if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
bvh_instance_motion_push(NULL, object, ray, &r.origin, &r.direction, &idir);
#else
# else
bvh_instance_push(NULL, object, ray, &r.origin, &r.direction, &idir);
#endif
# endif
}
if (prim_type & PRIMITIVE_POINT) {
@ -262,6 +267,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
return true;
}
}
#endif /* __POINTCLOUD__ */
return true;
}
@ -497,6 +503,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
isect->object = intersection.instance_id;
isect->t = intersection.distance;
}
# ifdef __HAIR__
else if (kernel_data.bvh.have_curves && intersection.type == intersection_type::curve) {
int prim = intersection.primitive_id + intersection.user_instance_id;
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
@ -517,6 +524,8 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
isect->v = 0.0f;
}
}
# endif
# ifdef __POINTCLOUD__
else if (kernel_data.bvh.have_points && intersection.type == intersection_type::bounding_box) {
const int object = intersection.instance_id;
const uint prim = intersection.primitive_id + intersection.user_instance_id;
@ -526,11 +535,11 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
float3 idir;
# if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
bvh_instance_motion_push(NULL, object, ray, &r.origin, &r.direction, &idir);
# else
# else
bvh_instance_push(NULL, object, ray, &r.origin, &r.direction, &idir);
# endif
# endif
}
if (prim_type & PRIMITIVE_POINT) {
@ -552,6 +561,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
return true;
}
}
# endif
return true;
}

40
intern/cycles/kernel/device/metal/kernel.metal
View File

@ -534,7 +534,6 @@ __anyhit__cycles_metalrt_visibility_test_box(const float ray_tmax [[max_distance
/* Primitive intersection functions. */
# ifdef __HAIR__
[[intersection(
curve, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
__intersection__curve(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
@ -548,10 +547,10 @@ __intersection__curve(constant KernelParamsMetal &launch_params_metal [[buffer(1
float u [[curve_parameter]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]]
# if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
,
const float time [[time]]
# endif
# endif
)
{
uint prim = primitive_id + primitive_id_offset;
@ -559,9 +558,9 @@ __intersection__curve(constant KernelParamsMetal &launch_params_metal [[buffer(1
Ray ray;
ray.P = ray_P;
ray.D = ray_D;
# if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
ray.time = time;
# endif
# endif
PrimitiveIntersectionResult result =
metalrt_visibility_test<PrimitiveIntersectionResult, METALRT_HIT_CURVE>(
@ -582,9 +581,9 @@ __intersection__curve_shadow(constant KernelParamsMetal &launch_params_metal [[b
const float3 ray_D [[direction]],
float u [[curve_parameter]],
float t [[distance]],
# if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
const float time [[time]],
# endif
# endif
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
@ -595,9 +594,9 @@ __intersection__curve_shadow(constant KernelParamsMetal &launch_params_metal [[b
Ray ray;
ray.P = ray_P;
ray.D = ray_D;
# if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
ray.time = time;
# endif
# endif
result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_CURVE>(
launch_params_metal, payload, object, prim, float2(u, 0), ray_tmax, t, &ray);
@ -605,7 +604,6 @@ __intersection__curve_shadow(constant KernelParamsMetal &launch_params_metal [[b
return result;
}
# endif /* __HAIR__ */
# ifdef __POINTCLOUD__
ccl_device_inline void metalrt_intersection_point(
@ -666,6 +664,8 @@ ccl_device_inline void metalrt_intersection_point_shadow(
}
}
# endif /* __POINTCLOUD__ */
[[intersection(bounding_box,
triangle_data,
curve_data,
@ -678,9 +678,9 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1
const uint primitive_id_offset [[user_instance_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
# if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
const float time [[time]],
# endif
# endif
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
@ -692,6 +692,8 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1
result.continue_search = true;
result.distance = ray_tmax;
# ifdef __POINTCLOUD__
metalrt_intersection_point(launch_params_metal,
payload,
object,
@ -708,6 +710,8 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1
ray_tmax,
result);
# endif /* __POINTCLOUD__ */
return result;
}
@ -724,9 +728,9 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b
const uint primitive_id_offset [[user_instance_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
# if defined(__METALRT_MOTION__)
# if defined(__METALRT_MOTION__)
const float time [[time]],
# endif
# endif
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
@ -738,6 +742,8 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b
result.continue_search = true;
result.distance = ray_tmax;
# ifdef __POINTCLOUD__
metalrt_intersection_point_shadow(launch_params_metal,
payload,
object,
@ -754,7 +760,9 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b
ray_tmax,
result);
# endif /* __POINTCLOUD__ */
return result;
}
# endif /* __POINTCLOUD__ */
#endif /* __METALRT__ */
#endif /* __METALRT__ */

2
intern/cycles/kernel/film/light_passes.h
View File

@ -571,7 +571,9 @@ ccl_device_inline void film_write_transparent(KernelGlobals kg,
film_write_pass_float(buffer + kernel_data.film.pass_combined + 3, transparent);
}
#ifdef __SHADOW_CATCHER__
film_write_shadow_catcher_transparent_only(kg, path_flag, transparent, buffer);
#endif
}
/* Write holdout to render buffer. */

11
intern/cycles/kernel/geom/curve.h
View File

@ -202,13 +202,16 @@ ccl_device float curve_thickness(KernelGlobals kg, ccl_private const ShaderData
float4 P_curve[2];
if (!(sd->type & PRIMITIVE_MOTION)) {
# ifdef __OBJECT_MOTION__
if (sd->type & PRIMITIVE_MOTION) {
motion_curve_keys_linear(kg, sd->object, sd->time, k0, k1, P_curve);
}
else
# endif
{
P_curve[0] = kernel_data_fetch(curve_keys, k0);
P_curve[1] = kernel_data_fetch(curve_keys, k1);
}
else {
motion_curve_keys_linear(kg, sd->object, sd->time, k0, k1, P_curve);
}
r = (P_curve[1].w - P_curve[0].w) * sd->u + P_curve[0].w;
}

7
intern/cycles/kernel/geom/object.h
View File

@ -75,12 +75,14 @@ ccl_device_inline Transform object_fetch_transform_motion(KernelGlobals kg, int
return tfm;
}
#endif /* __OBJECT_MOTION__ */
ccl_device_inline Transform object_fetch_transform_motion_test(KernelGlobals kg,
int object,
float time,
ccl_private Transform *itfm)
{
#ifdef __OBJECT_MOTION__
int object_flag = kernel_data_fetch(object_flag, object);
if (object_flag & SD_OBJECT_MOTION) {
/* if we do motion blur */
@ -91,7 +93,9 @@ ccl_device_inline Transform object_fetch_transform_motion_test(KernelGlobals kg,
return tfm;
}
else {
else
#endif /* __OBJECT_MOTION__ */
{
Transform tfm = object_fetch_transform(kg, object, OBJECT_TRANSFORM);
if (itfm)
*itfm = object_fetch_transform(kg, object, OBJECT_INVERSE_TRANSFORM);
@ -99,7 +103,6 @@ ccl_device_inline Transform object_fetch_transform_motion_test(KernelGlobals kg,
return tfm;
}
}
#endif
/* Get transform matrix for shading point. */

4
intern/cycles/kernel/geom/primitive.h
View File

@ -306,9 +306,9 @@ ccl_device_forceinline float4 primitive_motion_vector(KernelGlobals kg,
#if defined(__HAIR__) || defined(__POINTCLOUD__)
if (is_curve_or_point) {
motion_pre = float4_to_float3(curve_attribute_float4(kg, sd, desc, NULL, NULL));
motion_pre = float4_to_float3(primitive_surface_attribute_float4(kg, sd, desc, NULL, NULL));
desc.offset += numverts;
motion_post = float4_to_float3(curve_attribute_float4(kg, sd, desc, NULL, NULL));
motion_post = float4_to_float3(primitive_surface_attribute_float4(kg, sd, desc, NULL, NULL));
/* Curve */
if ((sd->object_flag & SD_OBJECT_HAS_VERTEX_MOTION) == 0) {

26
intern/cycles/kernel/geom/shader_data.h
View File

@ -14,17 +14,17 @@ CCL_NAMESPACE_BEGIN
/* ShaderData setup from incoming ray */
#ifdef __OBJECT_MOTION__
ccl_device void shader_setup_object_transforms(KernelGlobals kg,
ccl_private ShaderData *ccl_restrict sd,
float time)
{
#ifdef __OBJECT_MOTION__
if (sd->object_flag & SD_OBJECT_MOTION) {
sd->ob_tfm_motion = object_fetch_transform_motion(kg, sd->object, time);
sd->ob_itfm_motion = transform_inverse(sd->ob_tfm_motion);
}
}
#endif
}
/* TODO: break this up if it helps reduce register pressure to load data from
* global memory as we write it to shader-data. */
@ -273,6 +273,7 @@ ccl_device void shader_setup_from_displace(KernelGlobals kg,
/* ShaderData setup for point on curve. */
#ifdef __HAIR__
ccl_device void shader_setup_from_curve(KernelGlobals kg,
ccl_private ShaderData *ccl_restrict sd,
int object,
@ -296,9 +297,9 @@ ccl_device void shader_setup_from_curve(KernelGlobals kg,
/* Object */
sd->object = object;
sd->object_flag = kernel_data_fetch(object_flag, sd->object);
#ifdef __OBJECT_MOTION__
# ifdef __OBJECT_MOTION__
shader_setup_object_transforms(kg, sd, sd->time);
#endif
# endif
/* Get control points. */
KernelCurve kcurve = kernel_data_fetch(curves, prim);
@ -317,34 +318,35 @@ ccl_device void shader_setup_from_curve(KernelGlobals kg,
/* Interpolate position and tangent. */
sd->P = float4_to_float3(catmull_rom_basis_derivative(P_curve, sd->u));
#ifdef __DPDU__
# ifdef __DPDU__
sd->dPdu = float4_to_float3(catmull_rom_basis_derivative(P_curve, sd->u));
#endif
# endif
/* Transform into world space */
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
object_position_transform_auto(kg, sd, &sd->P);
#ifdef __DPDU__
# ifdef __DPDU__
object_dir_transform_auto(kg, sd, &sd->dPdu);
#endif
# endif
}
/* No view direction, normals or bitangent. */
sd->wi = zero_float3();
sd->N = zero_float3();
sd->Ng = zero_float3();
#ifdef __DPDU__
# ifdef __DPDU__
sd->dPdv = zero_float3();
#endif
# endif
/* No ray differentials currently. */
#ifdef __RAY_DIFFERENTIALS__
# ifdef __RAY_DIFFERENTIALS__
sd->dP = differential_zero_compact();
sd->dI = differential_zero_compact();
sd->du = differential_zero();
sd->dv = differential_zero();
#endif
# endif
}
#endif /* __HAIR__ */
/* ShaderData setup from ray into background */

2
intern/cycles/kernel/integrator/init_from_bake.h
View File

@ -328,7 +328,9 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg,
integrator_path_init_sorted(kg, state, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE, shader_index);
}
#ifdef __SHADOW_CATCHER__
integrator_split_shadow_catcher(kg, state, &isect, render_buffer);
#endif
}
return true;

8
intern/cycles/kernel/integrator/intersect_closest.h
View File

@ -43,9 +43,11 @@ ccl_device_forceinline bool integrator_intersect_terminate(KernelGlobals kg,
if (shader_flags & (SD_HAS_TRANSPARENT_SHADOW | SD_HAS_EMISSION)) {
INTEGRATOR_STATE_WRITE(state, path, flag) |= PATH_RAY_TERMINATE_AFTER_TRANSPARENT;
}
#ifdef __VOLUME__
else if (!integrator_state_volume_stack_is_empty(kg, state)) {
INTEGRATOR_STATE_WRITE(state, path, flag) |= PATH_RAY_TERMINATE_AFTER_VOLUME;
}
#endif
else {
return true;
}
@ -73,10 +75,12 @@ ccl_device_forceinline bool integrator_intersect_terminate(KernelGlobals kg,
/* Mark path to be terminated right after shader evaluation on the surface. */
INTEGRATOR_STATE_WRITE(state, path, flag) |= PATH_RAY_TERMINATE_ON_NEXT_SURFACE;
}
#ifdef __VOLUME__
else if (!integrator_state_volume_stack_is_empty(kg, state)) {
/* TODO: only do this for emissive volumes. */
INTEGRATOR_STATE_WRITE(state, path, flag) |= PATH_RAY_TERMINATE_IN_NEXT_VOLUME;
}
#endif
else {
return true;
}
@ -127,12 +131,14 @@ ccl_device_forceinline void integrator_split_shadow_catcher(
return;
}
# ifdef __VOLUME__
if (!integrator_state_volume_stack_is_empty(kg, state)) {
/* Volume stack is not empty. Re-init the volume stack to exclude any non-shadow catcher
* objects from it, and then continue shading volume and shadow catcher surface after. */
integrator_path_init(kg, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK);
return;
}
# endif
/* Continue with shading shadow catcher surface. */
const int shader = intersection_get_shader(kg, isect);
@ -189,6 +195,7 @@ template<DeviceKernel current_kernel>
ccl_device_forceinline void integrator_intersect_next_kernel_after_shadow_catcher_background(
KernelGlobals kg, IntegratorState state)
{
# ifdef __VOLUME__
/* Same logic as integrator_split_shadow_catcher, but using NEXT instead of INIT. */
if (!integrator_state_volume_stack_is_empty(kg, state)) {
/* Volume stack is not empty. Re-init the volume stack to exclude any non-shadow catcher
@ -197,6 +204,7 @@ ccl_device_forceinline void integrator_intersect_next_kernel_after_shadow_catche
kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK);
return;
}
# endif
/* Continue with shading shadow catcher surface. */
integrator_intersect_next_kernel_after_shadow_catcher_volume<current_kernel>(kg, state);

19
intern/cycles/kernel/integrator/intersect_volume_stack.h
View File

@ -10,6 +10,7 @@
CCL_NAMESPACE_BEGIN
#ifdef __VOLUME__
ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
IntegratorState state,
const float3 from_P,
@ -37,7 +38,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
const uint32_t visibility = SHADOW_CATCHER_PATH_VISIBILITY(path_flag, PATH_RAY_ALL_VISIBILITY);
#ifdef __VOLUME_RECORD_ALL__
# ifdef __VOLUME_RECORD_ALL__
Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1];
uint num_hits = scene_intersect_volume(kg, &volume_ray, hits, 2 * volume_stack_size, visibility);
if (num_hits > 0) {
@ -54,7 +55,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
volume_stack_enter_exit(kg, state, stack_sd);
}
}
#else
# else
Intersection isect;
int step = 0;
while (step < 2 * volume_stack_size &&
@ -71,7 +72,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
volume_ray.self.prim = isect.prim;
++step;
}
#endif
# endif
}
ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState state)
@ -114,7 +115,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
/* Store to avoid global fetches on every intersection step. */
const uint volume_stack_size = kernel_data.volume_stack_size;
#ifdef __VOLUME_RECORD_ALL__
# ifdef __VOLUME_RECORD_ALL__
Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1];
uint num_hits = scene_intersect_volume(kg, &volume_ray, hits, 2 * volume_stack_size, visibility);
if (num_hits > 0) {
@ -157,7 +158,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
}
}
}
#else
# else
/* CUDA does not support definition of a variable size arrays, so use the maximum possible. */
int enclosed_volumes[MAX_VOLUME_STACK_SIZE];
int step = 0;
@ -211,7 +212,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
volume_ray.self.prim = isect.prim;
++step;
}
#endif
# endif
/* Write terminator. */
const VolumeStack new_entry = {OBJECT_NONE, SHADER_NONE};
@ -222,12 +223,15 @@ ccl_device void integrator_intersect_volume_stack(KernelGlobals kg, IntegratorSt
{
integrator_volume_stack_init(kg, state);
# ifdef __SHADOW_CATCHER__
if (INTEGRATOR_STATE(state, path, flag) & PATH_RAY_SHADOW_CATCHER_PASS) {
/* Volume stack re-init for shadow catcher, continue with shading of hit. */
integrator_intersect_next_kernel_after_shadow_catcher_volume<
DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK>(kg, state);
}
else {
else
# endif
{
/* Volume stack init for camera rays, continue with intersection of camera ray. */
integrator_path_next(kg,
state,
@ -235,5 +239,6 @@ ccl_device void integrator_intersect_volume_stack(KernelGlobals kg, IntegratorSt
DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
}
}
#endif /* __VOLUME__ */
CCL_NAMESPACE_END

4
intern/cycles/kernel/integrator/shade_surface.h
View File

@ -206,8 +206,10 @@ integrate_direct_light_shadow_init_common(KernelGlobals kg,
IntegratorShadowState shadow_state = integrator_shadow_path_init(
kg, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, false);
#ifdef __VOLUME__
/* Copy volume stack and enter/exit volume. */
integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
#endif
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(shadow_state, ray);
@ -644,8 +646,10 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
IntegratorShadowState shadow_state = integrator_shadow_path_init(
kg, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, true);
# ifdef __VOLUME__
/* Copy volume stack and enter/exit volume. */
integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
# endif
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(shadow_state, &ray);

59
intern/cycles/kernel/integrator/state_util.h
View File

@ -115,6 +115,7 @@ ccl_device_forceinline void integrator_state_read_isect(
isect->t = INTEGRATOR_STATE(state, isect, t);
}
#ifdef __VOLUME__
ccl_device_forceinline VolumeStack integrator_state_read_volume_stack(ConstIntegratorState state,
int i)
{
@ -139,34 +140,6 @@ ccl_device_forceinline bool integrator_state_volume_stack_is_empty(KernelGlobals
true;
}
/* Shadow Intersection */
ccl_device_forceinline void integrator_state_write_shadow_isect(
IntegratorShadowState state,
ccl_private const Intersection *ccl_restrict isect,
const int index)
{
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, t) = isect->t;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, u) = isect->u;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, v) = isect->v;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, object) = isect->object;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, prim) = isect->prim;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, type) = isect->type;
}
ccl_device_forceinline void integrator_state_read_shadow_isect(
ConstIntegratorShadowState state,
ccl_private Intersection *ccl_restrict isect,
const int index)
{
isect->prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, prim);
isect->object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, object);
isect->type = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, type);
isect->u = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, u);
isect->v = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, v);
isect->t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, t);
}
ccl_device_forceinline void integrator_state_copy_volume_stack_to_shadow(
KernelGlobals kg, IntegratorShadowState shadow_state, ConstIntegratorState state)
{
@ -228,6 +201,36 @@ ccl_device_forceinline void integrator_state_write_shadow_volume_stack(Integrato
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_volume_stack, i, shader) = entry.shader;
}
#endif /* __VOLUME__*/
/* Shadow Intersection */
ccl_device_forceinline void integrator_state_write_shadow_isect(
IntegratorShadowState state,
ccl_private const Intersection *ccl_restrict isect,
const int index)
{
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, t) = isect->t;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, u) = isect->u;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, v) = isect->v;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, object) = isect->object;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, prim) = isect->prim;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, index, type) = isect->type;
}
ccl_device_forceinline void integrator_state_read_shadow_isect(
ConstIntegratorShadowState state,
ccl_private Intersection *ccl_restrict isect,
const int index)
{
isect->prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, prim);
isect->object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, object);
isect->type = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, type);
isect->u = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, u);
isect->v = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, v);
isect->t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, t);
}
#if defined(__KERNEL_GPU__)
ccl_device_inline void integrator_state_copy_only(KernelGlobals kg,
ConstIntegratorState to_state,

16
intern/cycles/kernel/integrator/volume_stack.h
View File

@ -6,13 +6,15 @@
CCL_NAMESPACE_BEGIN
#ifdef __VOLUME__
/* Volumetric read/write lambda functions - default implementations */
#ifndef VOLUME_READ_LAMBDA
# define VOLUME_READ_LAMBDA(function_call) \
auto volume_read_lambda_pass = [=](const int i) { return function_call; };
# define VOLUME_WRITE_LAMBDA(function_call) \
auto volume_write_lambda_pass = [=](const int i, VolumeStack entry) { function_call; };
#endif
# ifndef VOLUME_READ_LAMBDA
# define VOLUME_READ_LAMBDA(function_call) \
auto volume_read_lambda_pass = [=](const int i) { return function_call; };
# define VOLUME_WRITE_LAMBDA(function_call) \
auto volume_write_lambda_pass = [=](const int i, VolumeStack entry) { function_call; };
# endif
/* Volume Stack
*
@ -210,4 +212,6 @@ ccl_device VolumeSampleMethod volume_stack_sample_method(KernelGlobals kg, Integ
return method;
}
#endif /* __VOLUME__*/
CCL_NAMESPACE_END

6
intern/cycles/kernel/svm/voxel.h
View File

@ -37,12 +37,14 @@ ccl_device_noinline int svm_node_tex_voxel(
r = kernel_tex_image_interp_3d(kg, id, co, INTERPOLATION_NONE);
}
else if (space != NODE_TEX_VOXEL_SPACE_OBJECT) {
else
#endif /* __VOLUME__ */
if (space != NODE_TEX_VOXEL_SPACE_OBJECT)
{
read_node_float(kg, &offset);
read_node_float(kg, &offset);
read_node_float(kg, &offset);
}
#endif
if (stack_valid(density_out_offset)) {
stack_store_float(stack, density_out_offset, r.w);

250
intern/cycles/kernel/types.h
View File

@ -62,6 +62,109 @@ CCL_NAMESPACE_BEGIN
# define INTEGRATOR_SHADOW_ISECT_SIZE INTEGRATOR_SHADOW_ISECT_SIZE_CPU
#endif
/* Kernel Features */
/* Shader nodes. */
#define KERNEL_FEATURE_NODE_BSDF (1U << 0U)
#define KERNEL_FEATURE_NODE_EMISSION (1U << 1U)
#define KERNEL_FEATURE_NODE_VOLUME (1U << 2U)
#define KERNEL_FEATURE_NODE_BUMP (1U << 3U)
#define KERNEL_FEATURE_NODE_BUMP_STATE (1U << 4U)
#define KERNEL_FEATURE_NODE_VORONOI_EXTRA (1U << 5U)
#define KERNEL_FEATURE_NODE_RAYTRACE (1U << 6U)
#define KERNEL_FEATURE_NODE_AOV (1U << 7U)
#define KERNEL_FEATURE_NODE_LIGHT_PATH (1U << 8U)
#define KERNEL_FEATURE_NODE_PRINCIPLED_HAIR (1U << 9U)
/* Use path tracing kernels. */
#define KERNEL_FEATURE_PATH_TRACING (1U << 10U)
/* BVH/sampling kernel features. */
#define KERNEL_FEATURE_POINTCLOUD (1U << 11U)
#define KERNEL_FEATURE_HAIR (1U << 12U)
#define KERNEL_FEATURE_HAIR_THICK (1U << 13U)
#define KERNEL_FEATURE_OBJECT_MOTION (1U << 14U)
/* Denotes whether baking functionality is needed. */
#define KERNEL_FEATURE_BAKING (1U << 15U)
/* Use subsurface scattering materials. */
#define KERNEL_FEATURE_SUBSURFACE (1U << 16U)
/* Use volume materials. */
#define KERNEL_FEATURE_VOLUME (1U << 17U)
/* Use OpenSubdiv patch evaluation */
#define KERNEL_FEATURE_PATCH_EVALUATION (1U << 18U)
/* Use Transparent shadows */
#define KERNEL_FEATURE_TRANSPARENT (1U << 19U)
/* Use shadow catcher. */
#define KERNEL_FEATURE_SHADOW_CATCHER (1U << 20U)
/* Light render passes. */
#define KERNEL_FEATURE_LIGHT_PASSES (1U << 21U)
/* AO. */
#define KERNEL_FEATURE_AO_PASS (1U << 22U)
#define KERNEL_FEATURE_AO_ADDITIVE (1U << 23U)
#define KERNEL_FEATURE_AO (KERNEL_FEATURE_AO_PASS | KERNEL_FEATURE_AO_ADDITIVE)
/* MNEE. */
#define KERNEL_FEATURE_MNEE (1U << 24U)
/* Path guiding. */
#define KERNEL_FEATURE_PATH_GUIDING (1U << 25U)
/* OSL. */
#define KERNEL_FEATURE_OSL (1U << 26U)
/* Light and shadow linking. */
#define KERNEL_FEATURE_LIGHT_LINKING (1U << 27U)
#define KERNEL_FEATURE_SHADOW_LINKING (1U << 28U)
/* Use denoising kernels and output denoising passes. */
#define KERNEL_FEATURE_DENOISING (1U << 29U)
/* Light tree. */
#define KERNEL_FEATURE_LIGHT_TREE (1U << 30U)
/* Shader node feature mask, to specialize shader evaluation for kernels. */
#define KERNEL_FEATURE_NODE_MASK_SURFACE_LIGHT \
(KERNEL_FEATURE_NODE_EMISSION | KERNEL_FEATURE_NODE_VORONOI_EXTRA | \
KERNEL_FEATURE_NODE_LIGHT_PATH)
#define KERNEL_FEATURE_NODE_MASK_SURFACE_BACKGROUND \
(KERNEL_FEATURE_NODE_MASK_SURFACE_LIGHT | KERNEL_FEATURE_NODE_AOV)
#define KERNEL_FEATURE_NODE_MASK_SURFACE_SHADOW \
(KERNEL_FEATURE_NODE_BSDF | KERNEL_FEATURE_NODE_EMISSION | KERNEL_FEATURE_NODE_BUMP | \
KERNEL_FEATURE_NODE_BUMP_STATE | KERNEL_FEATURE_NODE_VORONOI_EXTRA | \
KERNEL_FEATURE_NODE_LIGHT_PATH | KERNEL_FEATURE_NODE_PRINCIPLED_HAIR)
#define KERNEL_FEATURE_NODE_MASK_SURFACE \
(KERNEL_FEATURE_NODE_MASK_SURFACE_SHADOW | KERNEL_FEATURE_NODE_RAYTRACE | \
KERNEL_FEATURE_NODE_AOV | KERNEL_FEATURE_NODE_LIGHT_PATH)
#define KERNEL_FEATURE_NODE_MASK_VOLUME \
(KERNEL_FEATURE_NODE_EMISSION | KERNEL_FEATURE_NODE_VOLUME | \
KERNEL_FEATURE_NODE_VORONOI_EXTRA | KERNEL_FEATURE_NODE_LIGHT_PATH)
#define KERNEL_FEATURE_NODE_MASK_DISPLACEMENT \
(KERNEL_FEATURE_NODE_VORONOI_EXTRA | KERNEL_FEATURE_NODE_BUMP | KERNEL_FEATURE_NODE_BUMP_STATE)
#define KERNEL_FEATURE_NODE_MASK_BUMP KERNEL_FEATURE_NODE_MASK_DISPLACEMENT
/* Must be constexpr on the CPU to avoid compile errors because the state types
* are different depending on the main, shadow or null path. For GPU we don't have
* C++17 everywhere so need to check it. */
#if __cplusplus < 201703L
# define IF_KERNEL_FEATURE(feature) if ((node_feature_mask & (KERNEL_FEATURE_##feature)) != 0U)
# define IF_KERNEL_NODES_FEATURE(feature) \
if ((node_feature_mask & (KERNEL_FEATURE_NODE_##feature)) != 0U)
#else
# define IF_KERNEL_FEATURE(feature) \
if constexpr ((node_feature_mask & (KERNEL_FEATURE_##feature)) != 0U)
# define IF_KERNEL_NODES_FEATURE(feature) \
if constexpr ((node_feature_mask & (KERNEL_FEATURE_NODE_##feature)) != 0U)
#endif
/* Kernel features */
#define __AO__
#define __CAUSTICS_TRICKS__
@ -122,35 +225,55 @@ CCL_NAMESPACE_BEGIN
# endif
#endif
/* Scene-based selective features compilation. */
/* Scene-based selective features compilation. */
#ifdef __KERNEL_FEATURES__
# if !(__KERNEL_FEATURES & KERNEL_FEATURE_OBJECT_MOTION)
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_OBJECT_MOTION)
# undef __OBJECT_MOTION__
# endif
# if !(__KERNEL_FEATURES & KERNEL_FEATURE_HAIR)
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_HAIR)
# undef __HAIR__
# endif
# if !(__KERNEL_FEATURES & KERNEL_FEATURE_POINTCLOUD)
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_POINTCLOUD)
# undef __POINTCLOUD__
# endif
# if !(__KERNEL_FEATURES & KERNEL_FEATURE_VOLUME)
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_VOLUME)
# undef __VOLUME__
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_TRANSPARENT)
# undef __TRANSPARENT_SHADOWS__
# undef __SHADOW_RECORD_ALL__
# endif
# endif
# if !(__KERNEL_FEATURES & KERNEL_FEATURE_SUBSURFACE)
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_SUBSURFACE)
# undef __SUBSURFACE__
# endif
# if !(__KERNEL_FEATURES & KERNEL_FEATURE_PATCH_EVALUATION)
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_PATCH_EVALUATION)
# undef __PATCH_EVAL__
# endif
# if !(__KERNEL_FEATURES & KERNEL_FEATURE_TRANSPARENT)
# undef __TRANSPARENT_SHADOWS__
# endif
# if !(__KERNEL_FEATURES & KERNEL_FEATURE_SHADOW_CATCHER)
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_SHADOW_CATCHER)
# undef __SHADOW_CATCHER__
# endif
# if !(__KERNEL_FEATURES & KERNEL_FEATURE_DENOISING)
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_DENOISING)
# undef __DENOISING_FEATURES__
# endif
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_AO)
# undef __AO__
# endif
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_MNEE)
# undef __MNEE__
# endif
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_PATH_GUIDING)
# undef __PATH_GUIDING__
# endif
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_NODE_PRINCIPLED_HAIR)
# undef __PRINCIPLED_HAIR__
# endif
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_LIGHT_LINKING)
# undef __LIGHT_LINKING__
# endif
# if !(__KERNEL_FEATURES__ & KERNEL_FEATURE_SHADOW_LINKING)
# undef __SHADOW_LINKING__
# endif
#endif
#ifdef WITH_CYCLES_DEBUG_NAN
@ -1687,109 +1810,4 @@ enum {
DEVICE_KERNEL_INTEGRATOR_NUM = DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL + 1,
};
/* Kernel Features */
enum KernelFeatureFlag : uint32_t {
/* Shader nodes. */
KERNEL_FEATURE_NODE_BSDF = (1U << 0U),
KERNEL_FEATURE_NODE_EMISSION = (1U << 1U),
KERNEL_FEATURE_NODE_VOLUME = (1U << 2U),
KERNEL_FEATURE_NODE_BUMP = (1U << 3U),
KERNEL_FEATURE_NODE_BUMP_STATE = (1U << 4U),
KERNEL_FEATURE_NODE_VORONOI_EXTRA = (1U << 5U),
KERNEL_FEATURE_NODE_RAYTRACE = (1U << 6U),
KERNEL_FEATURE_NODE_AOV = (1U << 7U),
KERNEL_FEATURE_NODE_LIGHT_PATH = (1U << 8U),
KERNEL_FEATURE_NODE_PRINCIPLED_HAIR = (1U << 9U),
/* Use path tracing kernels. */
KERNEL_FEATURE_PATH_TRACING = (1U << 10U),
/* BVH/sampling kernel features. */
KERNEL_FEATURE_POINTCLOUD = (1U << 11U),
KERNEL_FEATURE_HAIR = (1U << 12U),
KERNEL_FEATURE_HAIR_THICK = (1U << 13U),
KERNEL_FEATURE_OBJECT_MOTION = (1U << 14U),
/* Denotes whether baking functionality is needed. */
KERNEL_FEATURE_BAKING = (1U << 15U),
/* Use subsurface scattering materials. */
KERNEL_FEATURE_SUBSURFACE = (1U << 16U),
/* Use volume materials. */
KERNEL_FEATURE_VOLUME = (1U << 17U),
/* Use OpenSubdiv patch evaluation */
KERNEL_FEATURE_PATCH_EVALUATION = (1U << 18U),
/* Use Transparent shadows */
KERNEL_FEATURE_TRANSPARENT = (1U << 19U),
/* Use shadow catcher. */
KERNEL_FEATURE_SHADOW_CATCHER = (1U << 20U),
/* Light render passes. */
KERNEL_FEATURE_LIGHT_PASSES = (1U << 21U),
/* AO. */
KERNEL_FEATURE_AO_PASS = (1U << 22U),
KERNEL_FEATURE_AO_ADDITIVE = (1U << 23U),
KERNEL_FEATURE_AO = (KERNEL_FEATURE_AO_PASS | KERNEL_FEATURE_AO_ADDITIVE),
/* MNEE. */
KERNEL_FEATURE_MNEE = (1U << 24U),
/* Path guiding. */
KERNEL_FEATURE_PATH_GUIDING = (1U << 25U),
/* OSL. */
KERNEL_FEATURE_OSL = (1U << 26U),
/* Light and shadow linking. */
KERNEL_FEATURE_LIGHT_LINKING = (1U << 27U),
KERNEL_FEATURE_SHADOW_LINKING = (1U << 28U),
/* Use denoising kernels and output denoising passes. */
KERNEL_FEATURE_DENOISING = (1U << 29U),
/* Light tree. */
KERNEL_FEATURE_LIGHT_TREE = (1U << 30U),
};
/* Shader node feature mask, to specialize shader evaluation for kernels. */
#define KERNEL_FEATURE_NODE_MASK_SURFACE_LIGHT \
(KERNEL_FEATURE_NODE_EMISSION | KERNEL_FEATURE_NODE_VORONOI_EXTRA | \
KERNEL_FEATURE_NODE_LIGHT_PATH)
#define KERNEL_FEATURE_NODE_MASK_SURFACE_BACKGROUND \
(KERNEL_FEATURE_NODE_MASK_SURFACE_LIGHT | KERNEL_FEATURE_NODE_AOV)
#define KERNEL_FEATURE_NODE_MASK_SURFACE_SHADOW \
(KERNEL_FEATURE_NODE_BSDF | KERNEL_FEATURE_NODE_EMISSION | KERNEL_FEATURE_NODE_BUMP | \
KERNEL_FEATURE_NODE_BUMP_STATE | KERNEL_FEATURE_NODE_VORONOI_EXTRA | \
KERNEL_FEATURE_NODE_LIGHT_PATH | KERNEL_FEATURE_NODE_PRINCIPLED_HAIR)
#define KERNEL_FEATURE_NODE_MASK_SURFACE \
(KERNEL_FEATURE_NODE_MASK_SURFACE_SHADOW | KERNEL_FEATURE_NODE_RAYTRACE | \
KERNEL_FEATURE_NODE_AOV | KERNEL_FEATURE_NODE_LIGHT_PATH)
#define KERNEL_FEATURE_NODE_MASK_VOLUME \
(KERNEL_FEATURE_NODE_EMISSION | KERNEL_FEATURE_NODE_VOLUME | \
KERNEL_FEATURE_NODE_VORONOI_EXTRA | KERNEL_FEATURE_NODE_LIGHT_PATH)
#define KERNEL_FEATURE_NODE_MASK_DISPLACEMENT \
(KERNEL_FEATURE_NODE_VORONOI_EXTRA | KERNEL_FEATURE_NODE_BUMP | KERNEL_FEATURE_NODE_BUMP_STATE)
#define KERNEL_FEATURE_NODE_MASK_BUMP KERNEL_FEATURE_NODE_MASK_DISPLACEMENT
/* Must be constexpr on the CPU to avoid compile errors because the state types
* are different depending on the main, shadow or null path. For GPU we don't have
* C++17 everywhere so need to check it. */
#if __cplusplus < 201703L
# define IF_KERNEL_FEATURE(feature) if ((node_feature_mask & (KERNEL_FEATURE_##feature)) != 0U)
# define IF_KERNEL_NODES_FEATURE(feature) \
if ((node_feature_mask & (KERNEL_FEATURE_NODE_##feature)) != 0U)
#else
# define IF_KERNEL_FEATURE(feature) \
if constexpr ((node_feature_mask & (KERNEL_FEATURE_##feature)) != 0U)
# define IF_KERNEL_NODES_FEATURE(feature) \
if constexpr ((node_feature_mask & (KERNEL_FEATURE_NODE_##feature)) != 0U)
#endif
CCL_NAMESPACE_END