[Bf-blender-cvs] [9aa5aeee46a] cycles-x: Cycles X: restore AO pass support
Brecht Van Lommel
noreply at git.blender.org
Fri Jul 9 18:01:58 CEST 2021
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. @@
More information about the Bf-blender-cvs
mailing list