diff options
Diffstat (limited to 'intern/cycles/kernel/svm')
40 files changed, 261 insertions, 250 deletions
diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h index 871e370123e..472f3517839 100644 --- a/intern/cycles/kernel/svm/svm.h +++ b/intern/cycles/kernel/svm/svm.h @@ -107,15 +107,14 @@ ccl_device_inline bool stack_valid(uint a) /* Reading Nodes */ -ccl_device_inline uint4 read_node(ccl_global const KernelGlobals *kg, ccl_private int *offset) +ccl_device_inline uint4 read_node(KernelGlobals kg, ccl_private int *offset) { uint4 node = kernel_tex_fetch(__svm_nodes, *offset); (*offset)++; return node; } -ccl_device_inline float4 read_node_float(ccl_global const KernelGlobals *kg, - ccl_private int *offset) +ccl_device_inline float4 read_node_float(KernelGlobals kg, ccl_private int *offset) { uint4 node = kernel_tex_fetch(__svm_nodes, *offset); float4 f = make_float4(__uint_as_float(node.x), @@ -126,7 +125,7 @@ ccl_device_inline float4 read_node_float(ccl_global const KernelGlobals *kg, return f; } -ccl_device_inline float4 fetch_node_float(ccl_global const KernelGlobals *kg, int offset) +ccl_device_inline float4 fetch_node_float(KernelGlobals kg, int offset) { uint4 node = kernel_tex_fetch(__svm_nodes, offset); return make_float4(__uint_as_float(node.x), @@ -226,11 +225,12 @@ CCL_NAMESPACE_END CCL_NAMESPACE_BEGIN /* Main Interpreter Loop */ -template<uint node_feature_mask, ShaderType type> -ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, +template<uint node_feature_mask, ShaderType type, typename ConstIntegratorGenericState> +ccl_device void svm_eval_nodes(KernelGlobals kg, + ConstIntegratorGenericState state, ShaderData *sd, ccl_global float *render_buffer, - int path_flag) + uint32_t path_flag) { float stack[SVM_STACK_SIZE]; int offset = sd->shader & SHADER_MASK; @@ -257,12 +257,14 @@ ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, kg, sd, stack, node, path_flag, offset); break; case NODE_CLOSURE_EMISSION: - if (KERNEL_NODES_FEATURE(EMISSION)) { + IF_KERNEL_NODES_FEATURE(EMISSION) + { svm_node_closure_emission(sd, stack, node); } break; case NODE_CLOSURE_BACKGROUND: - if (KERNEL_NODES_FEATURE(EMISSION)) { + IF_KERNEL_NODES_FEATURE(EMISSION) + { svm_node_closure_background(sd, stack, node); } break; @@ -273,7 +275,8 @@ ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, svm_node_closure_weight(sd, stack, node.y); break; case NODE_EMISSION_WEIGHT: - if (KERNEL_NODES_FEATURE(EMISSION)) { + IF_KERNEL_NODES_FEATURE(EMISSION) + { svm_node_emission_weight(kg, sd, stack, node); } break; @@ -310,27 +313,32 @@ ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, svm_node_vertex_color(kg, sd, stack, node.y, node.z, node.w); break; case NODE_GEOMETRY_BUMP_DX: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { svm_node_geometry_bump_dx(kg, sd, stack, node.y, node.z); } break; case NODE_GEOMETRY_BUMP_DY: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { svm_node_geometry_bump_dy(kg, sd, stack, node.y, node.z); } break; case NODE_SET_DISPLACEMENT: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { svm_node_set_displacement(kg, sd, stack, node.y); } break; case NODE_DISPLACEMENT: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { svm_node_displacement(kg, sd, stack, node); } break; case NODE_VECTOR_DISPLACEMENT: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { offset = svm_node_vector_displacement(kg, sd, stack, node, offset); } break; @@ -344,52 +352,62 @@ ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, offset = svm_node_tex_noise(kg, sd, stack, node.y, node.z, node.w, offset); break; case NODE_SET_BUMP: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { svm_node_set_bump(kg, sd, stack, node); } break; case NODE_ATTR_BUMP_DX: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { svm_node_attr_bump_dx(kg, sd, stack, node); } break; case NODE_ATTR_BUMP_DY: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { svm_node_attr_bump_dy(kg, sd, stack, node); } break; case NODE_VERTEX_COLOR_BUMP_DX: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { svm_node_vertex_color_bump_dx(kg, sd, stack, node.y, node.z, node.w); } break; case NODE_VERTEX_COLOR_BUMP_DY: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { svm_node_vertex_color_bump_dy(kg, sd, stack, node.y, node.z, node.w); } break; case NODE_TEX_COORD_BUMP_DX: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { offset = svm_node_tex_coord_bump_dx(kg, sd, path_flag, stack, node, offset); } break; case NODE_TEX_COORD_BUMP_DY: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { offset = svm_node_tex_coord_bump_dy(kg, sd, path_flag, stack, node, offset); } break; case NODE_CLOSURE_SET_NORMAL: - if (KERNEL_NODES_FEATURE(BUMP)) { + IF_KERNEL_NODES_FEATURE(BUMP) + { svm_node_set_normal(kg, sd, stack, node.y, node.z); } break; case NODE_ENTER_BUMP_EVAL: - if (KERNEL_NODES_FEATURE(BUMP_STATE)) { + IF_KERNEL_NODES_FEATURE(BUMP_STATE) + { svm_node_enter_bump_eval(kg, sd, stack, node.y); } break; case NODE_LEAVE_BUMP_EVAL: - if (KERNEL_NODES_FEATURE(BUMP_STATE)) { + IF_KERNEL_NODES_FEATURE(BUMP_STATE) + { svm_node_leave_bump_eval(kg, sd, stack, node.y); } break; @@ -407,12 +425,14 @@ ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, svm_node_layer_weight(sd, stack, node); break; case NODE_CLOSURE_VOLUME: - if (KERNEL_NODES_FEATURE(VOLUME)) { + IF_KERNEL_NODES_FEATURE(VOLUME) + { svm_node_closure_volume<type>(kg, sd, stack, node); } break; case NODE_PRINCIPLED_VOLUME: - if (KERNEL_NODES_FEATURE(VOLUME)) { + IF_KERNEL_NODES_FEATURE(VOLUME) + { offset = svm_node_principled_volume<type>(kg, sd, stack, node, path_flag, offset); } break; @@ -432,7 +452,7 @@ ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, svm_node_brightness(sd, stack, node.y, node.z, node.w); break; case NODE_LIGHT_PATH: - svm_node_light_path(INTEGRATOR_STATE_PASS, sd, stack, node.y, node.z, path_flag); + svm_node_light_path<node_feature_mask>(kg, state, sd, stack, node.y, node.z, path_flag); break; case NODE_OBJECT_INFO: svm_node_object_info(kg, sd, stack, node.y, node.z); @@ -442,7 +462,8 @@ ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, break; #if defined(__HAIR__) case NODE_HAIR_INFO: - if (KERNEL_NODES_FEATURE(HAIR)) { + IF_KERNEL_NODES_FEATURE(HAIR) + { svm_node_hair_info(kg, sd, stack, node.y, node.z); } break; @@ -554,15 +575,16 @@ ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, break; #ifdef __SHADER_RAYTRACE__ case NODE_BEVEL: - svm_node_bevel<node_feature_mask>(INTEGRATOR_STATE_PASS, sd, stack, node); + svm_node_bevel<node_feature_mask>(kg, state, sd, stack, node); break; case NODE_AMBIENT_OCCLUSION: - svm_node_ao<node_feature_mask>(INTEGRATOR_STATE_PASS, sd, stack, node); + svm_node_ao<node_feature_mask>(kg, state, sd, stack, node); break; #endif case NODE_TEX_VOXEL: - if (KERNEL_NODES_FEATURE(VOLUME)) { + IF_KERNEL_NODES_FEATURE(VOLUME) + { offset = svm_node_tex_voxel(kg, sd, stack, node, offset); } break; @@ -572,10 +594,10 @@ ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, } break; case NODE_AOV_COLOR: - svm_node_aov_color(INTEGRATOR_STATE_PASS, sd, stack, node, render_buffer); + svm_node_aov_color<node_feature_mask>(kg, state, sd, stack, node, render_buffer); break; case NODE_AOV_VALUE: - svm_node_aov_value(INTEGRATOR_STATE_PASS, sd, stack, node, render_buffer); + svm_node_aov_value<node_feature_mask>(kg, state, sd, stack, node, render_buffer); break; default: kernel_assert(!"Unknown node type was passed to the SVM machine"); diff --git a/intern/cycles/kernel/svm/svm_ao.h b/intern/cycles/kernel/svm/svm_ao.h index 092f3817fd8..4cfef7bc204 100644 --- a/intern/cycles/kernel/svm/svm_ao.h +++ b/intern/cycles/kernel/svm/svm_ao.h @@ -21,15 +21,17 @@ CCL_NAMESPACE_BEGIN #ifdef __SHADER_RAYTRACE__ # ifdef __KERNEL_OPTIX__ -extern "C" __device__ float __direct_callable__svm_node_ao(INTEGRATOR_STATE_CONST_ARGS, +extern "C" __device__ float __direct_callable__svm_node_ao( # else -ccl_device float svm_ao(INTEGRATOR_STATE_CONST_ARGS, +ccl_device float svm_ao( # endif - ccl_private ShaderData *sd, - float3 N, - float max_dist, - int num_samples, - int flags) + KernelGlobals kg, + ConstIntegratorState state, + ccl_private ShaderData *sd, + float3 N, + float max_dist, + int num_samples, + int flags) { if (flags & NODE_AO_GLOBAL_RADIUS) { max_dist = kernel_data.integrator.ao_bounces_distance; @@ -54,7 +56,7 @@ ccl_device float svm_ao(INTEGRATOR_STATE_CONST_ARGS, /* TODO: support ray-tracing in shadow shader evaluation? */ RNGState rng_state; - path_state_rng_load(INTEGRATOR_STATE_PASS, &rng_state); + path_state_rng_load(state, &rng_state); int unoccluded = 0; for (int sample = 0; sample < num_samples; sample++) { @@ -89,14 +91,15 @@ ccl_device float svm_ao(INTEGRATOR_STATE_CONST_ARGS, return ((float)unoccluded) / num_samples; } -template<uint node_feature_mask> +template<uint node_feature_mask, typename ConstIntegratorGenericState> # if defined(__KERNEL_OPTIX__) ccl_device_inline # else ccl_device_noinline # endif void - svm_node_ao(INTEGRATOR_STATE_CONST_ARGS, + svm_node_ao(KernelGlobals kg, + ConstIntegratorGenericState state, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) @@ -107,16 +110,17 @@ ccl_device_noinline uint color_offset, out_color_offset, samples; svm_unpack_node_uchar3(node.z, &color_offset, &out_color_offset, &samples); - float dist = stack_load_float_default(stack, dist_offset, node.w); - float3 normal = stack_valid(normal_offset) ? stack_load_float3(stack, normal_offset) : sd->N; - float ao = 1.0f; - if (KERNEL_NODES_FEATURE(RAYTRACE)) { + IF_KERNEL_NODES_FEATURE(RAYTRACE) + { + float dist = stack_load_float_default(stack, dist_offset, node.w); + float3 normal = stack_valid(normal_offset) ? stack_load_float3(stack, normal_offset) : sd->N; + # ifdef __KERNEL_OPTIX__ - ao = optixDirectCall<float>(0, INTEGRATOR_STATE_PASS, sd, normal, dist, samples, flags); + ao = optixDirectCall<float>(0, kg, state, sd, normal, dist, samples, flags); # else - ao = svm_ao(INTEGRATOR_STATE_PASS, sd, normal, dist, samples, flags); + ao = svm_ao(kg, state, sd, normal, dist, samples, flags); # endif } diff --git a/intern/cycles/kernel/svm/svm_aov.h b/intern/cycles/kernel/svm/svm_aov.h index 640bec87ac9..833a6443b3c 100644 --- a/intern/cycles/kernel/svm/svm_aov.h +++ b/intern/cycles/kernel/svm/svm_aov.h @@ -18,23 +18,26 @@ CCL_NAMESPACE_BEGIN -ccl_device_inline bool svm_node_aov_check(const int path_flag, ccl_global float *render_buffer) +ccl_device_inline bool svm_node_aov_check(const uint32_t path_flag, + ccl_global float *render_buffer) { bool is_primary = (path_flag & PATH_RAY_CAMERA) && (!(path_flag & PATH_RAY_SINGLE_PASS_DONE)); return ((render_buffer != NULL) && is_primary); } -ccl_device void svm_node_aov_color(INTEGRATOR_STATE_CONST_ARGS, +template<uint node_feature_mask, typename ConstIntegratorGenericState> +ccl_device void svm_node_aov_color(KernelGlobals kg, + ConstIntegratorGenericState state, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, ccl_global float *render_buffer) { - float3 val = stack_load_float3(stack, node.y); - - if (render_buffer && !INTEGRATOR_STATE_IS_NULL) { - const uint32_t render_pixel_index = INTEGRATOR_STATE(path, render_pixel_index); + IF_KERNEL_NODES_FEATURE(AOV) + { + const float3 val = stack_load_float3(stack, node.y); + const uint32_t render_pixel_index = INTEGRATOR_STATE(state, path, render_pixel_index); const uint64_t render_buffer_offset = (uint64_t)render_pixel_index * kernel_data.film.pass_stride; ccl_global float *buffer = render_buffer + render_buffer_offset + @@ -43,16 +46,18 @@ ccl_device void svm_node_aov_color(INTEGRATOR_STATE_CONST_ARGS, } } -ccl_device void svm_node_aov_value(INTEGRATOR_STATE_CONST_ARGS, +template<uint node_feature_mask, typename ConstIntegratorGenericState> +ccl_device void svm_node_aov_value(KernelGlobals kg, + ConstIntegratorGenericState state, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, ccl_global float *render_buffer) { - float val = stack_load_float(stack, node.y); - - if (render_buffer && !INTEGRATOR_STATE_IS_NULL) { - const uint32_t render_pixel_index = INTEGRATOR_STATE(path, render_pixel_index); + IF_KERNEL_NODES_FEATURE(AOV) + { + const float val = stack_load_float(stack, node.y); + const uint32_t render_pixel_index = INTEGRATOR_STATE(state, path, render_pixel_index); const uint64_t render_buffer_offset = (uint64_t)render_pixel_index * kernel_data.film.pass_stride; ccl_global float *buffer = render_buffer + render_buffer_offset + diff --git a/intern/cycles/kernel/svm/svm_attribute.h b/intern/cycles/kernel/svm/svm_attribute.h index 9fd401ba1c3..b3c66d29f5c 100644 --- a/intern/cycles/kernel/svm/svm_attribute.h +++ b/intern/cycles/kernel/svm/svm_attribute.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Attribute Node */ -ccl_device AttributeDescriptor svm_node_attr_init(ccl_global const KernelGlobals *kg, +ccl_device AttributeDescriptor svm_node_attr_init(KernelGlobals kg, ccl_private ShaderData *sd, uint4 node, ccl_private NodeAttributeOutputType *type, @@ -48,7 +48,7 @@ ccl_device AttributeDescriptor svm_node_attr_init(ccl_global const KernelGlobals } template<uint node_feature_mask> -ccl_device_noinline void svm_node_attr(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_attr(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) @@ -58,7 +58,8 @@ ccl_device_noinline void svm_node_attr(ccl_global const KernelGlobals *kg, AttributeDescriptor desc = svm_node_attr_init(kg, sd, node, &type, &out_offset); #ifdef __VOLUME__ - if (KERNEL_NODES_FEATURE(VOLUME)) { + IF_KERNEL_NODES_FEATURE(VOLUME) + { /* Volumes * NOTE: moving this into its own node type might help improve performance. */ if (primitive_is_volume_attribute(sd, desc)) { @@ -148,7 +149,7 @@ ccl_device_noinline void svm_node_attr(ccl_global const KernelGlobals *kg, } } -ccl_device_noinline void svm_node_attr_bump_dx(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_attr_bump_dx(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) @@ -244,7 +245,7 @@ ccl_device_noinline void svm_node_attr_bump_dx(ccl_global const KernelGlobals *k } } -ccl_device_noinline void svm_node_attr_bump_dy(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_attr_bump_dy(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) diff --git a/intern/cycles/kernel/svm/svm_bevel.h b/intern/cycles/kernel/svm/svm_bevel.h index a76584e6bc8..292887beedf 100644 --- a/intern/cycles/kernel/svm/svm_bevel.h +++ b/intern/cycles/kernel/svm/svm_bevel.h @@ -99,13 +99,15 @@ ccl_device void svm_bevel_cubic_sample(const float radius, */ # ifdef __KERNEL_OPTIX__ -extern "C" __device__ float3 __direct_callable__svm_node_bevel(INTEGRATOR_STATE_CONST_ARGS, +extern "C" __device__ float3 __direct_callable__svm_node_bevel( # else -ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS, +ccl_device float3 svm_bevel( # endif - ccl_private ShaderData *sd, - float radius, - int num_samples) + KernelGlobals kg, + ConstIntegratorState state, + ccl_private ShaderData *sd, + float radius, + int num_samples) { /* Early out if no sampling needed. */ if (radius <= 0.0f || num_samples < 1 || sd->object == OBJECT_NONE) { @@ -118,15 +120,15 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS, } /* Don't bevel for blurry indirect rays. */ - if (INTEGRATOR_STATE(path, min_ray_pdf) < 8.0f) { + if (INTEGRATOR_STATE(state, path, min_ray_pdf) < 8.0f) { return sd->N; } /* Setup for multi intersection. */ LocalIntersection isect; - uint lcg_state = lcg_state_init(INTEGRATOR_STATE(path, rng_hash), - INTEGRATOR_STATE(path, rng_offset), - INTEGRATOR_STATE(path, sample), + uint lcg_state = lcg_state_init(INTEGRATOR_STATE(state, path, rng_hash), + INTEGRATOR_STATE(state, path, rng_offset), + INTEGRATOR_STATE(state, path, sample), 0x64c6a40e); /* Sample normals from surrounding points on surface. */ @@ -134,7 +136,7 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS, /* TODO: support ray-tracing in shadow shader evaluation? */ RNGState rng_state; - path_state_rng_load(INTEGRATOR_STATE_PASS, &rng_state); + path_state_rng_load(state, &rng_state); for (int sample = 0; sample < num_samples; sample++) { float disk_u, disk_v; @@ -280,14 +282,15 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS, return is_zero(N) ? sd->N : (sd->flag & SD_BACKFACING) ? -N : N; } -template<uint node_feature_mask> +template<uint node_feature_mask, typename ConstIntegratorGenericState> # if defined(__KERNEL_OPTIX__) ccl_device_inline # else ccl_device_noinline # endif void - svm_node_bevel(INTEGRATOR_STATE_CONST_ARGS, + svm_node_bevel(KernelGlobals kg, + ConstIntegratorGenericState state, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) @@ -295,15 +298,16 @@ ccl_device_noinline uint num_samples, radius_offset, normal_offset, out_offset; svm_unpack_node_uchar4(node.y, &num_samples, &radius_offset, &normal_offset, &out_offset); - float radius = stack_load_float(stack, radius_offset); - float3 bevel_N = sd->N; - if (KERNEL_NODES_FEATURE(RAYTRACE)) { + IF_KERNEL_NODES_FEATURE(RAYTRACE) + { + float radius = stack_load_float(stack, radius_offset); + # ifdef __KERNEL_OPTIX__ - bevel_N = optixDirectCall<float3>(1, INTEGRATOR_STATE_PASS, sd, radius, num_samples); + bevel_N = optixDirectCall<float3>(1, kg, state, sd, radius, num_samples); # else - bevel_N = svm_bevel(INTEGRATOR_STATE_PASS, sd, radius, num_samples); + bevel_N = svm_bevel(kg, state, sd, radius, num_samples); # endif if (stack_valid(normal_offset)) { diff --git a/intern/cycles/kernel/svm/svm_blackbody.h b/intern/cycles/kernel/svm/svm_blackbody.h index 521afb42adc..f1adb0e76af 100644 --- a/intern/cycles/kernel/svm/svm_blackbody.h +++ b/intern/cycles/kernel/svm/svm_blackbody.h @@ -34,7 +34,7 @@ CCL_NAMESPACE_BEGIN /* Blackbody Node */ -ccl_device_noinline void svm_node_blackbody(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_blackbody(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint temperature_offset, diff --git a/intern/cycles/kernel/svm/svm_brick.h b/intern/cycles/kernel/svm/svm_brick.h index 29a8350f1c1..9dc31ef37ec 100644 --- a/intern/cycles/kernel/svm/svm_brick.h +++ b/intern/cycles/kernel/svm/svm_brick.h @@ -72,11 +72,8 @@ ccl_device_noinline_cpu float2 svm_brick(float3 p, return make_float2(tint, mortar); } -ccl_device_noinline int svm_node_tex_brick(ccl_global const KernelGlobals *kg, - ccl_private ShaderData *sd, - ccl_private float *stack, - uint4 node, - int offset) +ccl_device_noinline int svm_node_tex_brick( + KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset) { uint4 node2 = read_node(kg, &offset); uint4 node3 = read_node(kg, &offset); diff --git a/intern/cycles/kernel/svm/svm_bump.h b/intern/cycles/kernel/svm/svm_bump.h index 70935c730f4..66e5b665532 100644 --- a/intern/cycles/kernel/svm/svm_bump.h +++ b/intern/cycles/kernel/svm/svm_bump.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Bump Eval Nodes */ -ccl_device_noinline void svm_node_enter_bump_eval(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_enter_bump_eval(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint offset) @@ -45,7 +45,7 @@ ccl_device_noinline void svm_node_enter_bump_eval(ccl_global const KernelGlobals } } -ccl_device_noinline void svm_node_leave_bump_eval(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_leave_bump_eval(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint offset) diff --git a/intern/cycles/kernel/svm/svm_camera.h b/intern/cycles/kernel/svm/svm_camera.h index 2b786757af8..787f11f38b5 100644 --- a/intern/cycles/kernel/svm/svm_camera.h +++ b/intern/cycles/kernel/svm/svm_camera.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -ccl_device_noinline void svm_node_camera(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_camera(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint out_vector, diff --git a/intern/cycles/kernel/svm/svm_checker.h b/intern/cycles/kernel/svm/svm_checker.h index e22367f4f59..9251d90c0e1 100644 --- a/intern/cycles/kernel/svm/svm_checker.h +++ b/intern/cycles/kernel/svm/svm_checker.h @@ -32,7 +32,7 @@ ccl_device float svm_checker(float3 p) return ((xi % 2 == yi % 2) == (zi % 2)) ? 1.0f : 0.0f; } -ccl_device_noinline void svm_node_tex_checker(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_tex_checker(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) diff --git a/intern/cycles/kernel/svm/svm_clamp.h b/intern/cycles/kernel/svm/svm_clamp.h index cb5224aebb2..5b5ea784f4a 100644 --- a/intern/cycles/kernel/svm/svm_clamp.h +++ b/intern/cycles/kernel/svm/svm_clamp.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Clamp Node */ -ccl_device_noinline int svm_node_clamp(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_clamp(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint value_stack_offset, diff --git a/intern/cycles/kernel/svm/svm_closure.h b/intern/cycles/kernel/svm/svm_closure.h index 87be73bb2cc..3378832c233 100644 --- a/intern/cycles/kernel/svm/svm_closure.h +++ b/intern/cycles/kernel/svm/svm_closure.h @@ -61,12 +61,25 @@ ccl_device void svm_node_glass_setup(ccl_private ShaderData *sd, } } +ccl_device_inline int svm_node_closure_bsdf_skip(KernelGlobals kg, int offset, uint type) +{ + if (type == CLOSURE_BSDF_PRINCIPLED_ID) { + /* Read all principled BSDF extra data to get the right offset. */ + read_node(kg, &offset); + read_node(kg, &offset); + read_node(kg, &offset); + read_node(kg, &offset); + } + + return offset; +} + template<uint node_feature_mask, ShaderType shader_type> -ccl_device_noinline int svm_node_closure_bsdf(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, - int path_flag, + uint32_t path_flag, int offset) { uint type, param1_offset, param2_offset; @@ -80,16 +93,15 @@ ccl_device_noinline int svm_node_closure_bsdf(ccl_global const KernelGlobals *kg uint4 data_node = read_node(kg, &offset); /* Only compute BSDF for surfaces, transparent variable is shared with volume extinction. */ - if ((!KERNEL_NODES_FEATURE(BSDF) || shader_type != SHADER_TYPE_SURFACE) || mix_weight == 0.0f) { - if (type == CLOSURE_BSDF_PRINCIPLED_ID) { - /* Read all principled BSDF extra data to get the right offset. */ - read_node(kg, &offset); - read_node(kg, &offset); - read_node(kg, &offset); - read_node(kg, &offset); + IF_KERNEL_NODES_FEATURE(BSDF) + { + if ((shader_type != SHADER_TYPE_SURFACE) || mix_weight == 0.0f) { + return svm_node_closure_bsdf_skip(kg, offset, type); } - - return offset; + } + else + { + return svm_node_closure_bsdf_skip(kg, offset, type); } float3 N = stack_valid(data_node.x) ? stack_load_float3(stack, data_node.x) : sd->N; @@ -944,7 +956,7 @@ ccl_device_noinline int svm_node_closure_bsdf(ccl_global const KernelGlobals *kg } template<ShaderType shader_type> -ccl_device_noinline void svm_node_closure_volume(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_closure_volume(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) @@ -999,11 +1011,11 @@ ccl_device_noinline void svm_node_closure_volume(ccl_global const KernelGlobals } template<ShaderType shader_type> -ccl_device_noinline int svm_node_principled_volume(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_principled_volume(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, - int path_flag, + uint32_t path_flag, int offset) { #ifdef __VOLUME__ @@ -1194,7 +1206,7 @@ ccl_device void svm_node_closure_weight(ccl_private ShaderData *sd, svm_node_closure_store_weight(sd, weight); } -ccl_device_noinline void svm_node_emission_weight(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_emission_weight(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) @@ -1232,7 +1244,7 @@ ccl_device_noinline void svm_node_mix_closure(ccl_private ShaderData *sd, /* (Bump) normal */ -ccl_device void svm_node_set_normal(ccl_global const KernelGlobals *kg, +ccl_device void svm_node_set_normal(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint in_direction, diff --git a/intern/cycles/kernel/svm/svm_convert.h b/intern/cycles/kernel/svm/svm_convert.h index 0d53779a5c8..ec5745dc78a 100644 --- a/intern/cycles/kernel/svm/svm_convert.h +++ b/intern/cycles/kernel/svm/svm_convert.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Conversion Nodes */ -ccl_device_noinline void svm_node_convert(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_convert(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint type, diff --git a/intern/cycles/kernel/svm/svm_displace.h b/intern/cycles/kernel/svm/svm_displace.h index 7a3c8a6d36d..f2446c3b3ef 100644 --- a/intern/cycles/kernel/svm/svm_displace.h +++ b/intern/cycles/kernel/svm/svm_displace.h @@ -20,7 +20,7 @@ CCL_NAMESPACE_BEGIN /* Bump Node */ -ccl_device_noinline void svm_node_set_bump(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_set_bump(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) @@ -88,7 +88,7 @@ ccl_device_noinline void svm_node_set_bump(ccl_global const KernelGlobals *kg, /* Displacement Node */ -ccl_device void svm_node_set_displacement(ccl_global const KernelGlobals *kg, +ccl_device void svm_node_set_displacement(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint fac_offset) @@ -97,7 +97,7 @@ ccl_device void svm_node_set_displacement(ccl_global const KernelGlobals *kg, sd->P += dP; } -ccl_device_noinline void svm_node_displacement(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_displacement(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) @@ -127,11 +127,8 @@ ccl_device_noinline void svm_node_displacement(ccl_global const KernelGlobals *k stack_store_float3(stack, node.z, dP); } -ccl_device_noinline int svm_node_vector_displacement(ccl_global const KernelGlobals *kg, - ccl_private ShaderData *sd, - ccl_private float *stack, - uint4 node, - int offset) +ccl_device_noinline int svm_node_vector_displacement( + KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset) { uint4 data_node = read_node(kg, &offset); uint space = data_node.x; diff --git a/intern/cycles/kernel/svm/svm_geometry.h b/intern/cycles/kernel/svm/svm_geometry.h index a94464d3a52..b29bfdbed07 100644 --- a/intern/cycles/kernel/svm/svm_geometry.h +++ b/intern/cycles/kernel/svm/svm_geometry.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Geometry Node */ -ccl_device_noinline void svm_node_geometry(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_geometry(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint type, @@ -54,7 +54,7 @@ ccl_device_noinline void svm_node_geometry(ccl_global const KernelGlobals *kg, stack_store_float3(stack, out_offset, data); } -ccl_device_noinline void svm_node_geometry_bump_dx(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_geometry_bump_dx(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint type, @@ -81,7 +81,7 @@ ccl_device_noinline void svm_node_geometry_bump_dx(ccl_global const KernelGlobal #endif } -ccl_device_noinline void svm_node_geometry_bump_dy(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_geometry_bump_dy(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint type, @@ -110,7 +110,7 @@ ccl_device_noinline void svm_node_geometry_bump_dy(ccl_global const KernelGlobal /* Object Info */ -ccl_device_noinline void svm_node_object_info(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_object_info(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint type, @@ -152,7 +152,7 @@ ccl_device_noinline void svm_node_object_info(ccl_global const KernelGlobals *kg /* Particle Info */ -ccl_device_noinline void svm_node_particle_info(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_particle_info(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint type, @@ -214,7 +214,7 @@ ccl_device_noinline void svm_node_particle_info(ccl_global const KernelGlobals * /* Hair Info */ -ccl_device_noinline void svm_node_hair_info(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_hair_info(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint type, diff --git a/intern/cycles/kernel/svm/svm_hsv.h b/intern/cycles/kernel/svm/svm_hsv.h index feb85eda122..978c4c2d781 100644 --- a/intern/cycles/kernel/svm/svm_hsv.h +++ b/intern/cycles/kernel/svm/svm_hsv.h @@ -19,7 +19,7 @@ CCL_NAMESPACE_BEGIN -ccl_device_noinline void svm_node_hsv(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_hsv(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) diff --git a/intern/cycles/kernel/svm/svm_ies.h b/intern/cycles/kernel/svm/svm_ies.h index 7d41205c9ef..0215670d062 100644 --- a/intern/cycles/kernel/svm/svm_ies.h +++ b/intern/cycles/kernel/svm/svm_ies.h @@ -19,7 +19,7 @@ CCL_NAMESPACE_BEGIN /* IES Light */ ccl_device_inline float interpolate_ies_vertical( - ccl_global const KernelGlobals *kg, int ofs, int v, int v_num, float v_frac, int h) + KernelGlobals kg, int ofs, int v, int v_num, float v_frac, int h) { /* Since lookups are performed in spherical coordinates, clamping the coordinates at the low end * of v (corresponding to the north pole) would result in artifacts. The proper way of dealing @@ -39,10 +39,7 @@ ccl_device_inline float interpolate_ies_vertical( return cubic_interp(a, b, c, d, v_frac); } -ccl_device_inline float kernel_ies_interp(ccl_global const KernelGlobals *kg, - int slot, - float h_angle, - float v_angle) +ccl_device_inline float kernel_ies_interp(KernelGlobals kg, int slot, float h_angle, float v_angle) { /* Find offset of the IES data in the table. */ int ofs = __float_as_int(kernel_tex_fetch(__ies, slot)); @@ -98,7 +95,7 @@ ccl_device_inline float kernel_ies_interp(ccl_global const KernelGlobals *kg, return max(cubic_interp(a, b, c, d, h_frac), 0.0f); } -ccl_device_noinline void svm_node_ies(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_ies(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) diff --git a/intern/cycles/kernel/svm/svm_image.h b/intern/cycles/kernel/svm/svm_image.h index 2de80d5fc29..68374fcfb0d 100644 --- a/intern/cycles/kernel/svm/svm_image.h +++ b/intern/cycles/kernel/svm/svm_image.h @@ -16,8 +16,7 @@ CCL_NAMESPACE_BEGIN -ccl_device float4 -svm_image_texture(ccl_global const KernelGlobals *kg, int id, float x, float y, uint flags) +ccl_device float4 svm_image_texture(KernelGlobals kg, int id, float x, float y, uint flags) { if (id == -1) { return make_float4( @@ -45,11 +44,8 @@ ccl_device_inline float3 texco_remap_square(float3 co) return (co - make_float3(0.5f, 0.5f, 0.5f)) * 2.0f; } -ccl_device_noinline int svm_node_tex_image(ccl_global const KernelGlobals *kg, - ccl_private ShaderData *sd, - ccl_private float *stack, - uint4 node, - int offset) +ccl_device_noinline int svm_node_tex_image( + KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset) { uint co_offset, out_offset, alpha_offset, flags; @@ -121,7 +117,7 @@ ccl_device_noinline int svm_node_tex_image(ccl_global const KernelGlobals *kg, return offset; } -ccl_device_noinline void svm_node_tex_image_box(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_tex_image_box(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) @@ -223,7 +219,7 @@ ccl_device_noinline void svm_node_tex_image_box(ccl_global const KernelGlobals * stack_store_float(stack, alpha_offset, f.w); } -ccl_device_noinline void svm_node_tex_environment(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_tex_environment(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) diff --git a/intern/cycles/kernel/svm/svm_light_path.h b/intern/cycles/kernel/svm/svm_light_path.h index aaff8376c7c..5e1fc4f671c 100644 --- a/intern/cycles/kernel/svm/svm_light_path.h +++ b/intern/cycles/kernel/svm/svm_light_path.h @@ -18,12 +18,14 @@ CCL_NAMESPACE_BEGIN /* Light Path Node */ -ccl_device_noinline void svm_node_light_path(INTEGRATOR_STATE_CONST_ARGS, +template<uint node_feature_mask, typename ConstIntegratorGenericState> +ccl_device_noinline void svm_node_light_path(KernelGlobals kg, + ConstIntegratorGenericState state, ccl_private const ShaderData *sd, ccl_private float *stack, uint type, uint out_offset, - int path_flag) + uint32_t path_flag) { float info = 0.0f; @@ -62,43 +64,43 @@ ccl_device_noinline void svm_node_light_path(INTEGRATOR_STATE_CONST_ARGS, /* Read bounce from difference location depending if this is a shadow * path. It's a bit dubious to have integrate state details leak into * this function but hard to avoid currently. */ - int bounce = (INTEGRATOR_STATE_IS_NULL) ? 0 : - (path_flag & PATH_RAY_SHADOW) ? INTEGRATOR_STATE(shadow_path, bounce) : - INTEGRATOR_STATE(path, bounce); + IF_KERNEL_NODES_FEATURE(LIGHT_PATH) + { + info = (float)integrator_state_bounce(state, path_flag); + } /* For background, light emission and shadow evaluation we from a * surface or volume we are effective one bounce further. */ if (path_flag & (PATH_RAY_SHADOW | PATH_RAY_EMISSION)) { - bounce++; + info += 1.0f; } - - info = (float)bounce; break; } - /* TODO */ case NODE_LP_ray_transparent: { - const int bounce = (INTEGRATOR_STATE_IS_NULL) ? - 0 : - (path_flag & PATH_RAY_SHADOW) ? - INTEGRATOR_STATE(shadow_path, transparent_bounce) : - INTEGRATOR_STATE(path, transparent_bounce); - - info = (float)bounce; + IF_KERNEL_NODES_FEATURE(LIGHT_PATH) + { + info = (float)integrator_state_transparent_bounce(state, path_flag); + } break; } -#if 0 case NODE_LP_ray_diffuse: - info = (float)state->diffuse_bounce; + IF_KERNEL_NODES_FEATURE(LIGHT_PATH) + { + info = (float)integrator_state_diffuse_bounce(state, path_flag); + } break; case NODE_LP_ray_glossy: - info = (float)state->glossy_bounce; + IF_KERNEL_NODES_FEATURE(LIGHT_PATH) + { + info = (float)integrator_state_glossy_bounce(state, path_flag); + } break; -#endif -#if 0 case NODE_LP_ray_transmission: - info = (float)state->transmission_bounce; + IF_KERNEL_NODES_FEATURE(LIGHT_PATH) + { + info = (float)integrator_state_transmission_bounce(state, path_flag); + } break; -#endif } stack_store_float(stack, out_offset, info); diff --git a/intern/cycles/kernel/svm/svm_magic.h b/intern/cycles/kernel/svm/svm_magic.h index 4c4f3bcf523..d3a429fec56 100644 --- a/intern/cycles/kernel/svm/svm_magic.h +++ b/intern/cycles/kernel/svm/svm_magic.h @@ -87,11 +87,8 @@ ccl_device_noinline_cpu float3 svm_magic(float3 p, int n, float distortion) return make_float3(0.5f - x, 0.5f - y, 0.5f - z); } -ccl_device_noinline int svm_node_tex_magic(ccl_global const KernelGlobals *kg, - ccl_private ShaderData *sd, - ccl_private float *stack, - uint4 node, - int offset) +ccl_device_noinline int svm_node_tex_magic( + KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset) { uint depth; uint scale_offset, distortion_offset, co_offset, fac_offset, color_offset; diff --git a/intern/cycles/kernel/svm/svm_map_range.h b/intern/cycles/kernel/svm/svm_map_range.h index f4f7d3ca76f..5e89947c6c7 100644 --- a/intern/cycles/kernel/svm/svm_map_range.h +++ b/intern/cycles/kernel/svm/svm_map_range.h @@ -24,7 +24,7 @@ ccl_device_inline float smootherstep(float edge0, float edge1, float x) return x * x * x * (x * (x * 6.0f - 15.0f) + 10.0f); } -ccl_device_noinline int svm_node_map_range(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_map_range(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint value_stack_offset, diff --git a/intern/cycles/kernel/svm/svm_mapping.h b/intern/cycles/kernel/svm/svm_mapping.h index 8102afc637e..ed420e5bc3d 100644 --- a/intern/cycles/kernel/svm/svm_mapping.h +++ b/intern/cycles/kernel/svm/svm_mapping.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Mapping Node */ -ccl_device_noinline void svm_node_mapping(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_mapping(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint type, @@ -43,7 +43,7 @@ ccl_device_noinline void svm_node_mapping(ccl_global const KernelGlobals *kg, /* Texture Mapping */ -ccl_device_noinline int svm_node_texture_mapping(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_texture_mapping(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint vec_offset, @@ -62,7 +62,7 @@ ccl_device_noinline int svm_node_texture_mapping(ccl_global const KernelGlobals return offset; } -ccl_device_noinline int svm_node_min_max(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_min_max(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint vec_offset, diff --git a/intern/cycles/kernel/svm/svm_math.h b/intern/cycles/kernel/svm/svm_math.h index 3897a453873..97f7d486c09 100644 --- a/intern/cycles/kernel/svm/svm_math.h +++ b/intern/cycles/kernel/svm/svm_math.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -ccl_device_noinline void svm_node_math(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_math(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint type, @@ -34,7 +34,7 @@ ccl_device_noinline void svm_node_math(ccl_global const KernelGlobals *kg, stack_store_float(stack, result_stack_offset, result); } -ccl_device_noinline int svm_node_vector_math(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_vector_math(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint type, diff --git a/intern/cycles/kernel/svm/svm_mix.h b/intern/cycles/kernel/svm/svm_mix.h index 0064c5e643c..568dda3dddc 100644 --- a/intern/cycles/kernel/svm/svm_mix.h +++ b/intern/cycles/kernel/svm/svm_mix.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Node */ -ccl_device_noinline int svm_node_mix(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_mix(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint fac_offset, diff --git a/intern/cycles/kernel/svm/svm_musgrave.h b/intern/cycles/kernel/svm/svm_musgrave.h index 8523f45b95f..decd29bbe13 100644 --- a/intern/cycles/kernel/svm/svm_musgrave.h +++ b/intern/cycles/kernel/svm/svm_musgrave.h @@ -700,7 +700,7 @@ ccl_device_noinline_cpu float noise_musgrave_ridged_multi_fractal_4d( return value; } -ccl_device_noinline int svm_node_tex_musgrave(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_tex_musgrave(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint offsets1, diff --git a/intern/cycles/kernel/svm/svm_noisetex.h b/intern/cycles/kernel/svm/svm_noisetex.h index 61da8227efa..3fe33f72b59 100644 --- a/intern/cycles/kernel/svm/svm_noisetex.h +++ b/intern/cycles/kernel/svm/svm_noisetex.h @@ -140,7 +140,7 @@ ccl_device void noise_texture_4d(float4 co, } } -ccl_device_noinline int svm_node_tex_noise(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_tex_noise(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint dimensions, diff --git a/intern/cycles/kernel/svm/svm_normal.h b/intern/cycles/kernel/svm/svm_normal.h index 0d1b4200d54..9bf64ed8823 100644 --- a/intern/cycles/kernel/svm/svm_normal.h +++ b/intern/cycles/kernel/svm/svm_normal.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -ccl_device_noinline int svm_node_normal(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_normal(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint in_normal_offset, diff --git a/intern/cycles/kernel/svm/svm_ramp.h b/intern/cycles/kernel/svm/svm_ramp.h index ef8b0d103c1..d2dddf4c6eb 100644 --- a/intern/cycles/kernel/svm/svm_ramp.h +++ b/intern/cycles/kernel/svm/svm_ramp.h @@ -21,18 +21,14 @@ CCL_NAMESPACE_BEGIN /* NOTE: svm_ramp.h, svm_ramp_util.h and node_ramp_util.h must stay consistent */ -ccl_device_inline float fetch_float(ccl_global const KernelGlobals *kg, int offset) +ccl_device_inline float fetch_float(KernelGlobals kg, int offset) { uint4 node = kernel_tex_fetch(__svm_nodes, offset); return __uint_as_float(node.x); } -ccl_device_inline float float_ramp_lookup(ccl_global const KernelGlobals *kg, - int offset, - float f, - bool interpolate, - bool extrapolate, - int table_size) +ccl_device_inline float float_ramp_lookup( + KernelGlobals kg, int offset, float f, bool interpolate, bool extrapolate, int table_size) { if ((f < 0.0f || f > 1.0f) && extrapolate) { float t0, dy; @@ -63,12 +59,8 @@ ccl_device_inline float float_ramp_lookup(ccl_global const KernelGlobals *kg, return a; } -ccl_device_inline float4 rgb_ramp_lookup(ccl_global const KernelGlobals *kg, - int offset, - float f, - bool interpolate, - bool extrapolate, - int table_size) +ccl_device_inline float4 rgb_ramp_lookup( + KernelGlobals kg, int offset, float f, bool interpolate, bool extrapolate, int table_size) { if ((f < 0.0f || f > 1.0f) && extrapolate) { float4 t0, dy; @@ -99,11 +91,8 @@ ccl_device_inline float4 rgb_ramp_lookup(ccl_global const KernelGlobals *kg, return a; } -ccl_device_noinline int svm_node_rgb_ramp(ccl_global const KernelGlobals *kg, - ccl_private ShaderData *sd, - ccl_private float *stack, - uint4 node, - int offset) +ccl_device_noinline int svm_node_rgb_ramp( + KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset) { uint fac_offset, color_offset, alpha_offset; uint interpolate = node.z; @@ -124,11 +113,8 @@ ccl_device_noinline int svm_node_rgb_ramp(ccl_global const KernelGlobals *kg, return offset; } -ccl_device_noinline int svm_node_curves(ccl_global const KernelGlobals *kg, - ccl_private ShaderData *sd, - ccl_private float *stack, - uint4 node, - int offset) +ccl_device_noinline int svm_node_curves( + KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset) { uint fac_offset, color_offset, out_offset; svm_unpack_node_uchar3(node.y, &fac_offset, &color_offset, &out_offset); @@ -153,11 +139,8 @@ ccl_device_noinline int svm_node_curves(ccl_global const KernelGlobals *kg, return offset; } -ccl_device_noinline int svm_node_curve(ccl_global const KernelGlobals *kg, - ccl_private ShaderData *sd, - ccl_private float *stack, - uint4 node, - int offset) +ccl_device_noinline int svm_node_curve( + KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset) { uint fac_offset, value_in_offset, out_offset; svm_unpack_node_uchar3(node.y, &fac_offset, &value_in_offset, &out_offset); diff --git a/intern/cycles/kernel/svm/svm_sepcomb_hsv.h b/intern/cycles/kernel/svm/svm_sepcomb_hsv.h index 3cd4ba87a55..bafa0456342 100644 --- a/intern/cycles/kernel/svm/svm_sepcomb_hsv.h +++ b/intern/cycles/kernel/svm/svm_sepcomb_hsv.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -ccl_device_noinline int svm_node_combine_hsv(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_combine_hsv(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint hue_in, @@ -39,7 +39,7 @@ ccl_device_noinline int svm_node_combine_hsv(ccl_global const KernelGlobals *kg, return offset; } -ccl_device_noinline int svm_node_separate_hsv(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_separate_hsv(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint color_in, diff --git a/intern/cycles/kernel/svm/svm_sky.h b/intern/cycles/kernel/svm/svm_sky.h index 04db8109170..3ab7bc89c66 100644 --- a/intern/cycles/kernel/svm/svm_sky.h +++ b/intern/cycles/kernel/svm/svm_sky.h @@ -37,7 +37,7 @@ ccl_device float sky_perez_function(ccl_private float *lam, float theta, float g (1.0f + lam[2] * expf(lam[3] * gamma) + lam[4] * cgamma * cgamma); } -ccl_device float3 sky_radiance_preetham(ccl_global const KernelGlobals *kg, +ccl_device float3 sky_radiance_preetham(KernelGlobals kg, float3 dir, float sunphi, float suntheta, @@ -90,7 +90,7 @@ ccl_device float sky_radiance_internal(ccl_private float *configuration, float t configuration[6] * mieM + configuration[7] * zenith); } -ccl_device float3 sky_radiance_hosek(ccl_global const KernelGlobals *kg, +ccl_device float3 sky_radiance_hosek(KernelGlobals kg, float3 dir, float sunphi, float suntheta, @@ -127,7 +127,7 @@ ccl_device float3 geographical_to_direction(float lat, float lon) return make_float3(cos(lat) * cos(lon), cos(lat) * sin(lon), sin(lat)); } -ccl_device float3 sky_radiance_nishita(ccl_global const KernelGlobals *kg, +ccl_device float3 sky_radiance_nishita(KernelGlobals kg, float3 dir, ccl_private float *nishita_data, uint texture_id) @@ -209,11 +209,8 @@ ccl_device float3 sky_radiance_nishita(ccl_global const KernelGlobals *kg, return xyz_to_rgb(kg, xyz); } -ccl_device_noinline int svm_node_tex_sky(ccl_global const KernelGlobals *kg, - ccl_private ShaderData *sd, - ccl_private float *stack, - uint4 node, - int offset) +ccl_device_noinline int svm_node_tex_sky( + KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset) { /* Load data */ uint dir_offset = node.y; diff --git a/intern/cycles/kernel/svm/svm_tex_coord.h b/intern/cycles/kernel/svm/svm_tex_coord.h index 295d5e9f65b..fe777eb34c8 100644 --- a/intern/cycles/kernel/svm/svm_tex_coord.h +++ b/intern/cycles/kernel/svm/svm_tex_coord.h @@ -22,9 +22,9 @@ CCL_NAMESPACE_BEGIN /* Texture Coordinate Node */ -ccl_device_noinline int svm_node_tex_coord(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_tex_coord(KernelGlobals kg, ccl_private ShaderData *sd, - int path_flag, + uint32_t path_flag, ccl_private float *stack, uint4 node, int offset) @@ -103,9 +103,9 @@ ccl_device_noinline int svm_node_tex_coord(ccl_global const KernelGlobals *kg, return offset; } -ccl_device_noinline int svm_node_tex_coord_bump_dx(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_tex_coord_bump_dx(KernelGlobals kg, ccl_private ShaderData *sd, - int path_flag, + uint32_t path_flag, ccl_private float *stack, uint4 node, int offset) @@ -188,9 +188,9 @@ ccl_device_noinline int svm_node_tex_coord_bump_dx(ccl_global const KernelGlobal #endif } -ccl_device_noinline int svm_node_tex_coord_bump_dy(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_tex_coord_bump_dy(KernelGlobals kg, ccl_private ShaderData *sd, - int path_flag, + uint32_t path_flag, ccl_private float *stack, uint4 node, int offset) @@ -273,7 +273,7 @@ ccl_device_noinline int svm_node_tex_coord_bump_dy(ccl_global const KernelGlobal #endif } -ccl_device_noinline void svm_node_normal_map(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_normal_map(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) @@ -366,7 +366,7 @@ ccl_device_noinline void svm_node_normal_map(ccl_global const KernelGlobals *kg, stack_store_float3(stack, normal_offset, N); } -ccl_device_noinline void svm_node_tangent(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_tangent(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) diff --git a/intern/cycles/kernel/svm/svm_value.h b/intern/cycles/kernel/svm/svm_value.h index d1038bc072d..cc72961d0f6 100644 --- a/intern/cycles/kernel/svm/svm_value.h +++ b/intern/cycles/kernel/svm/svm_value.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Value Nodes */ -ccl_device void svm_node_value_f(ccl_global const KernelGlobals *kg, +ccl_device void svm_node_value_f(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint ivalue, @@ -27,7 +27,7 @@ ccl_device void svm_node_value_f(ccl_global const KernelGlobals *kg, stack_store_float(stack, out_offset, __uint_as_float(ivalue)); } -ccl_device int svm_node_value_v(ccl_global const KernelGlobals *kg, +ccl_device int svm_node_value_v(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint out_offset, diff --git a/intern/cycles/kernel/svm/svm_vector_transform.h b/intern/cycles/kernel/svm/svm_vector_transform.h index b6c898c3952..4e0d36647da 100644 --- a/intern/cycles/kernel/svm/svm_vector_transform.h +++ b/intern/cycles/kernel/svm/svm_vector_transform.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Vector Transform */ -ccl_device_noinline void svm_node_vector_transform(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_vector_transform(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) diff --git a/intern/cycles/kernel/svm/svm_vertex_color.h b/intern/cycles/kernel/svm/svm_vertex_color.h index 3641f05ca43..a5fa15ee085 100644 --- a/intern/cycles/kernel/svm/svm_vertex_color.h +++ b/intern/cycles/kernel/svm/svm_vertex_color.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -ccl_device_noinline void svm_node_vertex_color(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_vertex_color(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint layer_id, @@ -35,7 +35,7 @@ ccl_device_noinline void svm_node_vertex_color(ccl_global const KernelGlobals *k } } -ccl_device_noinline void svm_node_vertex_color_bump_dx(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_vertex_color_bump_dx(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint layer_id, @@ -56,7 +56,7 @@ ccl_device_noinline void svm_node_vertex_color_bump_dx(ccl_global const KernelGl } } -ccl_device_noinline void svm_node_vertex_color_bump_dy(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_vertex_color_bump_dy(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint layer_id, diff --git a/intern/cycles/kernel/svm/svm_voronoi.h b/intern/cycles/kernel/svm/svm_voronoi.h index e7112087e17..b8067520770 100644 --- a/intern/cycles/kernel/svm/svm_voronoi.h +++ b/intern/cycles/kernel/svm/svm_voronoi.h @@ -917,7 +917,7 @@ ccl_device void voronoi_n_sphere_radius_4d(float4 coord, } template<uint node_feature_mask> -ccl_device_noinline int svm_node_tex_voronoi(ccl_global const KernelGlobals *kg, +ccl_device_noinline int svm_node_tex_voronoi(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint dimensions, @@ -1001,7 +1001,7 @@ ccl_device_noinline int svm_node_tex_voronoi(ccl_global const KernelGlobals *kg, } case 2: { float2 coord_2d = make_float2(coord.x, coord.y); - float2 position_out_2d; + float2 position_out_2d = zero_float2(); switch (voronoi_feature) { case NODE_VORONOI_F1: voronoi_f1_2d(coord_2d, @@ -1013,7 +1013,8 @@ ccl_device_noinline int svm_node_tex_voronoi(ccl_global const KernelGlobals *kg, &position_out_2d); break; case NODE_VORONOI_SMOOTH_F1: - if (KERNEL_NODES_FEATURE(VORONOI_EXTRA)) { + IF_KERNEL_NODES_FEATURE(VORONOI_EXTRA) + { voronoi_smooth_f1_2d(coord_2d, smoothness, exponent, @@ -1058,7 +1059,8 @@ ccl_device_noinline int svm_node_tex_voronoi(ccl_global const KernelGlobals *kg, &position_out); break; case NODE_VORONOI_SMOOTH_F1: - if (KERNEL_NODES_FEATURE(VORONOI_EXTRA)) { + IF_KERNEL_NODES_FEATURE(VORONOI_EXTRA) + { voronoi_smooth_f1_3d(coord, smoothness, exponent, @@ -1092,7 +1094,8 @@ ccl_device_noinline int svm_node_tex_voronoi(ccl_global const KernelGlobals *kg, } case 4: { - if (KERNEL_NODES_FEATURE(VORONOI_EXTRA)) { + IF_KERNEL_NODES_FEATURE(VORONOI_EXTRA) + { float4 coord_4d = make_float4(coord.x, coord.y, coord.z, w); float4 position_out_4d; switch (voronoi_feature) { diff --git a/intern/cycles/kernel/svm/svm_voxel.h b/intern/cycles/kernel/svm/svm_voxel.h index 764fb71ba72..be4bb315145 100644 --- a/intern/cycles/kernel/svm/svm_voxel.h +++ b/intern/cycles/kernel/svm/svm_voxel.h @@ -19,11 +19,8 @@ CCL_NAMESPACE_BEGIN /* TODO(sergey): Think of making it more generic volume-type attribute * sampler. */ -ccl_device_noinline int svm_node_tex_voxel(ccl_global const KernelGlobals *kg, - ccl_private ShaderData *sd, - ccl_private float *stack, - uint4 node, - int offset) +ccl_device_noinline int svm_node_tex_voxel( + KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset) { uint co_offset, density_out_offset, color_out_offset, space; svm_unpack_node_uchar4(node.z, &co_offset, &density_out_offset, &color_out_offset, &space); diff --git a/intern/cycles/kernel/svm/svm_wave.h b/intern/cycles/kernel/svm/svm_wave.h index 1ac130e2006..d04b7aa3476 100644 --- a/intern/cycles/kernel/svm/svm_wave.h +++ b/intern/cycles/kernel/svm/svm_wave.h @@ -82,11 +82,8 @@ ccl_device_noinline_cpu float svm_wave(NodeWaveType type, } } -ccl_device_noinline int svm_node_tex_wave(ccl_global const KernelGlobals *kg, - ccl_private ShaderData *sd, - ccl_private float *stack, - uint4 node, - int offset) +ccl_device_noinline int svm_node_tex_wave( + KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset) { uint4 node2 = read_node(kg, &offset); uint4 node3 = read_node(kg, &offset); diff --git a/intern/cycles/kernel/svm/svm_wavelength.h b/intern/cycles/kernel/svm/svm_wavelength.h index e891744f276..4ef041f68d5 100644 --- a/intern/cycles/kernel/svm/svm_wavelength.h +++ b/intern/cycles/kernel/svm/svm_wavelength.h @@ -34,7 +34,7 @@ CCL_NAMESPACE_BEGIN /* Wavelength to RGB */ -ccl_device_noinline void svm_node_wavelength(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_wavelength(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint wavelength, diff --git a/intern/cycles/kernel/svm/svm_white_noise.h b/intern/cycles/kernel/svm/svm_white_noise.h index ccc49bf1a7c..6c2c3d6a683 100644 --- a/intern/cycles/kernel/svm/svm_white_noise.h +++ b/intern/cycles/kernel/svm/svm_white_noise.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -ccl_device_noinline void svm_node_tex_white_noise(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_tex_white_noise(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint dimensions, diff --git a/intern/cycles/kernel/svm/svm_wireframe.h b/intern/cycles/kernel/svm/svm_wireframe.h index 70d1211aa4a..d75976d23e1 100644 --- a/intern/cycles/kernel/svm/svm_wireframe.h +++ b/intern/cycles/kernel/svm/svm_wireframe.h @@ -34,7 +34,7 @@ CCL_NAMESPACE_BEGIN /* Wireframe Node */ -ccl_device_inline float wireframe(ccl_global const KernelGlobals *kg, +ccl_device_inline float wireframe(KernelGlobals kg, ccl_private ShaderData *sd, float size, int pixel_size, @@ -91,7 +91,7 @@ ccl_device_inline float wireframe(ccl_global const KernelGlobals *kg, return 0.0f; } -ccl_device_noinline void svm_node_wireframe(ccl_global const KernelGlobals *kg, +ccl_device_noinline void svm_node_wireframe(KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node) |