[Bf-blender-cvs] [2bed80609a8] cycles-x: Cycles: add initial specialization of svm_eval_nodes
Brecht Van Lommel
noreply at git.blender.org
Fri Apr 30 17:38:07 CEST 2021
Commit: 2bed80609a87acd6886aace21cfebf86c31c3d84
Author: Brecht Van Lommel
Date: Thu Apr 29 21:23:55 2021 +0200
Branches: cycles-x
https://developer.blender.org/rB2bed80609a87acd6886aace21cfebf86c31c3d84
Cycles: add initial specialization of svm_eval_nodes
To leave out shaders nodes from rendering kernels where they are known to
be unused. Reduces GPU render time between 3-12% depending on the scene in
our tests.
Differential Revision: https://developer.blender.org/D11131
===================================================================
M intern/cycles/kernel/integrator/integrator_shade_background.h
M intern/cycles/kernel/integrator/integrator_shade_shadow.h
M intern/cycles/kernel/integrator/integrator_shade_surface.h
M intern/cycles/kernel/kernel_bake.h
M intern/cycles/kernel/kernel_emission.h
M intern/cycles/kernel/kernel_shader.h
M intern/cycles/kernel/svm/svm.h
M intern/cycles/kernel/svm/svm_attribute.h
M intern/cycles/kernel/svm/svm_types.h
M intern/cycles/kernel/svm/svm_voronoi.h
M intern/cycles/render/shader.cpp
===================================================================
diff --git a/intern/cycles/kernel/integrator/integrator_shade_background.h b/intern/cycles/kernel/integrator/integrator_shade_background.h
index 1319c812494..4401a27935b 100644
--- a/intern/cycles/kernel/integrator/integrator_shade_background.h
+++ b/intern/cycles/kernel/integrator/integrator_shade_background.h
@@ -59,7 +59,7 @@ ccl_device_noinline_cpu float3 integrator_eval_background_shader(
INTEGRATOR_STATE(ray, P),
INTEGRATOR_STATE(ray, D),
INTEGRATOR_STATE(ray, time));
- shader_eval_surface(
+ shader_eval_surface<NODE_FEATURE_MASK_LIGHT>(
INTEGRATOR_STATE_PASS, emission_sd, render_buffer, path_flag | PATH_RAY_EMISSION);
L = shader_background_eval(emission_sd);
diff --git a/intern/cycles/kernel/integrator/integrator_shade_shadow.h b/intern/cycles/kernel/integrator/integrator_shade_shadow.h
index 0017a398343..4abd510ca52 100644
--- a/intern/cycles/kernel/integrator/integrator_shade_shadow.h
+++ b/intern/cycles/kernel/integrator/integrator_shade_shadow.h
@@ -41,7 +41,8 @@ ccl_device_inline float3 integrate_transparent_shadow_shader_eval(INTEGRATOR_STA
/* Evaluate shader. */
if (!(shadow_sd->flag & SD_HAS_ONLY_VOLUME)) {
- shader_eval_surface(INTEGRATOR_STATE_PASS, shadow_sd, NULL, PATH_RAY_SHADOW);
+ shader_eval_surface<NODE_FEATURE_MASK_SURFACE>(
+ INTEGRATOR_STATE_PASS, shadow_sd, NULL, PATH_RAY_SHADOW);
}
/* Compute transparency from closures. */
diff --git a/intern/cycles/kernel/integrator/integrator_shade_surface.h b/intern/cycles/kernel/integrator/integrator_shade_surface.h
index 05ea0b0c082..1b54e9e7f9e 100644
--- a/intern/cycles/kernel/integrator/integrator_shade_surface.h
+++ b/intern/cycles/kernel/integrator/integrator_shade_surface.h
@@ -330,7 +330,8 @@ ccl_device_inline bool integrate_surface(INTEGRATOR_STATE_ARGS,
#endif
{
/* Evaluate shader. */
- shader_eval_surface(INTEGRATOR_STATE_PASS, &sd, render_buffer, path_flag);
+ shader_eval_surface<NODE_FEATURE_MASK_SURFACE>(
+ INTEGRATOR_STATE_PASS, &sd, render_buffer, path_flag);
}
#ifdef __SUBSURFACE__
diff --git a/intern/cycles/kernel/kernel_bake.h b/intern/cycles/kernel/kernel_bake.h
index 058371c5f4c..70db3c193cc 100644
--- a/intern/cycles/kernel/kernel_bake.h
+++ b/intern/cycles/kernel/kernel_bake.h
@@ -455,7 +455,7 @@ ccl_device void kernel_background_evaluate(const KernelGlobals *kg,
/* Evaluate shader.
* This is being evaluated for all BSDFs, so path flag does not contain a specific type. */
const int path_flag = PATH_RAY_EMISSION;
- shader_eval_surface(INTEGRATOR_STATE_PASS_NULL, &sd, NULL, path_flag);
+ shader_eval_surface<NODE_FEATURE_MASK_LIGHT>(INTEGRATOR_STATE_PASS_NULL, &sd, NULL, path_flag);
const float3 color = shader_background_eval(&sd);
/* Write output. */
diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h
index 9196d88104f..a0f6cc3cb01 100644
--- a/intern/cycles/kernel/kernel_emission.h
+++ b/intern/cycles/kernel/kernel_emission.h
@@ -67,7 +67,8 @@ ccl_device_noinline_cpu float3 light_sample_shader_eval(INTEGRATOR_STATE_ARGS,
/* No proper path flag, we're evaluating this for all closures. that's
* weak but we'd have to do multiple evaluations otherwise. */
- shader_eval_surface(INTEGRATOR_STATE_PASS, emission_sd, NULL, PATH_RAY_EMISSION);
+ shader_eval_surface<NODE_FEATURE_MASK_LIGHT>(
+ INTEGRATOR_STATE_PASS, emission_sd, NULL, PATH_RAY_EMISSION);
/* Evaluate closures. */
#ifdef __BACKGROUND_MIS__
diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h
index e5026926d54..7f9fdbace01 100644
--- a/intern/cycles/kernel/kernel_shader.h
+++ b/intern/cycles/kernel/kernel_shader.h
@@ -1025,6 +1025,7 @@ ccl_device float3 shader_holdout_apply(const KernelGlobals *kg, ShaderData *sd)
/* Surface Evaluation */
+template<uint node_feature_mask>
ccl_device void shader_eval_surface(INTEGRATOR_STATE_CONST_ARGS,
ShaderData *sd,
ccl_global float *buffer,
@@ -1059,7 +1060,8 @@ ccl_device void shader_eval_surface(INTEGRATOR_STATE_CONST_ARGS,
#endif
{
#ifdef __SVM__
- svm_eval_nodes(INTEGRATOR_STATE_PASS, sd, buffer, SHADER_TYPE_SURFACE, path_flag);
+ svm_eval_nodes<node_feature_mask, SHADER_TYPE_SURFACE>(
+ INTEGRATOR_STATE_PASS, sd, buffer, path_flag);
#else
if (sd->object == OBJECT_NONE) {
sd->closure_emission_background = make_float3(0.8f, 0.8f, 0.8f);
@@ -1076,7 +1078,7 @@ ccl_device void shader_eval_surface(INTEGRATOR_STATE_CONST_ARGS,
#endif
}
- if (sd->flag & SD_BSDF_NEEDS_LCG) {
+ if (NODES_FEATURE(BSDF) && (sd->flag & SD_BSDF_NEEDS_LCG)) {
sd->lcg_state = lcg_state_init(INTEGRATOR_STATE(path, rng_hash),
INTEGRATOR_STATE(path, rng_offset),
INTEGRATOR_STATE(path, sample),
@@ -1275,7 +1277,8 @@ ccl_device_inline void shader_eval_volume(INTEGRATOR_STATE_CONST_ARGS,
else
# endif
{
- svm_eval_nodes(INTEGRATOR_STATE_PASS, sd, NULL, SHADER_TYPE_VOLUME, path_flag);
+ svm_eval_nodes<NODE_FEATURE_MASK_VOLUME, SHADER_TYPE_VOLUME>(
+ INTEGRATOR_STATE_PASS, sd, NULL, path_flag);
}
# endif
@@ -1302,7 +1305,8 @@ ccl_device void shader_eval_displacement(INTEGRATOR_STATE_CONST_ARGS, ShaderData
else
# endif
{
- svm_eval_nodes(INTEGRATOR_STATE_PASS, sd, NULL, SHADER_TYPE_DISPLACEMENT, 0);
+ svm_eval_nodes<NODE_FEATURE_MASK_DISPLACEMENT, SHADER_TYPE_DISPLACEMENT>(
+ INTEGRATOR_STATE_PASS, sd, NULL, 0);
}
#endif
}
diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h
index a49199a1cec..aee6e2a5823 100644
--- a/intern/cycles/kernel/svm/svm.h
+++ b/intern/cycles/kernel/svm/svm.h
@@ -230,10 +230,10 @@ ccl_device_inline void svm_eval_nodes(const KernelGlobals *kg,
extern "C" __device__ void __direct_callable__svm_eval_nodes(
#endif
+template<uint node_feature_mask, ShaderType type>
ccl_device_noinline void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS,
ShaderData *sd,
ccl_global float *buffer,
- ShaderType type,
int path_flag)
{
float stack[SVM_STACK_SIZE];
@@ -258,13 +258,19 @@ ccl_device_noinline void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS,
break;
}
case NODE_CLOSURE_BSDF:
- svm_node_closure_bsdf(kg, sd, stack, node, type, path_flag, &offset);
+ if (NODES_FEATURE(BSDF)) {
+ svm_node_closure_bsdf(kg, sd, stack, node, type, path_flag, &offset);
+ }
break;
case NODE_CLOSURE_EMISSION:
- svm_node_closure_emission(sd, stack, node);
+ if (NODES_FEATURE(EMISSION)) {
+ svm_node_closure_emission(sd, stack, node);
+ }
break;
case NODE_CLOSURE_BACKGROUND:
- svm_node_closure_background(sd, stack, node);
+ if (NODES_FEATURE(EMISSION)) {
+ svm_node_closure_background(sd, stack, node);
+ }
break;
case NODE_CLOSURE_SET_WEIGHT:
svm_node_closure_set_weight(sd, node.y, node.z, node.w);
@@ -273,7 +279,9 @@ ccl_device_noinline void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS,
svm_node_closure_weight(sd, stack, node.y);
break;
case NODE_EMISSION_WEIGHT:
- svm_node_emission_weight(kg, sd, stack, node);
+ if (NODES_FEATURE(EMISSION)) {
+ svm_node_emission_weight(kg, sd, stack, node);
+ }
break;
case NODE_MIX_CLOSURE:
svm_node_mix_closure(sd, stack, node);
@@ -302,28 +310,36 @@ ccl_device_noinline void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS,
svm_node_value_v(kg, sd, stack, node.y, &offset);
break;
case NODE_ATTR:
- svm_node_attr(kg, sd, stack, node);
+ svm_node_attr<node_feature_mask>(kg, sd, stack, node);
break;
case NODE_VERTEX_COLOR:
svm_node_vertex_color(kg, sd, stack, node.y, node.z, node.w);
break;
-# if NODES_FEATURE(NODE_FEATURE_BUMP)
case NODE_GEOMETRY_BUMP_DX:
- svm_node_geometry_bump_dx(kg, sd, stack, node.y, node.z);
+ if (NODES_FEATURE(BUMP)) {
+ svm_node_geometry_bump_dx(kg, sd, stack, node.y, node.z);
+ }
break;
case NODE_GEOMETRY_BUMP_DY:
- svm_node_geometry_bump_dy(kg, sd, stack, node.y, node.z);
+ if (NODES_FEATURE(BUMP)) {
+ svm_node_geometry_bump_dy(kg, sd, stack, node.y, node.z);
+ }
break;
case NODE_SET_DISPLACEMENT:
- svm_node_set_displacement(kg, sd, stack, node.y);
+ if (NODES_FEATURE(BUMP)) {
+ svm_node_set_displacement(kg, sd, stack, node.y);
+ }
break;
case NODE_DISPLACEMENT:
- svm_node_displacement(kg, sd, stack, node);
+ if (NODES_FEATURE(BUMP)) {
+ svm_node_displacement(kg, sd, stack, node);
+ }
break;
case NODE_VECTOR_DISPLACEMENT:
- svm_node_vector_displacement(kg, sd, stack, node, &offset);
+ if (NODES_FEATURE(BUMP)) {
+ svm_node_vector_displacement(kg, sd, stack, node, &offset);
+ }
break;
-# endif /* NODES_FEATURE(NODE_FEATURE_BUMP) */
case NODE_TEX_IMAGE:
svm_node_tex_image(kg, sd, stack, node, &offset);
break;
@@ -333,40 +349,56 @@ ccl_device_noinline void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS,
case NODE_TEX_NOISE:
svm_node_tex_noise(kg, sd, stack, node.y, node.z, node.w, &offset);
break;
-# if NODES_FEATURE(NODE_FEATURE_BUMP)
case NODE_SET_BUMP:
- svm_node_set_bump(kg, sd, stack, node);
+ if (NODES_FEATURE(BUMP)) {
+ svm_node_set_bump(kg, sd, stack, node);
+ }
break;
case NODE_ATTR_BUMP_DX:
- svm_node_attr_bump_dx(kg, sd, stack, node);
+ if (NODES_FEATURE(BUMP)) {
+ svm_node_attr_bump_dx(kg, sd, stack, node);
+ }
break;
case NODE_
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list