Commit: 9aa5aeee46aeb45f5ea191fd29998daca79da343 Author: Brecht Van Lommel Date: Fri Jul 9 16:48:40 2021 +0200 Branches: cycles-x https://developer.blender.org/rB9aa5aeee46aeb45f5ea191fd29998daca79da343
Cycles X: restore AO pass support This uses shader ray-tracing, which is not ideal for performance or runtime compile times. For that reason it might be moved into its own kernel later, but this at least brings back the feature. Differential Revision: https://developer.blender.org/D11873 =================================================================== M intern/cycles/device/optix/device_impl.cpp M intern/cycles/device/optix/device_impl.h M intern/cycles/kernel/CMakeLists.txt M intern/cycles/kernel/integrator/integrator_init_from_bake.h M intern/cycles/kernel/integrator/integrator_intersect_closest.h M intern/cycles/kernel/integrator/integrator_shade_surface.h M intern/cycles/kernel/integrator/integrator_subsurface.h M intern/cycles/kernel/kernel_passes.h D intern/cycles/kernel/kernel_path.h M intern/cycles/kernel/kernel_shader.h M intern/cycles/render/scene.cpp =================================================================== diff --git a/intern/cycles/device/optix/device_impl.cpp b/intern/cycles/device/optix/device_impl.cpp index e265bbba726..4907be6570f 100644 --- a/intern/cycles/device/optix/device_impl.cpp +++ b/intern/cycles/device/optix/device_impl.cpp @@ -393,6 +393,9 @@ bool OptiXDevice::load_kernels(const DeviceRequestedFeatures &requested_features group_descs[PG_CALL_SVM_BEVEL].callables.moduleDC = optix_module; group_descs[PG_CALL_SVM_BEVEL].callables.entryFunctionNameDC = "__direct_callable__svm_node_bevel"; + group_descs[PG_CALL_AO_PASS].kind = OPTIX_PROGRAM_GROUP_KIND_CALLABLES; + group_descs[PG_CALL_AO_PASS].callables.moduleDC = optix_module; + group_descs[PG_CALL_AO_PASS].callables.entryFunctionNameDC = "__direct_callable__ao_pass"; } optix_assert(optixProgramGroupCreate( diff --git a/intern/cycles/device/optix/device_impl.h b/intern/cycles/device/optix/device_impl.h index 12251736ed0..85e7b0aa6a8 100644 --- a/intern/cycles/device/optix/device_impl.h +++ b/intern/cycles/device/optix/device_impl.h @@ -46,6 +46,7 @@ enum { # endif PG_CALL_SVM_AO, PG_CALL_SVM_BEVEL, + PG_CALL_AO_PASS, NUM_PROGRAM_GROUPS }; @@ -58,7 +59,7 @@ static const int NUM_HIT_PROGRAM_GROUPS = 5; static const int NUM_HIT_PROGRAM_GROUPS = 3; # endif static const int CALLABLE_PROGRAM_GROUPS_BASE = PG_CALL_SVM_AO; -static const int NUM_CALLABLE_PROGRAM_GROUPS = 2; +static const int NUM_CALLABLE_PROGRAM_GROUPS = 3; /* List of OptiX pipelines. */ enum { PIP_SHADE_RAYTRACE, PIP_INTERSECT, NUM_PIPELINES }; diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index a1dc2d8b902..44175fec111 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -81,7 +81,6 @@ set(SRC_HEADERS kernel_math.h kernel_montecarlo.h kernel_passes.h - kernel_path.h kernel_path_state.h kernel_profiling.h kernel_projection.h diff --git a/intern/cycles/kernel/integrator/integrator_init_from_bake.h b/intern/cycles/kernel/integrator/integrator_init_from_bake.h index 2048f75c617..98ba0708e60 100644 --- a/intern/cycles/kernel/integrator/integrator_init_from_bake.h +++ b/intern/cycles/kernel/integrator/integrator_init_from_bake.h @@ -164,8 +164,8 @@ ccl_device bool integrator_init_from_bake(INTEGRATOR_STATE_ARGS, integrator_state_write_isect(INTEGRATOR_STATE_PASS, &isect); /* Setup next kernel to execute. */ - const int flags = kernel_tex_fetch(__shaders, shader).flags; - if (flags & SD_HAS_RAYTRACE) { + const int shader_flags = kernel_tex_fetch(__shaders, shader).flags; + if ((shader_flags & SD_HAS_RAYTRACE) || (kernel_data.film.pass_ao != PASS_UNUSED)) { INTEGRATOR_PATH_INIT_SORTED(DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE, shader); } else { diff --git a/intern/cycles/kernel/integrator/integrator_intersect_closest.h b/intern/cycles/kernel/integrator/integrator_intersect_closest.h index d49068d9823..dad08de8590 100644 --- a/intern/cycles/kernel/integrator/integrator_intersect_closest.h +++ b/intern/cycles/kernel/integrator/integrator_intersect_closest.h @@ -91,7 +91,7 @@ ccl_device_forceinline void integrator_intersect_shader_next_kernel( const int shader_flags) { /* Setup next kernel to execute. */ - if (shader_flags & SD_HAS_RAYTRACE) { + if ((shader_flags & SD_HAS_RAYTRACE) || (kernel_data.film.pass_ao != PASS_UNUSED)) { INTEGRATOR_PATH_NEXT_SORTED( current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE, shader); } diff --git a/intern/cycles/kernel/integrator/integrator_shade_surface.h b/intern/cycles/kernel/integrator/integrator_shade_surface.h index 0557e3a2c58..3f006192501 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_surface.h +++ b/intern/cycles/kernel/integrator/integrator_shade_surface.h @@ -281,6 +281,50 @@ ccl_device_forceinline bool integrate_surface_volume_only_bounce(INTEGRATOR_STAT } #endif +#if defined(__AO__) && defined(__SHADER_RAYTRACE__) +ccl_device_forceinline void integrate_surface_ao_pass(INTEGRATOR_STATE_CONST_ARGS, + const ShaderData *ccl_restrict sd, + const RNGState *ccl_restrict rng_state, + ccl_global float *ccl_restrict render_buffer) +{ +# ifdef __KERNEL_OPTIX__ + optixDirectCall<void>(2, INTEGRATOR_STATE_PASS, sd, rng_state, render_buffer); +} + +extern "C" __device__ void __direct_callable__ao_pass(INTEGRATOR_STATE_CONST_ARGS, + const ShaderData *ccl_restrict sd, + const RNGState *ccl_restrict rng_state, + ccl_global float *ccl_restrict render_buffer) +{ +# endif /* __KERNEL_OPTIX__ */ + float bsdf_u, bsdf_v; + path_state_rng_2D(kg, rng_state, PRNG_BSDF_U, &bsdf_u, &bsdf_v); + + const float3 ao_N = shader_bsdf_ao_normal(kg, sd); + float3 ao_D; + float ao_pdf; + sample_cos_hemisphere(ao_N, bsdf_u, bsdf_v, &ao_D, &ao_pdf); + + if (dot(sd->Ng, ao_D) > 0.0f && ao_pdf != 0.0f) { + Ray ray ccl_optional_struct_init; + ray.P = ray_offset(sd->P, sd->Ng); + ray.D = ao_D; + ray.t = kernel_data.integrator.ao_bounces_distance; + ray.time = sd->time; + ray.dP = differential_zero_compact(); + ray.dD = differential_zero_compact(); + + Intersection isect ccl_optional_struct_init; + if (!scene_intersect(kg, &ray, PATH_RAY_SHADOW_OPAQUE, &isect)) { + ccl_global float *buffer = kernel_pass_pixel_render_buffer(INTEGRATOR_STATE_PASS, + render_buffer); + const float3 throughput = INTEGRATOR_STATE(path, throughput); + kernel_write_pass_float3(buffer + kernel_data.film.pass_ao, throughput); + } + } +} +#endif /* defined(__AO__) && defined(__SHADER_RAYTRACE__) */ + template<uint node_feature_mask> ccl_device bool integrate_surface(INTEGRATOR_STATE_ARGS, ccl_global float *ccl_restrict render_buffer) @@ -370,14 +414,14 @@ ccl_device bool integrate_surface(INTEGRATOR_STATE_ARGS, kernel_write_shadow_catcher_bounce_data(INTEGRATOR_STATE_PASS, &sd, render_buffer); #endif - /* TODO */ -#if 0 -# ifdef __AO__ - /* ambient occlusion */ - if (kernel_data.integrator.use_ambient_occlusion) { - kernel_path_ao(kg, &sd, emission_sd, L, state, throughput, shader_bsdf_alpha(kg, &sd)); - } -# endif /* __AO__ */ +#if defined(__AO__) && defined(__SHADER_RAYTRACE__) + /* Ambient occlusion pass. */ + if (node_feature_mask & NODE_FEATURE_RAYTRACE) { + if ((kernel_data.film.pass_ao != PASS_UNUSED) && + (INTEGRATOR_STATE(path, flag) & PATH_RAY_CAMERA)) { + integrate_surface_ao_pass(INTEGRATOR_STATE_PASS, &sd, &rng_state, render_buffer); + } + } #endif continue_path_label = integrate_surface_bsdf_bssrdf_bounce( @@ -387,14 +431,12 @@ ccl_device bool integrate_surface(INTEGRATOR_STATE_ARGS, else { continue_path_label = integrate_surface_volume_only_bounce(INTEGRATOR_STATE_PASS, &sd); } -#endif if (continue_path_label & LABEL_TRANSMIT) { /* Enter/Exit volume. */ -#ifdef __VOLUME__ volume_stack_enter_exit(INTEGRATOR_STATE_PASS, &sd); -#endif } +#endif return continue_path_label != 0; } diff --git a/intern/cycles/kernel/integrator/integrator_subsurface.h b/intern/cycles/kernel/integrator/integrator_subsurface.h index 3a282f023d8..a8ebbdd6a0c 100644 --- a/intern/cycles/kernel/integrator/integrator_subsurface.h +++ b/intern/cycles/kernel/integrator/integrator_subsurface.h @@ -531,8 +531,8 @@ ccl_device_inline bool subsurface_random_walk(INTEGRATOR_STATE_ARGS) INTEGRATOR_STATE_WRITE(path, throughput) = throughput; const int shader = intersection_get_shader(kg, &ss_isect.hits[0]); - const int flags = kernel_tex_fetch(__shaders, shader).flags; - if (flags & SD_HAS_RAYTRACE) { + const int shader_flags = kernel_tex_fetch(__shaders, shader).flags; + if ((shader_flags & SD_HAS_RAYTRACE) || (kernel_data.film.pass_ao != PASS_UNUSED)) { INTEGRATOR_PATH_NEXT_SORTED(DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE, shader); diff --git a/intern/cycles/kernel/kernel_passes.h b/intern/cycles/kernel/kernel_passes.h index da4d12d8891..3ac76774add 100644 --- a/intern/cycles/kernel/kernel_passes.h +++ b/intern/cycles/kernel/kernel_passes.h @@ -309,78 +309,4 @@ ccl_device_inline void kernel_write_data_passes(INTEGRATOR_STATE_ARGS, #endif } -#if 0 -ccl_device_inline void kernel_write_light_passes(const KernelGlobals *ccl_restrict kg, - ccl_global float *ccl_restrict buffer, - PathRadiance *L) -{ -# ifdef __PASSES__ - int light_flag = kernel_data.film.light_pass_flag; - - if (!kernel_data.film.use_light_pass) - return; - - if (light_flag & PASSMASK(AO)) - kernel_write_pass_float3(buffer + kernel_data.film.pass_ao, L->ao); -# endif -} -#endif - -#if 0 -ccl_device_inline void kernel_write_result(const KernelGlobals *ccl_restrict kg, - ccl_global float *ccl_restrict buffer, - int sample, - PathRadiance *L) -{ - PROFILING_INIT(kg, PROFILING_WRITE_RESULT); - PROFILING_OBJECT(PRIM_NONE); - - float alpha; - float3 L_sum = path_radiance_clamp_and_sum(kg, L, &alpha); - - if (kernel_data.film.light_pass_flag & PASSMASK(COMBINED)) { - kernel_write_pass_float4(buffer, make_float4(L_sum.x, L_sum.y, L_sum.z, alpha)); - } - - - /* Adaptive Sampling. Fill the additional buffer with the odd samples and calculate our stopping - criteria. This is the heuristic from "A hierarchical automatic stopping condition for Monte - Carlo global illumination" except that here it is applied per pixel and not in hierarchical - tiles. */ - if (kernel_data.film.pass_adaptive_aux_buffer != PASS_UNUSED) { - if @@ Diff output truncated at 10240 characters. @@ _______________________________________________ Bf-blender-cvs mailing list [email protected] List details, subscription details or unsubscribe: https://lists.blender.org/mailman/listinfo/bf-blender-cvs
