diff options
Diffstat (limited to 'intern/cycles/kernel/svm/svm.h')
| -rw-r--r-- | intern/cycles/kernel/svm/svm.h | 227 |
1 files changed, 120 insertions, 107 deletions
diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h index 000da1fa615..4aee1ef11b3 100644 --- a/intern/cycles/kernel/svm/svm.h +++ b/intern/cycles/kernel/svm/svm.h @@ -48,16 +48,18 @@ ccl_device_inline float3 stack_load_float3(float *stack, uint a) { kernel_assert(a + 2 < SVM_STACK_SIZE); - return make_float3(stack[a + 0], stack[a + 1], stack[a + 2]); + float *stack_a = stack + a; + return make_float3(stack_a[0], stack_a[1], stack_a[2]); } ccl_device_inline void stack_store_float3(float *stack, uint a, float3 f) { kernel_assert(a + 2 < SVM_STACK_SIZE); - stack[a + 0] = f.x; - stack[a + 1] = f.y; - stack[a + 2] = f.z; + float *stack_a = stack + a; + stack_a[0] = f.x; + stack_a[1] = f.y; + stack_a[2] = f.z; } ccl_device_inline float stack_load_float(float *stack, uint a) @@ -105,14 +107,14 @@ ccl_device_inline bool stack_valid(uint a) /* Reading Nodes */ -ccl_device_inline uint4 read_node(KernelGlobals *kg, int *offset) +ccl_device_inline uint4 read_node(const KernelGlobals *kg, int *offset) { uint4 node = kernel_tex_fetch(__svm_nodes, *offset); (*offset)++; return node; } -ccl_device_inline float4 read_node_float(KernelGlobals *kg, int *offset) +ccl_device_inline float4 read_node_float(const KernelGlobals *kg, int *offset) { uint4 node = kernel_tex_fetch(__svm_nodes, *offset); float4 f = make_float4(__uint_as_float(node.x), @@ -123,7 +125,7 @@ ccl_device_inline float4 read_node_float(KernelGlobals *kg, int *offset) return f; } -ccl_device_inline float4 fetch_node_float(KernelGlobals *kg, int offset) +ccl_device_inline float4 fetch_node_float(const KernelGlobals *kg, int offset) { uint4 node = kernel_tex_fetch(__svm_nodes, offset); return make_float4(__uint_as_float(node.x), @@ -217,26 +219,11 @@ CCL_NAMESPACE_END CCL_NAMESPACE_BEGIN /* Main Interpreter Loop */ -#if defined(__KERNEL_OPTIX__) && defined(__SHADER_RAYTRACE__) -ccl_device_inline void svm_eval_nodes(KernelGlobals *kg, - ShaderData *sd, - ccl_addr_space PathState *state, - ccl_global float *buffer, - ShaderType type, - int path_flag) -{ - optixDirectCall<void>(0, kg, sd, state, buffer, type, path_flag); -} -extern "C" __device__ void __direct_callable__svm_eval_nodes( -#else -ccl_device_noinline void svm_eval_nodes( -#endif - KernelGlobals *kg, - ShaderData *sd, - ccl_addr_space PathState *state, - ccl_global float *buffer, - ShaderType type, - int path_flag) +template<uint node_feature_mask, ShaderType type> +ccl_device void svm_eval_nodes(INTEGRATOR_STATE_CONST_ARGS, + ShaderData *sd, + ccl_global float *render_buffer, + int path_flag) { float stack[SVM_STACK_SIZE]; int offset = sd->shader & SHADER_MASK; @@ -247,7 +234,6 @@ ccl_device_noinline void svm_eval_nodes( switch (node.x) { case NODE_END: return; -#if NODES_GROUP(NODE_GROUP_LEVEL_0) case NODE_SHADER_JUMP: { if (type == SHADER_TYPE_SURFACE) offset = node.y; @@ -260,13 +246,18 @@ ccl_device_noinline void svm_eval_nodes( break; } case NODE_CLOSURE_BSDF: - svm_node_closure_bsdf(kg, sd, stack, node, type, path_flag, &offset); + offset = svm_node_closure_bsdf<node_feature_mask, type>( + kg, sd, stack, node, path_flag, offset); break; case NODE_CLOSURE_EMISSION: - svm_node_closure_emission(sd, stack, node); + if (KERNEL_NODES_FEATURE(EMISSION)) { + svm_node_closure_emission(sd, stack, node); + } break; case NODE_CLOSURE_BACKGROUND: - svm_node_closure_background(sd, stack, node); + if (KERNEL_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); @@ -275,7 +266,9 @@ ccl_device_noinline void svm_eval_nodes( svm_node_closure_weight(sd, stack, node.y); break; case NODE_EMISSION_WEIGHT: - svm_node_emission_weight(kg, sd, stack, node); + if (KERNEL_NODES_FEATURE(EMISSION)) { + svm_node_emission_weight(kg, sd, stack, node); + } break; case NODE_MIX_CLOSURE: svm_node_mix_closure(sd, stack, node); @@ -295,86 +288,108 @@ ccl_device_noinline void svm_eval_nodes( svm_node_convert(kg, sd, stack, node.y, node.z, node.w); break; case NODE_TEX_COORD: - svm_node_tex_coord(kg, sd, path_flag, stack, node, &offset); + offset = svm_node_tex_coord(kg, sd, path_flag, stack, node, offset); break; case NODE_VALUE_F: svm_node_value_f(kg, sd, stack, node.y, node.z); break; case NODE_VALUE_V: - svm_node_value_v(kg, sd, stack, node.y, &offset); + offset = 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 (KERNEL_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 (KERNEL_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 (KERNEL_NODES_FEATURE(BUMP)) { + svm_node_set_displacement(kg, sd, stack, node.y); + } break; case NODE_DISPLACEMENT: - svm_node_displacement(kg, sd, stack, node); + if (KERNEL_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 (KERNEL_NODES_FEATURE(BUMP)) { + offset = 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); + offset = svm_node_tex_image(kg, sd, stack, node, offset); break; case NODE_TEX_IMAGE_BOX: svm_node_tex_image_box(kg, sd, stack, node); break; case NODE_TEX_NOISE: - svm_node_tex_noise(kg, sd, stack, node.y, node.z, node.w, &offset); + offset = 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 (KERNEL_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 (KERNEL_NODES_FEATURE(BUMP)) { + svm_node_attr_bump_dx(kg, sd, stack, node); + } break; case NODE_ATTR_BUMP_DY: - svm_node_attr_bump_dy(kg, sd, stack, node); + if (KERNEL_NODES_FEATURE(BUMP)) { + svm_node_attr_bump_dy(kg, sd, stack, node); + } break; case NODE_VERTEX_COLOR_BUMP_DX: - svm_node_vertex_color_bump_dx(kg, sd, stack, node.y, node.z, node.w); + 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: - svm_node_vertex_color_bump_dy(kg, sd, stack, node.y, node.z, node.w); + 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: - svm_node_tex_coord_bump_dx(kg, sd, path_flag, stack, node, &offset); + 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: - svm_node_tex_coord_bump_dy(kg, sd, path_flag, stack, node, &offset); + 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: - svm_node_set_normal(kg, sd, stack, node.y, node.z); + if (KERNEL_NODES_FEATURE(BUMP)) { + svm_node_set_normal(kg, sd, stack, node.y, node.z); + } break; -# if NODES_FEATURE(NODE_FEATURE_BUMP_STATE) case NODE_ENTER_BUMP_EVAL: - svm_node_enter_bump_eval(kg, sd, stack, node.y); + if (KERNEL_NODES_FEATURE(BUMP_STATE)) { + svm_node_enter_bump_eval(kg, sd, stack, node.y); + } break; case NODE_LEAVE_BUMP_EVAL: - svm_node_leave_bump_eval(kg, sd, stack, node.y); + if (KERNEL_NODES_FEATURE(BUMP_STATE)) { + svm_node_leave_bump_eval(kg, sd, stack, node.y); + } break; -# endif /* NODES_FEATURE(NODE_FEATURE_BUMP_STATE) */ -# endif /* NODES_FEATURE(NODE_FEATURE_BUMP) */ case NODE_HSV: - svm_node_hsv(kg, sd, stack, node, &offset); + svm_node_hsv(kg, sd, stack, node); break; -#endif /* NODES_GROUP(NODE_GROUP_LEVEL_0) */ -#if NODES_GROUP(NODE_GROUP_LEVEL_1) case NODE_CLOSURE_HOLDOUT: svm_node_closure_holdout(sd, stack, node); break; @@ -384,22 +399,24 @@ ccl_device_noinline void svm_eval_nodes( case NODE_LAYER_WEIGHT: svm_node_layer_weight(sd, stack, node); break; -# if NODES_FEATURE(NODE_FEATURE_VOLUME) case NODE_CLOSURE_VOLUME: - svm_node_closure_volume(kg, sd, stack, node, type); + if (KERNEL_NODES_FEATURE(VOLUME)) { + svm_node_closure_volume<type>(kg, sd, stack, node); + } break; case NODE_PRINCIPLED_VOLUME: - svm_node_principled_volume(kg, sd, stack, node, type, path_flag, &offset); + if (KERNEL_NODES_FEATURE(VOLUME)) { + offset = svm_node_principled_volume<type>(kg, sd, stack, node, path_flag, offset); + } break; -# endif /* NODES_FEATURE(NODE_FEATURE_VOLUME) */ case NODE_MATH: - svm_node_math(kg, sd, stack, node.y, node.z, node.w, &offset); + svm_node_math(kg, sd, stack, node.y, node.z, node.w); break; case NODE_VECTOR_MATH: - svm_node_vector_math(kg, sd, stack, node.y, node.z, node.w, &offset); + offset = svm_node_vector_math(kg, sd, stack, node.y, node.z, node.w, offset); break; case NODE_RGB_RAMP: - svm_node_rgb_ramp(kg, sd, stack, node, &offset); + offset = svm_node_rgb_ramp(kg, sd, stack, node, offset); break; case NODE_GAMMA: svm_node_gamma(sd, stack, node.y, node.z, node.w); @@ -408,7 +425,7 @@ ccl_device_noinline void svm_eval_nodes( svm_node_brightness(sd, stack, node.y, node.z, node.w); break; case NODE_LIGHT_PATH: - svm_node_light_path(sd, state, stack, node.y, node.z, path_flag); + svm_node_light_path(INTEGRATOR_STATE_PASS, sd, stack, node.y, node.z, path_flag); break; case NODE_OBJECT_INFO: svm_node_object_info(kg, sd, stack, node.y, node.z); @@ -416,22 +433,22 @@ ccl_device_noinline void svm_eval_nodes( case NODE_PARTICLE_INFO: svm_node_particle_info(kg, sd, stack, node.y, node.z); break; -# if defined(__HAIR__) && NODES_FEATURE(NODE_FEATURE_HAIR) +#if defined(__HAIR__) case NODE_HAIR_INFO: - svm_node_hair_info(kg, sd, stack, node.y, node.z); + if (KERNEL_NODES_FEATURE(HAIR)) { + svm_node_hair_info(kg, sd, stack, node.y, node.z); + } break; -# endif /* NODES_FEATURE(NODE_FEATURE_HAIR) */ -#endif /* NODES_GROUP(NODE_GROUP_LEVEL_1) */ +#endif -#if NODES_GROUP(NODE_GROUP_LEVEL_2) case NODE_TEXTURE_MAPPING: - svm_node_texture_mapping(kg, sd, stack, node.y, node.z, &offset); + offset = svm_node_texture_mapping(kg, sd, stack, node.y, node.z, offset); break; case NODE_MAPPING: - svm_node_mapping(kg, sd, stack, node.y, node.z, node.w, &offset); + svm_node_mapping(kg, sd, stack, node.y, node.z, node.w); break; case NODE_MIN_MAX: - svm_node_min_max(kg, sd, stack, node.y, node.z, &offset); + offset = svm_node_min_max(kg, sd, stack, node.y, node.z, offset); break; case NODE_CAMERA: svm_node_camera(kg, sd, stack, node.y, node.z, node.w); @@ -440,47 +457,46 @@ ccl_device_noinline void svm_eval_nodes( svm_node_tex_environment(kg, sd, stack, node); break; case NODE_TEX_SKY: - svm_node_tex_sky(kg, sd, stack, node, &offset); + offset = svm_node_tex_sky(kg, sd, stack, node, offset); break; case NODE_TEX_GRADIENT: svm_node_tex_gradient(sd, stack, node); break; case NODE_TEX_VORONOI: - svm_node_tex_voronoi(kg, sd, stack, node.y, node.z, node.w, &offset); + offset = svm_node_tex_voronoi<node_feature_mask>( + kg, sd, stack, node.y, node.z, node.w, offset); break; case NODE_TEX_MUSGRAVE: - svm_node_tex_musgrave(kg, sd, stack, node.y, node.z, node.w, &offset); + offset = svm_node_tex_musgrave(kg, sd, stack, node.y, node.z, node.w, offset); break; case NODE_TEX_WAVE: - svm_node_tex_wave(kg, sd, stack, node, &offset); + offset = svm_node_tex_wave(kg, sd, stack, node, offset); break; case NODE_TEX_MAGIC: - svm_node_tex_magic(kg, sd, stack, node, &offset); + offset = svm_node_tex_magic(kg, sd, stack, node, offset); break; case NODE_TEX_CHECKER: svm_node_tex_checker(kg, sd, stack, node); break; case NODE_TEX_BRICK: - svm_node_tex_brick(kg, sd, stack, node, &offset); + offset = svm_node_tex_brick(kg, sd, stack, node, offset); break; case NODE_TEX_WHITE_NOISE: - svm_node_tex_white_noise(kg, sd, stack, node.y, node.z, node.w, &offset); + svm_node_tex_white_noise(kg, sd, stack, node.y, node.z, node.w); break; case NODE_NORMAL: - svm_node_normal(kg, sd, stack, node.y, node.z, node.w, &offset); + offset = svm_node_normal(kg, sd, stack, node.y, node.z, node.w, offset); break; case NODE_LIGHT_FALLOFF: svm_node_light_falloff(sd, stack, node); break; case NODE_IES: - svm_node_ies(kg, sd, stack, node, &offset); + svm_node_ies(kg, sd, stack, node); break; -#endif /* NODES_GROUP(NODE_GROUP_LEVEL_2) */ -#if NODES_GROUP(NODE_GROUP_LEVEL_3) case NODE_RGB_CURVES: case NODE_VECTOR_CURVES: - svm_node_curves(kg, sd, stack, node, &offset); + offset = svm_node_curves(kg, sd, stack, node, offset); break; case NODE_TANGENT: svm_node_tangent(kg, sd, stack, node); @@ -492,7 +508,7 @@ ccl_device_noinline void svm_eval_nodes( svm_node_invert(sd, stack, node.y, node.z, node.w); break; case NODE_MIX: - svm_node_mix(kg, sd, stack, node.y, node.z, node.w, &offset); + offset = svm_node_mix(kg, sd, stack, node.y, node.z, node.w, offset); break; case NODE_SEPARATE_VECTOR: svm_node_separate_vector(sd, stack, node.y, node.z, node.w); @@ -501,10 +517,10 @@ ccl_device_noinline void svm_eval_nodes( svm_node_combine_vector(sd, stack, node.y, node.z, node.w); break; case NODE_SEPARATE_HSV: - svm_node_separate_hsv(kg, sd, stack, node.y, node.z, node.w, &offset); + offset = svm_node_separate_hsv(kg, sd, stack, node.y, node.z, node.w, offset); break; case NODE_COMBINE_HSV: - svm_node_combine_hsv(kg, sd, stack, node.y, node.z, node.w, &offset); + offset = svm_node_combine_hsv(kg, sd, stack, node.y, node.z, node.w, offset); break; case NODE_VECTOR_ROTATE: svm_node_vector_rotate(sd, stack, node.y, node.z, node.w); @@ -522,39 +538,36 @@ ccl_device_noinline void svm_eval_nodes( svm_node_blackbody(kg, sd, stack, node.y, node.z); break; case NODE_MAP_RANGE: - svm_node_map_range(kg, sd, stack, node.y, node.z, node.w, &offset); + offset = svm_node_map_range(kg, sd, stack, node.y, node.z, node.w, offset); break; case NODE_CLAMP: - svm_node_clamp(kg, sd, stack, node.y, node.z, node.w, &offset); + offset = svm_node_clamp(kg, sd, stack, node.y, node.z, node.w, offset); break; -# ifdef __SHADER_RAYTRACE__ +#ifdef __SHADER_RAYTRACE__ case NODE_BEVEL: - svm_node_bevel(kg, sd, state, stack, node); + svm_node_bevel<node_feature_mask>(INTEGRATOR_STATE_PASS, sd, stack, node); break; case NODE_AMBIENT_OCCLUSION: - svm_node_ao(kg, sd, state, stack, node); + svm_node_ao<node_feature_mask>(INTEGRATOR_STATE_PASS, sd, stack, node); break; -# endif /* __SHADER_RAYTRACE__ */ -#endif /* NODES_GROUP(NODE_GROUP_LEVEL_3) */ +#endif -#if NODES_GROUP(NODE_GROUP_LEVEL_4) -# if NODES_FEATURE(NODE_FEATURE_VOLUME) case NODE_TEX_VOXEL: - svm_node_tex_voxel(kg, sd, stack, node, &offset); + if (KERNEL_NODES_FEATURE(VOLUME)) { + offset = svm_node_tex_voxel(kg, sd, stack, node, offset); + } break; -# endif /* NODES_FEATURE(NODE_FEATURE_VOLUME) */ case NODE_AOV_START: - if (!svm_node_aov_check(state, buffer)) { + if (!svm_node_aov_check(path_flag, render_buffer)) { return; } break; case NODE_AOV_COLOR: - svm_node_aov_color(kg, sd, stack, node, buffer); + svm_node_aov_color(INTEGRATOR_STATE_PASS, sd, stack, node, render_buffer); break; case NODE_AOV_VALUE: - svm_node_aov_value(kg, sd, stack, node, buffer); + svm_node_aov_value(INTEGRATOR_STATE_PASS, sd, stack, node, render_buffer); break; -#endif /* NODES_GROUP(NODE_GROUP_LEVEL_4) */ default: kernel_assert(!"Unknown node type was passed to the SVM machine"); return; |