[Bf-blender-cvs] [f69f542] cycles_kernel_split: De-duplicate shadow_blocked function
varunsundar08
noreply at git.blender.org
Tue May 5 09:17:15 CEST 2015
Commit: f69f5422f9049aea45d45470a9a5e9b348a143d5
Author: varunsundar08
Date: Tue May 5 00:55:59 2015 +0530
Branches: cycles_kernel_split
https://developer.blender.org/rBf69f5422f9049aea45d45470a9a5e9b348a143d5
De-duplicate shadow_blocked function
===================================================================
M intern/cycles/kernel/kernel_ShadowBlocked.cl
M intern/cycles/kernel/kernel_path.h
M intern/cycles/kernel/kernel_path_surface.h
M intern/cycles/kernel/kernel_path_volume.h
M intern/cycles/kernel/kernel_shadow.h
===================================================================
diff --git a/intern/cycles/kernel/kernel_ShadowBlocked.cl b/intern/cycles/kernel/kernel_ShadowBlocked.cl
index 24a204d..6f380ae 100644
--- a/intern/cycles/kernel/kernel_ShadowBlocked.cl
+++ b/intern/cycles/kernel/kernel_ShadowBlocked.cl
@@ -117,7 +117,7 @@ __kernel void kernel_ocl_path_trace_ShadowBlocked_DirectLighting_SPLIT_KERNEL(
ccl_global Intersection *isect_global = RAY_SHADOW_RAY_CAST_AO ? isect_ao_global : isect_dl_global;
float3 shadow;
- update_path_radiance = !(shadow_blocked_SPLIT_KERNEL(kg, state, light_ray_global, &shadow, sd_shadow, isect_global));
+ update_path_radiance = !(shadow_blocked(kg, state, light_ray_global, &shadow, sd_shadow, isect_global));
/* We use light_ray_global's P and t to store shadow and update_path_radiance */
light_ray_global->P = shadow;
diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h
index 7b9a51c..ead9ea7 100644
--- a/intern/cycles/kernel/kernel_path.h
+++ b/intern/cycles/kernel/kernel_path.h
@@ -250,7 +250,7 @@ ccl_device void kernel_path_indirect(ccl_addr_space KernelGlobals *kg, RNG *rng,
light_ray.dP = sd.dP;
light_ray.dD = differential3_zero();
- if(!shadow_blocked(kg, &state, &light_ray, &ao_shadow))
+ if(!shadow_blocked(kg, &state, &light_ray, &ao_shadow, NULL, NULL))
path_radiance_accum_ao(L, throughput, ao_alpha, ao_bsdf, ao_shadow, state.bounce);
}
}
@@ -320,7 +320,7 @@ ccl_device void kernel_path_ao(ccl_addr_space KernelGlobals *kg, ShaderData *sd,
light_ray.dP = sd->dP;
light_ray.dD = differential3_zero();
- if(!shadow_blocked(kg, state, &light_ray, &ao_shadow))
+ if(!shadow_blocked(kg, state, &light_ray, &ao_shadow, NULL, NULL))
path_radiance_accum_ao(L, throughput, ao_alpha, ao_bsdf, ao_shadow, state->bounce);
}
}
@@ -356,7 +356,7 @@ ccl_device void kernel_branched_path_ao(ccl_addr_space KernelGlobals *kg, Shader
light_ray.dP = sd->dP;
light_ray.dD = differential3_zero();
- if(!shadow_blocked(kg, state, &light_ray, &ao_shadow))
+ if(!shadow_blocked(kg, state, &light_ray, &ao_shadow, NULL, NULL))
path_radiance_accum_ao(L, throughput*num_samples_inv, ao_alpha, ao_bsdf, ao_shadow, state->bounce);
}
}
diff --git a/intern/cycles/kernel/kernel_path_surface.h b/intern/cycles/kernel/kernel_path_surface.h
index 2eea539..4bd430f 100644
--- a/intern/cycles/kernel/kernel_path_surface.h
+++ b/intern/cycles/kernel/kernel_path_surface.h
@@ -59,7 +59,7 @@ ccl_device void kernel_branched_path_surface_connect_light(ccl_addr_space Kernel
/* trace shadow ray */
float3 shadow;
- if(!shadow_blocked(kg, state, &light_ray, &shadow)) {
+ if(!shadow_blocked(kg, state, &light_ray, &shadow, NULL, NULL)) {
/* accumulate */
path_radiance_accum_light(L, throughput*num_samples_inv, &L_light, shadow, num_samples_inv, state->bounce, is_lamp);
}
@@ -91,7 +91,7 @@ ccl_device void kernel_branched_path_surface_connect_light(ccl_addr_space Kernel
/* trace shadow ray */
float3 shadow;
- if(!shadow_blocked(kg, state, &light_ray, &shadow)) {
+ if(!shadow_blocked(kg, state, &light_ray, &shadow, NULL, NULL)) {
/* accumulate */
path_radiance_accum_light(L, throughput*num_samples_inv, &L_light, shadow, num_samples_inv, state->bounce, is_lamp);
}
@@ -113,7 +113,7 @@ ccl_device void kernel_branched_path_surface_connect_light(ccl_addr_space Kernel
/* trace shadow ray */
float3 shadow;
- if(!shadow_blocked(kg, state, &light_ray, &shadow)) {
+ if(!shadow_blocked(kg, state, &light_ray, &shadow, NULL, NULL)) {
/* accumulate */
path_radiance_accum_light(L, throughput*num_samples_adjust, &L_light, shadow, num_samples_adjust, state->bounce, is_lamp);
}
@@ -210,7 +210,7 @@ ccl_device_inline void kernel_path_surface_connect_light(ccl_addr_space KernelGl
/* trace shadow ray */
float3 shadow;
- if(!shadow_blocked(kg, state, &light_ray, &shadow)) {
+ if(!shadow_blocked(kg, state, &light_ray, &shadow, NULL, NULL)) {
/* accumulate */
path_radiance_accum_light(L, throughput, &L_light, shadow, 1.0f, state->bounce, is_lamp);
}
diff --git a/intern/cycles/kernel/kernel_path_volume.h b/intern/cycles/kernel/kernel_path_volume.h
index 064645a..773f075 100644
--- a/intern/cycles/kernel/kernel_path_volume.h
+++ b/intern/cycles/kernel/kernel_path_volume.h
@@ -48,7 +48,7 @@ ccl_device void kernel_path_volume_connect_light(ccl_addr_space KernelGlobals *k
/* trace shadow ray */
float3 shadow;
- if(!shadow_blocked(kg, state, &light_ray, &shadow)) {
+ if(!shadow_blocked(kg, state, &light_ray, &shadow, NULL, NULL)) {
/* accumulate */
path_radiance_accum_light(L, throughput, &L_light, shadow, 1.0f, state->bounce, is_lamp);
}
@@ -164,7 +164,7 @@ ccl_device void kernel_branched_path_volume_connect_light(ccl_addr_space KernelG
/* trace shadow ray */
float3 shadow;
- if(!shadow_blocked(kg, state, &light_ray, &shadow)) {
+ if(!shadow_blocked(kg, state, &light_ray, &shadow, NULL, NULL)) {
/* accumulate */
path_radiance_accum_light(L, tp*num_samples_inv, &L_light, shadow, num_samples_inv, state->bounce, is_lamp);
}
@@ -215,7 +215,7 @@ ccl_device void kernel_branched_path_volume_connect_light(ccl_addr_space KernelG
/* trace shadow ray */
float3 shadow;
- if(!shadow_blocked(kg, state, &light_ray, &shadow)) {
+ if(!shadow_blocked(kg, state, &light_ray, &shadow, NULL, NULL)) {
/* accumulate */
path_radiance_accum_light(L, tp*num_samples_inv, &L_light, shadow, num_samples_inv, state->bounce, is_lamp);
}
@@ -255,7 +255,7 @@ ccl_device void kernel_branched_path_volume_connect_light(ccl_addr_space KernelG
/* trace shadow ray */
float3 shadow;
- if(!shadow_blocked(kg, state, &light_ray, &shadow)) {
+ if(!shadow_blocked(kg, state, &light_ray, &shadow, NULL, NULL)) {
/* accumulate */
path_radiance_accum_light(L, tp, &L_light, shadow, 1.0f, state->bounce, is_lamp);
}
diff --git a/intern/cycles/kernel/kernel_shadow.h b/intern/cycles/kernel/kernel_shadow.h
index 26bafeb..37f3b12 100644
--- a/intern/cycles/kernel/kernel_shadow.h
+++ b/intern/cycles/kernel/kernel_shadow.h
@@ -41,7 +41,8 @@ CCL_NAMESPACE_BEGIN
#define STACK_MAX_HITS 64
-ccl_device_inline bool shadow_blocked(ccl_addr_space KernelGlobals *kg, PathState *state, Ray *ray, float3 *shadow)
+ /* dummy_arg_1 and dummy_arg_2 are declared just to match the function signature of the GPU variant */
+ccl_device_inline bool shadow_blocked(ccl_addr_space KernelGlobals *kg, PathState *state, Ray *ray, float3 *shadow, void *dummy_arg_1, void *dummy_arg_2)
{
*shadow = make_float3(1.0f, 1.0f, 1.0f);
@@ -180,8 +181,8 @@ ccl_device_inline bool shadow_blocked(ccl_addr_space KernelGlobals *kg, PathStat
* potentially transparent, and only in that case start marching. this gives
* one extra ray cast for the cases were we do want transparency. */
-#ifdef __SPLIT_KERNEL__
-ccl_device_inline bool shadow_blocked_SPLIT_KERNEL(ccl_addr_space KernelGlobals *kg, ccl_addr_space PathState *state, ccl_addr_space Ray *ray_input, float3 *shadow, ccl_addr_space ShaderData *sd_input, ccl_addr_space Intersection *isect_input)
+/* The arguments sd_mem and isect_mem are meaningful only for OpenCL split kernel. Other uses can just pass a NULL */
+ccl_device_inline bool shadow_blocked(ccl_addr_space KernelGlobals *kg, ccl_addr_space PathState *state, ccl_addr_space Ray *ray_input, float3 *shadow, ccl_addr_space ShaderData *sd_mem, ccl_addr_space Intersection *isect_mem)
{
*shadow = make_float3(1.0f, 1.0f, 1.0f);
@@ -196,7 +197,7 @@ ccl_device_inline bool shadow_blocked_SPLIT_KERNEL(ccl_addr_space KernelGlobals
#endif
#ifdef __SPLIT_KERNEL__
- ccl_addr_space Intersection *isect = isect_global;
+ ccl_addr_space Intersection *isect = isect_mem;
#else
Intersection isect_object;
Intersection *isect = &isect_object;
@@ -245,7 +246,7 @@ ccl_device_inline bool shadow_blocked_SPLIT_KERNEL(ccl_addr_space KernelGlobals
/* setup shader data at surface */
#ifdef __SPLIT_KERNEL__
- ccl_addr_space ShaderData *sd = sd_input;
+ ccl_addr_space ShaderData *sd = sd_mem;
#else
ShaderData sd_object;
ShaderData *sd = &sd_object;
@@ -285,97 +286,6 @@ ccl_device_inline bool shadow_blocked_SPLIT_KERNEL(ccl_addr_space KernelGlobals
return blocked;
}
-#else
-
-ccl_device_inline bool shadow_blocked(ccl_addr_space KernelGlobals *kg, PathState *state, Ray *ray, float3 *shadow)
-{
- *shadow = make_float3(1.0f, 1.0f, 1.0f);
-
- if(ray->t == 0.0f)
- return false;
-
- Intersection isect;
- bool blocked = scene_intersect(kg, ray, PATH_RAY_SHADOW_OPAQUE, &isect, NULL, 0.0f, 0.0f);
-
-#ifdef __TRANSPARENT_SHADOWS__
- if(blocked && kernel_data.integrator.transparent_shadows) {
- if(shader_transparent_shadow(kg, &isect)) {
- float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
- float3 Pend = ray->P + ray->D*ray->t;
- int bounce = state->transparent_bounce;
-#ifdef __VOLUME__
- PathState ps = *state;
-#endif
-
- for(;;) {
- if(bounce >= kernel_data.integrator.transparent_max_bounce)
- return true;
-
- if(!scene_intersect(kg, ray, PATH_RAY_SHADOW_TRANSPARENT, &isect, NULL, 0.0f, 0.0f))
- {
-
-#ifdef __VOLUME__
- /* attenuation for last line segment towards light */
- if(ps.volume_stack[0].shader != SHADER_NONE)
- kernel_volume_shadow(kg, &ps, ray, &throughput);
-#endif
-
- *shadow *= throughput;
-
- return false;
- }
-
- if(!shader_transparent_shadow(kg, &isect))
- return true;
-
-#ifdef __VOLUME__
- /* attenuation between last surface and next surface */
- if(ps.volume_stack[0].shader != SHADER_NONE) {
- Ray segment_ray = *ray;
- segment_ray.t = isect.t;
- kernel_volume_shadow(kg, &ps, &segment_ray, &throughput);
- }
-#endif
-
- /* setup shader data at surface */
- ShaderData sd;
- shader_setup_from_ray(kg, &sd, &isect, ray, state->bounce+1, bounce);
-
- /* attenuation from transparent surface */
- if(!(sd.flag & SD_HAS_ONLY_VOLUME)) {
- shader_eval_surface(kg, &sd, 0.0f, PATH_RAY_SHADOW, SHADER_CONTEXT_SHADOW);
- throughput *= shader_bsdf_transpa
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list