kannonier8/blender
1
0
Fork 0
forked from bartvdbraak/blender
Code Pull Requests Activity

Cycles: Refactor how we pass bounce info to light path node.

Browse Source 
This commit changes the way how we pass bounce information to the Light
Path node. Instead of manualy copying the bounces into ShaderData, we now
directly pass PathState. This reduces the arguments that we need to pass
around and also makes it easier to extend the feature.

This commit also exposes the Transmission Bounce Depth to the Light Path
node. It works similar to the Transparent Depth Output: Replace a
Transmission lightpath after X bounces with another shader, e.g a Diffuse
one. This can be used to avoid black surfaces, due to low amount of max
bounces.

Reviewed by Sergey and Brecht, thanks for some hlp with this.

I tested compilation and usage on CPU (SVM and OSL), CUDA, OpenCL Split
and Mega kernel. Hopefully this covers all devices. :)
This commit is contained in:
Thomas Dinges 2016-01-06 23:38:13 +01:00
parent be28706bac
commit 83e73a2100
29 changed files with 197 additions and 204 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

22
intern/cycles/device/device_opencl.cpp
View File

@ -1920,10 +1920,6 @@ public:
cl_mem time_sd_DL_shadow;
cl_mem ray_length_sd;
cl_mem ray_length_sd_DL_shadow;
cl_mem ray_depth_sd;
cl_mem ray_depth_sd_DL_shadow;
cl_mem transparent_depth_sd;
cl_mem transparent_depth_sd_DL_shadow;
/* Ray differentials. */
cl_mem dP_sd, dI_sd;
@ -2073,10 +2069,6 @@ public:
time_sd_DL_shadow = NULL;
ray_length_sd = NULL;
ray_length_sd_DL_shadow = NULL;
ray_depth_sd = NULL;
ray_depth_sd_DL_shadow = NULL;
transparent_depth_sd = NULL;
transparent_depth_sd_DL_shadow = NULL;
/* Ray differentials. */
dP_sd = NULL;
@ -2417,10 +2409,6 @@ public:
release_mem_object_safe(time_sd_DL_shadow);
release_mem_object_safe(ray_length_sd);
release_mem_object_safe(ray_length_sd_DL_shadow);
release_mem_object_safe(ray_depth_sd);
release_mem_object_safe(ray_depth_sd_DL_shadow);
release_mem_object_safe(transparent_depth_sd);
release_mem_object_safe(transparent_depth_sd_DL_shadow);
/* Ray differentials. */
release_mem_object_safe(dP_sd);
@ -2619,10 +2607,6 @@ public:
time_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(float));
ray_length_sd = mem_alloc(num_global_elements * sizeof(float));
ray_length_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(float));
ray_depth_sd = mem_alloc(num_global_elements * sizeof(int));
ray_depth_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(int));
transparent_depth_sd = mem_alloc(num_global_elements * sizeof(int));
transparent_depth_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(int));
/* Ray differentials. */
dP_sd = mem_alloc(num_global_elements * sizeof(differential3));
@ -2725,11 +2709,7 @@ public:
time_sd,
time_sd_DL_shadow,
ray_length_sd,
ray_length_sd_DL_shadow,
ray_depth_sd,
ray_depth_sd_DL_shadow,
transparent_depth_sd,
transparent_depth_sd_DL_shadow);
ray_length_sd_DL_shadow);
/* Ray differentials. */
start_arg_index +=

46
intern/cycles/kernel/kernel_bake.h
View File

@ -46,7 +46,7 @@ ccl_device void compute_light_pass(KernelGlobals *kg, ShaderData *sd, PathRadian
/* evaluate surface shader */
float rbsdf = path_state_rng_1D(kg, &rng, &state, PRNG_BSDF);
shader_eval_surface(kg, sd, rbsdf, state.flag, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, sd, &state, rbsdf, state.flag, SHADER_CONTEXT_MAIN);
/* TODO, disable the closures we won't need */
@ -212,6 +212,7 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
ShaderEvalType type, int i, int offset, int sample)
{
ShaderData sd;
PathState state = {0};
uint4 in = input[i * 2];
uint4 diff = input[i * 2 + 1];
@ -262,13 +263,11 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
float3 I = Ng;
float t = 0.0f;
float time = TIME_INVALID;
int bounce = 0;
int transparent_bounce = 0;
/* light passes */
PathRadiance L;
shader_setup_from_sample(kg, &sd, P, Ng, I, shader, object, prim, u, v, t, time, bounce, transparent_bounce);
shader_setup_from_sample(kg, &sd, P, Ng, I, shader, object, prim, u, v, t, time);
sd.I = sd.N;
/* update differentials */
@ -294,7 +293,7 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
case SHADER_EVAL_NORMAL:
{
if((sd.flag & SD_HAS_BUMP)) {
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
}
/* compression: normal = (2 * color) - 1 */
@ -308,33 +307,33 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
}
case SHADER_EVAL_DIFFUSE_COLOR:
{
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = shader_bsdf_diffuse(kg, &sd);
break;
}
case SHADER_EVAL_GLOSSY_COLOR:
{
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = shader_bsdf_glossy(kg, &sd);
break;
}
case SHADER_EVAL_TRANSMISSION_COLOR:
{
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = shader_bsdf_transmission(kg, &sd);
break;
}
case SHADER_EVAL_SUBSURFACE_COLOR:
{
#ifdef __SUBSURFACE__
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = shader_bsdf_subsurface(kg, &sd);
#endif
break;
}
case SHADER_EVAL_EMISSION:
{
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_EMISSION);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_EMISSION);
out = shader_emissive_eval(kg, &sd);
break;
}
@ -358,52 +357,52 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
}
case SHADER_EVAL_DIFFUSE_DIRECT:
{
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = safe_divide_color(L.direct_diffuse, shader_bsdf_diffuse(kg, &sd));
break;
}
case SHADER_EVAL_GLOSSY_DIRECT:
{
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = safe_divide_color(L.direct_glossy, shader_bsdf_glossy(kg, &sd));
break;
}
case SHADER_EVAL_TRANSMISSION_DIRECT:
{
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = safe_divide_color(L.direct_transmission, shader_bsdf_transmission(kg, &sd));
break;
}
case SHADER_EVAL_SUBSURFACE_DIRECT:
{
#ifdef __SUBSURFACE__
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = safe_divide_color(L.direct_subsurface, shader_bsdf_subsurface(kg, &sd));
#endif
break;
}
case SHADER_EVAL_DIFFUSE_INDIRECT:
{
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = safe_divide_color(L.indirect_diffuse, shader_bsdf_diffuse(kg, &sd));
break;
}
case SHADER_EVAL_GLOSSY_INDIRECT:
{
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = safe_divide_color(L.indirect_glossy, shader_bsdf_glossy(kg, &sd));
break;
}
case SHADER_EVAL_TRANSMISSION_INDIRECT:
{
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = safe_divide_color(L.indirect_transmission, shader_bsdf_transmission(kg, &sd));
break;
}
case SHADER_EVAL_SUBSURFACE_INDIRECT:
{
#ifdef __SUBSURFACE__
shader_eval_surface(kg, &sd, 0.f, 0, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, 0.f, 0, SHADER_CONTEXT_MAIN);
out = safe_divide_color(L.indirect_subsurface, shader_bsdf_subsurface(kg, &sd));
#endif
break;
@ -429,11 +428,11 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
#endif
/* setup shader data */
shader_setup_from_background(kg, &sd, &ray, 0, 0);
shader_setup_from_background(kg, &sd, &ray);
/* evaluate */
int flag = 0; /* we can't know which type of BSDF this is for */
out = shader_eval_background(kg, &sd, flag, SHADER_CONTEXT_MAIN);
out = shader_eval_background(kg, &sd, &state, flag, SHADER_CONTEXT_MAIN);
break;
}
default:
@ -462,6 +461,7 @@ ccl_device void kernel_shader_evaluate(KernelGlobals *kg,
int sample)
{
ShaderData sd;
PathState state = {0};
uint4 in = input[i];
float3 out;
@ -476,7 +476,7 @@ ccl_device void kernel_shader_evaluate(KernelGlobals *kg,
/* evaluate */
float3 P = sd.P;
shader_eval_displacement(kg, &sd, SHADER_CONTEXT_MAIN);
shader_eval_displacement(kg, &sd, &state, SHADER_CONTEXT_MAIN);
out = sd.P - P;
}
else { // SHADER_EVAL_BACKGROUND
@ -498,11 +498,11 @@ ccl_device void kernel_shader_evaluate(KernelGlobals *kg,
#endif
/* setup shader data */
shader_setup_from_background(kg, &sd, &ray, 0, 0);
shader_setup_from_background(kg, &sd, &ray);
/* evaluate */
int flag = 0; /* we can't know which type of BSDF this is for */
out = shader_eval_background(kg, &sd, flag, SHADER_CONTEXT_MAIN);
out = shader_eval_background(kg, &sd, &state, flag, SHADER_CONTEXT_MAIN);
}
/* write output */

41
intern/cycles/kernel/kernel_emission.h
View File

@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN
/* Direction Emission */
ccl_device_noinline float3 direct_emissive_eval(KernelGlobals *kg,
LightSample *ls, float3 I, differential3 dI, float t, float time, int bounce, int transparent_bounce
LightSample *ls, ccl_addr_space PathState *state, float3 I, differential3 dI, float t, float time
#ifdef __SPLIT_KERNEL__
,ShaderData *sd_input
#endif
@ -45,19 +45,24 @@ ccl_device_noinline float3 direct_emissive_eval(KernelGlobals *kg,
ray.dP = differential3_zero();
ray.dD = dI;
shader_setup_from_background(kg, sd, &ray, bounce+1, transparent_bounce);
eval = shader_eval_background(kg, sd, 0, SHADER_CONTEXT_EMISSION);
shader_setup_from_background(kg, sd, &ray);
path_state_modify_bounce(state, true);
eval = shader_eval_background(kg, sd, state, 0, SHADER_CONTEXT_EMISSION);
path_state_modify_bounce(state, false);
}
else
#endif
{
shader_setup_from_sample(kg, sd, ls->P, ls->Ng, I, ls->shader, ls->object, ls->prim, ls->u, ls->v, t, time, bounce+1, transparent_bounce);
shader_setup_from_sample(kg, sd, ls->P, ls->Ng, I, ls->shader, ls->object, ls->prim, ls->u, ls->v, t, time);
ls->Ng = ccl_fetch(sd, Ng);
/* no path flag, we're evaluating this for all closures. that's weak but
* we'd have to do multiple evaluations otherwise */
shader_eval_surface(kg, sd, 0.0f, 0, SHADER_CONTEXT_EMISSION);
path_state_modify_bounce(state, true);
shader_eval_surface(kg, sd, state, 0.0f, 0, SHADER_CONTEXT_EMISSION);
path_state_modify_bounce(state, false);
/* evaluate emissive closure */
if(ccl_fetch(sd, flag) & SD_EMISSION)
@ -72,8 +77,7 @@ ccl_device_noinline float3 direct_emissive_eval(KernelGlobals *kg,
}
ccl_device_noinline bool direct_emission(KernelGlobals *kg, ShaderData *sd,
LightSample *ls, Ray *ray, BsdfEval *eval, bool *is_lamp,
int bounce, int transparent_bounce
LightSample *ls, ccl_addr_space PathState *state, Ray *ray, BsdfEval *eval, bool *is_lamp
#ifdef __SPLIT_KERNEL__
, ShaderData *sd_DL
#endif
@ -87,9 +91,7 @@ ccl_device_noinline bool direct_emission(KernelGlobals *kg, ShaderData *sd,
/* evaluate closure */
float3 light_eval = direct_emissive_eval(kg, ls, -ls->D, dD, ls->t, ccl_fetch(sd, time),
bounce,
transparent_bounce
float3 light_eval = direct_emissive_eval(kg, ls, state, -ls->D, dD, ls->t, ccl_fetch(sd, time)
#ifdef __SPLIT_KERNEL__
,sd_DL
#endif
@ -191,7 +193,7 @@ ccl_device_noinline float3 indirect_primitive_emission(KernelGlobals *kg, Shader
/* Indirect Lamp Emission */
ccl_device_noinline bool indirect_lamp_emission(KernelGlobals *kg, PathState *state, Ray *ray, float3 *emission
ccl_device_noinline bool indirect_lamp_emission(KernelGlobals *kg, ccl_addr_space PathState *state, Ray *ray, float3 *emission
#ifdef __SPLIT_KERNEL__
,ShaderData *sd
#endif
@ -219,9 +221,7 @@ ccl_device_noinline bool indirect_lamp_emission(KernelGlobals *kg, PathState *st
}
#endif
float3 L = direct_emissive_eval(kg, &ls, -ray->D, ray->dD, ls.t, ray->time,
state->bounce,
state->transparent_bounce
float3 L = direct_emissive_eval(kg, &ls, state, -ray->D, ray->dD, ls.t, ray->time
#ifdef __SPLIT_KERNEL__
,sd
#endif
@ -277,13 +277,18 @@ ccl_device_noinline float3 indirect_background(KernelGlobals *kg, ccl_addr_space
#ifdef __SPLIT_KERNEL__
/* evaluate background closure */
Ray priv_ray = *ray;
shader_setup_from_background(kg, sd_global, &priv_ray, state->bounce+1, state->transparent_bounce);
float3 L = shader_eval_background(kg, sd_global, state->flag, SHADER_CONTEXT_EMISSION);
shader_setup_from_background(kg, sd_global, &priv_ray);
path_state_modify_bounce(state, true);
float3 L = shader_eval_background(kg, sd_global, state, state->flag, SHADER_CONTEXT_EMISSION);
path_state_modify_bounce(state, false);
#else
ShaderData sd;
shader_setup_from_background(kg, &sd, ray, state->bounce+1, state->transparent_bounce);
shader_setup_from_background(kg, &sd, ray);
float3 L = shader_eval_background(kg, &sd, state->flag, SHADER_CONTEXT_EMISSION);
path_state_modify_bounce(state, true);
float3 L = shader_eval_background(kg, &sd, state, state->flag, SHADER_CONTEXT_EMISSION);
path_state_modify_bounce(state, false);
#endif
#ifdef __BACKGROUND_MIS__

18
intern/cycles/kernel/kernel_path.h
View File

@ -119,9 +119,7 @@ ccl_device void kernel_path_indirect(KernelGlobals *kg,
shader_setup_from_volume(kg,
&volume_sd,
&volume_ray,
state->bounce,
state->transparent_bounce);
&volume_ray);
kernel_volume_decoupled_record(kg,
state,
&volume_ray,
@ -252,11 +250,9 @@ ccl_device void kernel_path_indirect(KernelGlobals *kg,
shader_setup_from_ray(kg,
&sd,
&isect,
ray,
state->bounce,
state->transparent_bounce);
ray);
float rbsdf = path_state_rng_1D_for_decision(kg, rng, state, PRNG_BSDF);
shader_eval_surface(kg, &sd, rbsdf, state->flag, SHADER_CONTEXT_INDIRECT);
shader_eval_surface(kg, &sd, state, rbsdf, state->flag, SHADER_CONTEXT_INDIRECT);
#ifdef __BRANCHED_PATH__
shader_merge_closures(&sd);
#endif
@ -366,6 +362,7 @@ ccl_device void kernel_path_indirect(KernelGlobals *kg,
&bssrdf_u, &bssrdf_v);
subsurface_scatter_step(kg,
&sd,
state,
state->flag,
sc,
&lcg_state,
@ -481,6 +478,7 @@ ccl_device bool kernel_path_subsurface_scatter(
&ss_isect,
hit,
sd,
state,
state->flag,
sc,
false);
@ -684,7 +682,7 @@ ccl_device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample,
VolumeSegment volume_segment;
ShaderData volume_sd;
shader_setup_from_volume(kg, &volume_sd, &volume_ray, state.bounce, state.transparent_bounce);
shader_setup_from_volume(kg, &volume_sd, &volume_ray);
kernel_volume_decoupled_record(kg, &state,
&volume_ray, &volume_sd, &volume_segment, heterogeneous);
@ -774,9 +772,9 @@ ccl_device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample,
/* setup shading */
ShaderData sd;
shader_setup_from_ray(kg, &sd, &isect, &ray, state.bounce, state.transparent_bounce);
shader_setup_from_ray(kg, &sd, &isect, &ray);
float rbsdf = path_state_rng_1D_for_decision(kg, rng, &state, PRNG_BSDF);
shader_eval_surface(kg, &sd, rbsdf, state.flag, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, &sd, &state, rbsdf, state.flag, SHADER_CONTEXT_MAIN);
/* holdout */
#ifdef __HOLDOUT__

7
intern/cycles/kernel/kernel_path_branched.h
View File

@ -168,6 +168,7 @@ ccl_device void kernel_branched_path_subsurface_scatter(KernelGlobals *kg,
&ss_isect,
hit,
&bssrdf_sd,
state,
state->flag,
sc,
true);
@ -288,7 +289,7 @@ ccl_device float4 kernel_branched_path_integrate(KernelGlobals *kg, RNG *rng, in
VolumeSegment volume_segment;
ShaderData volume_sd;
shader_setup_from_volume(kg, &volume_sd, &volume_ray, state.bounce, state.transparent_bounce);
shader_setup_from_volume(kg, &volume_sd, &volume_ray);
kernel_volume_decoupled_record(kg, &state,
&volume_ray, &volume_sd, &volume_segment, heterogeneous);
@ -440,8 +441,8 @@ ccl_device float4 kernel_branched_path_integrate(KernelGlobals *kg, RNG *rng, in
/* setup shading */
ShaderData sd;
shader_setup_from_ray(kg, &sd, &isect, &ray, state.bounce, state.transparent_bounce);
shader_eval_surface(kg, &sd, 0.0f, state.flag, SHADER_CONTEXT_MAIN);
shader_setup_from_ray(kg, &sd, &isect, &ray);
shader_eval_surface(kg, &sd, &state, 0.0f, state.flag, SHADER_CONTEXT_MAIN);
shader_merge_closures(&sd);
/* holdout */

10
intern/cycles/kernel/kernel_path_state.h
View File

@ -168,5 +168,15 @@ ccl_device_inline float path_state_terminate_probability(KernelGlobals *kg, ccl_
return average(throughput); /* todo: try using max here */
}
/* TODO(DingTo): Find more meaningful name for this */
ccl_device_inline void path_state_modify_bounce(ccl_addr_space PathState *state, bool increase)
{
/* Modify bounce temporarily for shader eval */
if(increase)
state->bounce += 1;
else
state->bounce -= 1;
}
CCL_NAMESPACE_END

8
intern/cycles/kernel/kernel_path_surface.h
View File

@ -55,7 +55,7 @@ ccl_device void kernel_branched_path_surface_connect_light(KernelGlobals *kg, RN
LightSample ls;
lamp_light_sample(kg, i, light_u, light_v, ccl_fetch(sd, P), &ls);
if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp, state->bounce, state->transparent_bounce)) {
if(direct_emission(kg, sd, &ls, state, &light_ray, &L_light, &is_lamp)) {
/* trace shadow ray */
float3 shadow;
@ -87,7 +87,7 @@ ccl_device void kernel_branched_path_surface_connect_light(KernelGlobals *kg, RN
LightSample ls;
light_sample(kg, light_t, light_u, light_v, ccl_fetch(sd, time), ccl_fetch(sd, P), state->bounce, &ls);
if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp, state->bounce, state->transparent_bounce)) {
if(direct_emission(kg, sd, &ls, state, &light_ray, &L_light, &is_lamp)) {
/* trace shadow ray */
float3 shadow;
@ -109,7 +109,7 @@ ccl_device void kernel_branched_path_surface_connect_light(KernelGlobals *kg, RN
light_sample(kg, light_t, light_u, light_v, ccl_fetch(sd, time), ccl_fetch(sd, P), state->bounce, &ls);
/* sample random light */
if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp, state->bounce, state->transparent_bounce)) {
if(direct_emission(kg, sd, &ls, state, &light_ray, &L_light, &is_lamp)) {
/* trace shadow ray */
float3 shadow;
@ -206,7 +206,7 @@ ccl_device_inline void kernel_path_surface_connect_light(KernelGlobals *kg, ccl_
LightSample ls;
light_sample(kg, light_t, light_u, light_v, ccl_fetch(sd, time), ccl_fetch(sd, P), state->bounce, &ls);
if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp, state->bounce, state->transparent_bounce)) {
if(direct_emission(kg, sd, &ls, state, &light_ray, &L_light, &is_lamp)) {
/* trace shadow ray */
float3 shadow;

8
intern/cycles/kernel/kernel_path_volume.h
View File

@ -44,7 +44,7 @@ ccl_device void kernel_path_volume_connect_light(KernelGlobals *kg, RNG *rng,
if(ls.pdf == 0.0f)
return;
if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp, state->bounce, state->transparent_bounce)) {
if(direct_emission(kg, sd, &ls, state, &light_ray, &L_light, &is_lamp)) {
/* trace shadow ray */
float3 shadow;
@ -160,7 +160,7 @@ ccl_device void kernel_branched_path_volume_connect_light(KernelGlobals *kg, RNG
if(ls.pdf == 0.0f)
continue;
if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp, state->bounce, state->transparent_bounce)) {
if(direct_emission(kg, sd, &ls, state, &light_ray, &L_light, &is_lamp)) {
/* trace shadow ray */
float3 shadow;
@ -211,7 +211,7 @@ ccl_device void kernel_branched_path_volume_connect_light(KernelGlobals *kg, RNG
if(ls.pdf == 0.0f)
continue;
if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp, state->bounce, state->transparent_bounce)) {
if(direct_emission(kg, sd, &ls, state, &light_ray, &L_light, &is_lamp)) {
/* trace shadow ray */
float3 shadow;
@ -251,7 +251,7 @@ ccl_device void kernel_branched_path_volume_connect_light(KernelGlobals *kg, RNG
return;
/* sample random light */
if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp, state->bounce, state->transparent_bounce)) {
if(direct_emission(kg, sd, &ls, state, &light_ray, &L_light, &is_lamp)) {
/* trace shadow ray */
float3 shadow;

41
intern/cycles/kernel/kernel_shader.h
View File

@ -49,7 +49,7 @@ ccl_device void shader_setup_object_transforms(KernelGlobals *kg, ShaderData *sd
#endif
ccl_device void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd,
const Intersection *isect, const Ray *ray, int bounce, int transparent_bounce)
const Intersection *isect, const Ray *ray)
{
#ifdef __INSTANCING__
ccl_fetch(sd, object) = (isect->object == PRIM_NONE)? kernel_tex_fetch(__prim_object, isect->prim): isect->object;
@ -66,8 +66,6 @@ ccl_device void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd,
ccl_fetch(sd, prim) = kernel_tex_fetch(__prim_index, isect->prim);
ccl_fetch(sd, ray_length) = isect->t;
ccl_fetch(sd, ray_depth) = bounce;
ccl_fetch(sd, transparent_depth) = transparent_bounce;
#ifdef __UV__
ccl_fetch(sd, u) = isect->u;
@ -227,7 +225,7 @@ ccl_device_inline void shader_setup_from_subsurface(KernelGlobals *kg, ShaderDat
ccl_device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd,
const float3 P, const float3 Ng, const float3 I,
int shader, int object, int prim, float u, float v, float t, float time, int bounce, int transparent_bounce)
int shader, int object, int prim, float u, float v, float t, float time)
{
/* vectors */
ccl_fetch(sd, P) = P;
@ -248,8 +246,6 @@ ccl_device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd,
ccl_fetch(sd, v) = v;
#endif
ccl_fetch(sd, ray_length) = t;
ccl_fetch(sd, ray_depth) = bounce;
ccl_fetch(sd, transparent_depth) = transparent_bounce;
/* detect instancing, for non-instanced the object index is -object-1 */
#ifdef __INSTANCING__
@ -347,12 +343,12 @@ ccl_device void shader_setup_from_displace(KernelGlobals *kg, ShaderData *sd,
/* watch out: no instance transform currently */
shader_setup_from_sample(kg, sd, P, Ng, I, shader, object, prim, u, v, 0.0f, TIME_INVALID, 0, 0);
shader_setup_from_sample(kg, sd, P, Ng, I, shader, object, prim, u, v, 0.0f, TIME_INVALID);
}
/* ShaderData setup from ray into background */
ccl_device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData *sd, const Ray *ray, int bounce, int transparent_bounce)
ccl_device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData *sd, const Ray *ray)
{
/* vectors */
ccl_fetch(sd, P) = ray->D;
@ -365,8 +361,6 @@ ccl_device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderDat
ccl_fetch(sd, time) = ray->time;
#endif
ccl_fetch(sd, ray_length) = 0.0f;
ccl_fetch(sd, ray_depth) = bounce;
ccl_fetch(sd, transparent_depth) = transparent_bounce;
#ifdef __INSTANCING__
ccl_fetch(sd, object) = PRIM_NONE;
@ -395,7 +389,7 @@ ccl_device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderDat
/* ShaderData setup from point inside volume */
#ifdef __VOLUME__
ccl_device_inline void shader_setup_from_volume(KernelGlobals *kg, ShaderData *sd, const Ray *ray, int bounce, int transparent_bounce)
ccl_device_inline void shader_setup_from_volume(KernelGlobals *kg, ShaderData *sd, const Ray *ray)
{
/* vectors */
sd->P = ray->P;
@ -408,8 +402,6 @@ ccl_device_inline void shader_setup_from_volume(KernelGlobals *kg, ShaderData *s
sd->time = ray->time;
#endif
sd->ray_length = 0.0f; /* todo: can we set this to some useful value? */
sd->ray_depth = bounce;
sd->transparent_depth = transparent_bounce;
#ifdef __INSTANCING__
sd->object = PRIM_NONE; /* todo: fill this for texture coordinates */
@ -826,19 +818,19 @@ ccl_device float3 shader_holdout_eval(KernelGlobals *kg, ShaderData *sd)
/* Surface Evaluation */
ccl_device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd,
float randb, int path_flag, ShaderContext ctx)
ccl_addr_space PathState *state, float randb, int path_flag, ShaderContext ctx)
{
ccl_fetch(sd, num_closure) = 0;
ccl_fetch(sd, randb_closure) = randb;
#ifdef __OSL__
if(kg->osl)
OSLShader::eval_surface(kg, sd, path_flag, ctx);
OSLShader::eval_surface(kg, sd, state, path_flag, ctx);
else
#endif
{
#ifdef __SVM__
svm_eval_nodes(kg, sd, SHADER_TYPE_SURFACE, path_flag);
svm_eval_nodes(kg, sd, state, SHADER_TYPE_SURFACE, path_flag);
#else
ccl_fetch_array(sd, closure, 0)->weight = make_float3(0.8f, 0.8f, 0.8f);
ccl_fetch_array(sd, closure, 0)->N = ccl_fetch(sd, N);
@ -851,21 +843,22 @@ ccl_device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd,
/* Background Evaluation */
ccl_device float3 shader_eval_background(KernelGlobals *kg, ShaderData *sd, int path_flag, ShaderContext ctx)
ccl_device float3 shader_eval_background(KernelGlobals *kg, ShaderData *sd,
ccl_addr_space PathState *state, int path_flag, ShaderContext ctx)
{
ccl_fetch(sd, num_closure) = 0;
ccl_fetch(sd, randb_closure) = 0.0f;
#ifdef __OSL__
if(kg->osl) {
return OSLShader::eval_background(kg, sd, path_flag, ctx);
return OSLShader::eval_background(kg, sd, state, path_flag, ctx);
}
else
#endif
{
#ifdef __SVM__
svm_eval_nodes(kg, sd, SHADER_TYPE_SURFACE, path_flag);
svm_eval_nodes(kg, sd, state, SHADER_TYPE_SURFACE, path_flag);
float3 eval = make_float3(0.0f, 0.0f, 0.0f);
@ -992,7 +985,7 @@ ccl_device int shader_phase_sample_closure(KernelGlobals *kg, const ShaderData *
/* Volume Evaluation */
ccl_device void shader_eval_volume(KernelGlobals *kg, ShaderData *sd,
VolumeStack *stack, int path_flag, ShaderContext ctx)
PathState *state, VolumeStack *stack, int path_flag, ShaderContext ctx)
{
/* reset closures once at the start, we will be accumulating the closures
* for all volumes in the stack into a single array of closures */
@ -1022,12 +1015,12 @@ ccl_device void shader_eval_volume(KernelGlobals *kg, ShaderData *sd,
#ifdef __SVM__
#ifdef __OSL__
if(kg->osl) {
OSLShader::eval_volume(kg, sd, path_flag, ctx);
OSLShader::eval_volume(kg, sd, state, path_flag, ctx);
}
else
#endif
{
svm_eval_nodes(kg, sd, SHADER_TYPE_VOLUME, path_flag);
svm_eval_nodes(kg, sd, state, SHADER_TYPE_VOLUME, path_flag);
}
#endif
@ -1041,7 +1034,7 @@ ccl_device void shader_eval_volume(KernelGlobals *kg, ShaderData *sd,
/* Displacement Evaluation */
ccl_device void shader_eval_displacement(KernelGlobals *kg, ShaderData *sd, ShaderContext ctx)
ccl_device void shader_eval_displacement(KernelGlobals *kg, ShaderData *sd, ccl_addr_space PathState *state, ShaderContext ctx)
{
ccl_fetch(sd, num_closure) = 0;
ccl_fetch(sd, randb_closure) = 0.0f;
@ -1054,7 +1047,7 @@ ccl_device void shader_eval_displacement(KernelGlobals *kg, ShaderData *sd, Shad
else
#endif
{
svm_eval_nodes(kg, sd, SHADER_TYPE_DISPLACEMENT, 0);
svm_eval_nodes(kg, sd, state, SHADER_TYPE_DISPLACEMENT, 0);
}
#endif
}

7
intern/cycles/kernel/kernel_shaderdata_vars.h
View File

@ -53,12 +53,6 @@ SD_VAR(float, time)
/* length of the ray being shaded */
SD_VAR(float, ray_length)
/* ray bounce depth */
SD_VAR(int, ray_depth)
/* ray transparent depth */
SD_VAR(int, transparent_depth)
#ifdef __RAY_DIFFERENTIALS__
/* differential of P. these are orthogonal to Ng, not N */
SD_VAR(differential3, dP)
@ -93,6 +87,7 @@ SD_VAR(differential3, ray_dP)
#ifdef __OSL__
SD_VAR(struct KernelGlobals *, osl_globals)
SD_VAR(struct PathState *, osl_path_state)
#endif
#undef SD_VAR

14
intern/cycles/kernel/kernel_shadow.h
View File

@ -107,11 +107,14 @@ ccl_device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *
/* setup shader data at surface */
ShaderData sd;
shader_setup_from_ray(kg, &sd, isect, ray, state->bounce+1, bounce);
shader_setup_from_ray(kg, &sd, isect, ray);
/* attenuation from transparent surface */
if(!(sd.flag & SD_HAS_ONLY_VOLUME)) {
shader_eval_surface(kg, &sd, 0.0f, PATH_RAY_SHADOW, SHADER_CONTEXT_SHADOW);
path_state_modify_bounce(state, true);
shader_eval_surface(kg, &sd, state, 0.0f, PATH_RAY_SHADOW, SHADER_CONTEXT_SHADOW);
path_state_modify_bounce(state, false);
throughput *= shader_bsdf_transparency(kg, &sd);
}
@ -253,11 +256,14 @@ ccl_device_inline bool shadow_blocked(KernelGlobals *kg, ccl_addr_space PathStat
ShaderData sd_object;
ShaderData *sd = &sd_object;
#endif
shader_setup_from_ray(kg, sd, isect, ray, state->bounce+1, bounce);
shader_setup_from_ray(kg, sd, isect, ray);
/* attenuation from transparent surface */
if(!(ccl_fetch(sd, flag) & SD_HAS_ONLY_VOLUME)) {
shader_eval_surface(kg, sd, 0.0f, PATH_RAY_SHADOW, SHADER_CONTEXT_SHADOW);
path_state_modify_bounce(state, true);
shader_eval_surface(kg, sd, state, 0.0f, PATH_RAY_SHADOW, SHADER_CONTEXT_SHADOW);
path_state_modify_bounce(state, false);
throughput *= shader_bsdf_transparency(kg, sd);
}

10
intern/cycles/kernel/kernel_subsurface.h
View File

@ -181,6 +181,7 @@ ccl_device float3 subsurface_color_pow(float3 color, float exponent)
ccl_device void subsurface_color_bump_blur(KernelGlobals *kg,
ShaderData *sd,
PathState *state,
int state_flag,
float3 *eval,
float3 *N)
@ -194,7 +195,7 @@ ccl_device void subsurface_color_bump_blur(KernelGlobals *kg,
if(bump || texture_blur > 0.0f) {
/* average color and normal at incoming point */
shader_eval_surface(kg, sd, 0.0f, state_flag, SHADER_CONTEXT_SSS);
shader_eval_surface(kg, sd, state, 0.0f, state_flag, SHADER_CONTEXT_SSS);
float3 in_color = shader_bssrdf_sum(sd, (bump)? N: NULL, NULL);
/* we simply divide out the average color and multiply with the average
@ -329,6 +330,7 @@ ccl_device void subsurface_scatter_multi_setup(KernelGlobals *kg,
SubsurfaceIntersection* ss_isect,
int hit,
ShaderData *sd,
PathState *state,
int state_flag,
ShaderClosure *sc,
bool all)
@ -339,14 +341,14 @@ ccl_device void subsurface_scatter_multi_setup(KernelGlobals *kg,
/* Optionally blur colors and bump mapping. */
float3 weight = ss_isect->weight[hit];
float3 N = sd->N;
subsurface_color_bump_blur(kg, sd, state_flag, &weight, &N);
subsurface_color_bump_blur(kg, sd, state, state_flag, &weight, &N);
/* Setup diffuse BSDF. */
subsurface_scatter_setup_diffuse_bsdf(sd, weight, true, N);
}
/* subsurface scattering step, from a point on the surface to another nearby point on the same object */
ccl_device void subsurface_scatter_step(KernelGlobals *kg, ShaderData *sd,
ccl_device void subsurface_scatter_step(KernelGlobals *kg, ShaderData *sd, PathState *state,
int state_flag, ShaderClosure *sc, uint *lcg_state, float disk_u, float disk_v, bool all)
{
float3 eval = make_float3(0.0f, 0.0f, 0.0f);
@ -429,7 +431,7 @@ ccl_device void subsurface_scatter_step(KernelGlobals *kg, ShaderData *sd,
/* optionally blur colors and bump mapping */
float3 N = sd->N;
subsurface_color_bump_blur(kg, sd, state_flag, &eval, &N);
subsurface_color_bump_blur(kg, sd, state, state_flag, &eval, &N);
/* setup diffuse bsdf */
subsurface_scatter_setup_diffuse_bsdf(sd, eval, (ss_isect.num_hits > 0), N);

16
intern/cycles/kernel/kernel_volume.h
View File

@ -39,7 +39,7 @@ typedef struct VolumeShaderCoefficients {
ccl_device bool volume_shader_extinction_sample(KernelGlobals *kg, ShaderData *sd, PathState *state, float3 P, float3 *extinction)
{
sd->P = P;
shader_eval_volume(kg, sd, state->volume_stack, PATH_RAY_SHADOW, SHADER_CONTEXT_SHADOW);
shader_eval_volume(kg, sd, state, state->volume_stack, PATH_RAY_SHADOW, SHADER_CONTEXT_SHADOW);
if(!(sd->flag & (SD_ABSORPTION|SD_SCATTER)))
return false;
@ -61,7 +61,7 @@ ccl_device bool volume_shader_extinction_sample(KernelGlobals *kg, ShaderData *s
ccl_device bool volume_shader_sample(KernelGlobals *kg, ShaderData *sd, PathState *state, float3 P, VolumeShaderCoefficients *coeff)
{
sd->P = P;
shader_eval_volume(kg, sd, state->volume_stack, state->flag, SHADER_CONTEXT_VOLUME);
shader_eval_volume(kg, sd, state, state->volume_stack, state->flag, SHADER_CONTEXT_VOLUME);
if(!(sd->flag & (SD_ABSORPTION|SD_SCATTER|SD_EMISSION)))
return false;
@ -222,7 +222,7 @@ ccl_device void kernel_volume_shadow_heterogeneous(KernelGlobals *kg, PathState
ccl_device_noinline void kernel_volume_shadow(KernelGlobals *kg, PathState *state, Ray *ray, float3 *throughput)
{
ShaderData sd;
shader_setup_from_volume(kg, &sd, ray, state->bounce, state->transparent_bounce);
shader_setup_from_volume(kg, &sd, ray);
if(volume_stack_is_heterogeneous(kg, state->volume_stack))
kernel_volume_shadow_heterogeneous(kg, state, ray, &sd, throughput);
@ -567,7 +567,7 @@ ccl_device_noinline VolumeIntegrateResult kernel_volume_integrate(KernelGlobals
* performance of rendering without volumes */
RNG tmp_rng = cmj_hash(*rng, state->rng_offset);
shader_setup_from_volume(kg, sd, ray, state->bounce, state->transparent_bounce);
shader_setup_from_volume(kg, sd, ray);
if(heterogeneous)
return kernel_volume_integrate_heterogeneous_distance(kg, state, ray, sd, L, throughput, &tmp_rng);
@ -1008,7 +1008,7 @@ ccl_device void kernel_volume_stack_init(KernelGlobals *kg,
for(uint hit = 0; hit < num_hits; ++hit, ++isect) {
ShaderData sd;
shader_setup_from_ray(kg, &sd, isect, &volume_ray, 0, 0);
shader_setup_from_ray(kg, &sd, isect, &volume_ray);
if(sd.flag & SD_BACKFACING) {
/* If ray exited the volume and never entered to that volume
* it means that camera is inside such a volume.
@ -1048,7 +1048,7 @@ ccl_device void kernel_volume_stack_init(KernelGlobals *kg,
}
ShaderData sd;
shader_setup_from_ray(kg, &sd, &isect, &volume_ray, 0, 0);
shader_setup_from_ray(kg, &sd, &isect, &volume_ray);
if(sd.flag & SD_BACKFACING) {
/* If ray exited the volume and never entered to that volume
* it means that camera is inside such a volume.
@ -1162,7 +1162,7 @@ ccl_device void kernel_volume_stack_update_for_subsurface(KernelGlobals *kg,
for(uint hit = 0; hit < num_hits; ++hit, ++isect) {
ShaderData sd;
shader_setup_from_ray(kg, &sd, isect, &volume_ray, 0, 0);
shader_setup_from_ray(kg, &sd, isect, &volume_ray);
kernel_volume_stack_enter_exit(kg, &sd, stack);
}
}
@ -1173,7 +1173,7 @@ ccl_device void kernel_volume_stack_update_for_subsurface(KernelGlobals *kg,
scene_intersect_volume(kg, &volume_ray, &isect))
{
ShaderData sd;
shader_setup_from_ray(kg, &sd, &isect, &volume_ray, 0, 0);
shader_setup_from_ray(kg, &sd, &isect, &volume_ray);
kernel_volume_stack_enter_exit(kg, &sd, stack);
/* Move ray forward. */

10
intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
View File

@ -60,12 +60,6 @@ __kernel void kernel_ocl_path_trace_data_init(
ccl_global float *ray_length_sd,
ccl_global float *ray_length_sd_DL_shadow,
ccl_global int *ray_depth_sd,
ccl_global int *ray_depth_sd_DL_shadow,
ccl_global int *transparent_depth_sd,
ccl_global int *transparent_depth_sd_DL_shadow,
/* Ray differentials. */
ccl_global differential3 *dP_sd,
ccl_global differential3 *dP_sd_DL_shadow,
@ -170,10 +164,6 @@ __kernel void kernel_ocl_path_trace_data_init(
time_sd_DL_shadow,
ray_length_sd,
ray_length_sd_DL_shadow,
ray_depth_sd,
ray_depth_sd_DL_shadow,
transparent_depth_sd,
transparent_depth_sd_DL_shadow,
/* Ray differentials. */
dP_sd,

19
intern/cycles/kernel/osl/osl_services.cpp
View File

@ -101,6 +101,7 @@ ustring OSLRenderServices::u_curve_tangent_normal("geom:curve_tangent_normal");
ustring OSLRenderServices::u_path_ray_length("path:ray_length");
ustring OSLRenderServices::u_path_ray_depth("path:ray_depth");
ustring OSLRenderServices::u_path_transparent_depth("path:transparent_depth");
ustring OSLRenderServices::u_path_transmission_depth("path:transmission_depth");
ustring OSLRenderServices::u_trace("trace");
ustring OSLRenderServices::u_hit("hit");
ustring OSLRenderServices::u_hitdist("hitdist");
@ -726,12 +727,20 @@ bool OSLRenderServices::get_background_attribute(KernelGlobals *kg, ShaderData *
}
else if(name == u_path_ray_depth) {
/* Ray Depth */
int f = sd->ray_depth;
PathState *state = sd->osl_path_state;
int f = state->bounce;
return set_attribute_int(f, type, derivatives, val);
}
else if(name == u_path_transparent_depth) {
/* Transparent Ray Depth */
int f = sd->transparent_depth;
PathState *state = sd->osl_path_state;
int f = state->transparent_bounce;
return set_attribute_int(f, type, derivatives, val);
}
else if(name == u_path_transmission_depth) {
/* Transmission Ray Depth */
PathState *state = sd->osl_path_state;
int f = state->transmission_bounce;
return set_attribute_int(f, type, derivatives, val);
}
else if(name == u_ndc) {
@ -1257,11 +1266,7 @@ bool OSLRenderServices::getmessage(OSL::ShaderGlobals *sg, ustring source, ustri
if(!tracedata->setup) {
/* lazy shader data setup */
ShaderData *original_sd = (ShaderData *)(sg->renderstate);
int bounce = original_sd->ray_depth + 1;
int transparent_bounce = original_sd->transparent_depth;
shader_setup_from_ray(kg, sd, &tracedata->isect, &tracedata->ray, bounce, transparent_bounce);
shader_setup_from_ray(kg, sd, &tracedata->isect, &tracedata->ray);
tracedata->setup = true;
}

1
intern/cycles/kernel/osl/osl_services.h
View File

@ -170,6 +170,7 @@ public:
static ustring u_path_ray_length;
static ustring u_path_ray_depth;
static ustring u_path_transparent_depth;
static ustring u_path_transmission_depth;
static ustring u_trace;
static ustring u_hit;
static ustring u_hitdist;

20
intern/cycles/kernel/osl/osl_shader.cpp
View File

@ -92,7 +92,7 @@ void OSLShader::thread_free(KernelGlobals *kg)
/* Globals */
static void shaderdata_to_shaderglobals(KernelGlobals *kg, ShaderData *sd,
static void shaderdata_to_shaderglobals(KernelGlobals *kg, ShaderData *sd, PathState *state,
int path_flag, OSLThreadData *tdata)
{
OSL::ShaderGlobals *globals = &tdata->globals;
@ -136,6 +136,7 @@ static void shaderdata_to_shaderglobals(KernelGlobals *kg, ShaderData *sd,
/* used by renderservices */
sd->osl_globals = kg;
sd->osl_path_state = state;
}
/* Surface */
@ -317,11 +318,11 @@ static void flatten_surface_closure_tree(ShaderData *sd, int path_flag,
}
}
void OSLShader::eval_surface(KernelGlobals *kg, ShaderData *sd, int path_flag, ShaderContext ctx)
void OSLShader::eval_surface(KernelGlobals *kg, ShaderData *sd, PathState *state, int path_flag, ShaderContext ctx)
{
/* setup shader globals from shader data */
OSLThreadData *tdata = kg->osl_tdata;
shaderdata_to_shaderglobals(kg, sd, path_flag, tdata);
shaderdata_to_shaderglobals(kg, sd, state, path_flag, tdata);
/* execute shader for this point */
OSL::ShadingSystem *ss = (OSL::ShadingSystem*)kg->osl_ss;
@ -377,11 +378,11 @@ static float3 flatten_background_closure_tree(const OSL::ClosureColor *closure)
return make_float3(0.0f, 0.0f, 0.0f);
}
float3 OSLShader::eval_background(KernelGlobals *kg, ShaderData *sd, int path_flag, ShaderContext ctx)
float3 OSLShader::eval_background(KernelGlobals *kg, ShaderData *sd, PathState *state, int path_flag, ShaderContext ctx)
{
/* setup shader globals from shader data */
OSLThreadData *tdata = kg->osl_tdata;
shaderdata_to_shaderglobals(kg, sd, path_flag, tdata);
shaderdata_to_shaderglobals(kg, sd, state, path_flag, tdata);
/* execute shader for this point */
OSL::ShadingSystem *ss = (OSL::ShadingSystem*)kg->osl_ss;
@ -486,11 +487,11 @@ static void flatten_volume_closure_tree(ShaderData *sd,
}
}
void OSLShader::eval_volume(KernelGlobals *kg, ShaderData *sd, int path_flag, ShaderContext ctx)
void OSLShader::eval_volume(KernelGlobals *kg, ShaderData *sd, PathState *state, int path_flag, ShaderContext ctx)
{
/* setup shader globals from shader data */
OSLThreadData *tdata = kg->osl_tdata;
shaderdata_to_shaderglobals(kg, sd, path_flag, tdata);
shaderdata_to_shaderglobals(kg, sd, state, path_flag, tdata);
/* execute shader */
OSL::ShadingSystem *ss = (OSL::ShadingSystem*)kg->osl_ss;
@ -512,7 +513,10 @@ void OSLShader::eval_displacement(KernelGlobals *kg, ShaderData *sd, ShaderConte
{
/* setup shader globals from shader data */
OSLThreadData *tdata = kg->osl_tdata;
shaderdata_to_shaderglobals(kg, sd, 0, tdata);
PathState state = {0};
shaderdata_to_shaderglobals(kg, sd, &state, 0, tdata);
/* execute shader */
OSL::ShadingSystem *ss = (OSL::ShadingSystem*)kg->osl_ss;

6
intern/cycles/kernel/osl/osl_shader.h
View File

@ -53,9 +53,9 @@ public:
static void thread_free(KernelGlobals *kg);
/* eval */
static void eval_surface(KernelGlobals *kg, ShaderData *sd, int path_flag, ShaderContext ctx);
static float3 eval_background(KernelGlobals *kg, ShaderData *sd, int path_flag, ShaderContext ctx);
static void eval_volume(KernelGlobals *kg, ShaderData *sd, int path_flag, ShaderContext ctx);
static void eval_surface(KernelGlobals *kg, ShaderData *sd, PathState *state, int path_flag, ShaderContext ctx);
static float3 eval_background(KernelGlobals *kg, ShaderData *sd, PathState *state, int path_flag, ShaderContext ctx);
static void eval_volume(KernelGlobals *kg, ShaderData *sd, PathState *state, int path_flag, ShaderContext ctx);
static void eval_displacement(KernelGlobals *kg, ShaderData *sd, ShaderContext ctx);
/* sample & eval */

7
intern/cycles/kernel/shaders/node_light_path.osl
View File

@ -27,7 +27,8 @@ shader node_light_path(
output float IsVolumeScatterRay = 0.0,
output float RayLength = 0.0,
output float RayDepth = 0.0,
output float TransparentDepth = 0.0)
output float TransparentDepth = 0.0,
output float TransmissionDepth = 0.0)
{
IsCameraRay = raytype("camera");
IsShadowRay = raytype("shadow");
@ -47,5 +48,9 @@ shader node_light_path(
int transparent_depth;
getattribute("path:transparent_depth", transparent_depth);
TransparentDepth = (float)transparent_depth;
int transmission_depth;
getattribute("path:transmission_depth", transmission_depth);
TransmissionDepth = (float)transmission_depth;
}

12
intern/cycles/kernel/split/kernel_data_init.h
View File

@ -94,12 +94,6 @@ ccl_device void kernel_data_init(
ccl_global float *ray_length_sd,
ccl_global float *ray_length_sd_DL_shadow,
ccl_global int *ray_depth_sd,
ccl_global int *ray_depth_sd_DL_shadow,
ccl_global int *transparent_depth_sd,
ccl_global int *transparent_depth_sd_DL_shadow,
/* Ray differentials. */
ccl_global differential3 *dP_sd,
ccl_global differential3 *dP_sd_DL_shadow,
@ -219,12 +213,6 @@ ccl_device void kernel_data_init(
sd->ray_length = ray_length_sd;
sd_DL_shadow->ray_length = ray_length_sd_DL_shadow;
sd->ray_depth = ray_depth_sd;
sd_DL_shadow->ray_depth = ray_depth_sd_DL_shadow;
sd->transparent_depth = transparent_depth_sd;
sd_DL_shadow->transparent_depth = transparent_depth_sd_DL_shadow;
#ifdef __RAY_DIFFERENTIALS__
sd->dP = dP_sd;
sd_DL_shadow->dP = dP_sd_DL_shadow;

4
intern/cycles/kernel/split/kernel_direct_lighting.h
View File

@ -90,8 +90,8 @@ ccl_device char kernel_direct_lighting(
BsdfEval L_light;
bool is_lamp;
if(direct_emission(kg, sd, &ls, &light_ray, &L_light, &is_lamp,
state->bounce, state->transparent_bounce, sd_DL))
if(direct_emission(kg, sd, &ls, state, &light_ray, &L_light, &is_lamp,
sd_DL))
{
/* Write intermediate data to global memory to access from
* the next kernel.

42
intern/cycles/kernel/split/kernel_lamp_emission.h
View File

@ -59,28 +59,28 @@ ccl_device void kernel_lamp_emission(
IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND))
{
PathRadiance *L = &PathRadiance_coop[ray_index];
ccl_global PathState *state = &PathState_coop[ray_index];
float3 throughput = throughput_coop[ray_index];
Ray ray = Ray_coop[ray_index];
PathState state = PathState_coop[ray_index];
#ifdef __LAMP_MIS__
if(kernel_data.integrator.use_lamp_mis && !(state.flag & PATH_RAY_CAMERA)) {
if(kernel_data.integrator.use_lamp_mis && !(state->flag & PATH_RAY_CAMERA)) {
/* ray starting from previous non-transparent bounce */
Ray light_ray;
light_ray.P = ray.P - state.ray_t*ray.D;
state.ray_t += Intersection_coop[ray_index].t;
light_ray.P = ray.P - state->ray_t*ray.D;
state->ray_t += Intersection_coop[ray_index].t;
light_ray.D = ray.D;
light_ray.t = state.ray_t;
light_ray.t = state->ray_t;
light_ray.time = ray.time;
light_ray.dD = ray.dD;
light_ray.dP = ray.dP;
/* intersect with lamp */
float3 emission;
if(indirect_lamp_emission(kg, &state, &light_ray, &emission, sd)) {
path_radiance_accum_emission(L, throughput, emission, state.bounce);
if(indirect_lamp_emission(kg, state, &light_ray, &emission, sd)) {
path_radiance_accum_emission(L, throughput, emission, state->bounce);
}
}
#endif /* __LAMP_MIS__ */
@ -89,14 +89,14 @@ ccl_device void kernel_lamp_emission(
#if 0
#ifdef __VOLUME__
/* volume attenuation, emission, scatter */
if(state.volume_stack[0].shader != SHADER_NONE) {
if(state->volume_stack[0].shader != SHADER_NONE) {
Ray volume_ray = ray;
volume_ray.t = (hit)? isect.t: FLT_MAX;
bool heterogeneous = volume_stack_is_heterogeneous(kg, state.volume_stack);
bool heterogeneous = volume_stack_is_heterogeneous(kg, state->volume_stack);
#ifdef __VOLUME_DECOUPLED__
int sampling_method = volume_stack_sampling_method(kg, state.volume_stack);
int sampling_method = volume_stack_sampling_method(kg, state->volume_stack);
bool decoupled = kernel_volume_use_decoupled(kg, heterogeneous, true, sampling_method);
if(decoupled) {
@ -104,15 +104,15 @@ ccl_device void kernel_lamp_emission(
VolumeSegment volume_segment;
ShaderData volume_sd;
shader_setup_from_volume(kg, &volume_sd, &volume_ray, state.bounce, state.transparent_bounce);
kernel_volume_decoupled_record(kg, &state,
shader_setup_from_volume(kg, &volume_sd, &volume_ray);
kernel_volume_decoupled_record(kg, state,
&volume_ray, &volume_sd, &volume_segment, heterogeneous);
volume_segment.sampling_method = sampling_method;
/* emission */
if(volume_segment.closure_flag & SD_EMISSION)
path_radiance_accum_emission(&L, throughput, volume_segment.accum_emission, state.bounce);
path_radiance_accum_emission(&L, throughput, volume_segment.accum_emission, state->bounce);
/* scattering */
VolumeIntegrateResult result = VOLUME_PATH_ATTENUATED;
@ -122,16 +122,16 @@ ccl_device void kernel_lamp_emission(
/* direct light sampling */
kernel_branched_path_volume_connect_light(kg, rng, &volume_sd,
throughput, &state, &L, 1.0f, all, &volume_ray, &volume_segment);
throughput, state, &L, 1.0f, all, &volume_ray, &volume_segment);
/* indirect sample. if we use distance sampling and take just
* one sample for direct and indirect light, we could share
* this computation, but makes code a bit complex */
float rphase = path_state_rng_1D_for_decision(kg, rng, &state, PRNG_PHASE);
float rscatter = path_state_rng_1D_for_decision(kg, rng, &state, PRNG_SCATTER_DISTANCE);
float rphase = path_state_rng_1D_for_decision(kg, rng, state, PRNG_PHASE);
float rscatter = path_state_rng_1D_for_decision(kg, rng, state, PRNG_SCATTER_DISTANCE);
result = kernel_volume_decoupled_scatter(kg,
&state, &volume_ray, &volume_sd, &throughput,
state, &volume_ray, &volume_sd, &throughput,
rphase, rscatter, &volume_segment, NULL, true);
}
@ -142,7 +142,7 @@ ccl_device void kernel_lamp_emission(
kernel_volume_decoupled_free(kg, &volume_segment);
if(result == VOLUME_PATH_SCATTERED) {
if(kernel_path_volume_bounce(kg, rng, &volume_sd, &throughput, &state, &L, &ray))
if(kernel_path_volume_bounce(kg, rng, &volume_sd, &throughput, state, &L, &ray))
continue;
else
break;
@ -154,15 +154,15 @@ ccl_device void kernel_lamp_emission(
/* integrate along volume segment with distance sampling */
ShaderData volume_sd;
VolumeIntegrateResult result = kernel_volume_integrate(
kg, &state, &volume_sd, &volume_ray, &L, &throughput, rng, heterogeneous);
kg, state, &volume_sd, &volume_ray, &L, &throughput, rng, heterogeneous);
#ifdef __VOLUME_SCATTER__
if(result == VOLUME_PATH_SCATTERED) {
/* direct lighting */
kernel_path_volume_connect_light(kg, rng, &volume_sd, throughput, &state, &L);
kernel_path_volume_connect_light(kg, rng, &volume_sd, throughput, state, &L);
/* indirect light bounce */
if(kernel_path_volume_bounce(kg, rng, &volume_sd, &throughput, &state, &L, &ray))
if(kernel_path_volume_bounce(kg, rng, &volume_sd, &throughput, state, &L, &ray))
continue;
else
break;

6
intern/cycles/kernel/split/kernel_shader_eval.h
View File

@ -63,10 +63,8 @@ ccl_device void kernel_shader_eval(
shader_setup_from_ray(kg,
sd,
isect,
&ray,
state->bounce,
state->transparent_bounce);
&ray);
float rbsdf = path_state_rng_1D_for_decision(kg, rng, state, PRNG_BSDF);
shader_eval_surface(kg, sd, rbsdf, state->flag, SHADER_CONTEXT_MAIN);
shader_eval_surface(kg, sd, state, rbsdf, state->flag, SHADER_CONTEXT_MAIN);
}
}

4
intern/cycles/kernel/svm/svm.h
View File

@ -187,7 +187,7 @@ CCL_NAMESPACE_BEGIN
#define NODES_FEATURE(feature) ((__NODES_FEATURES__ & (feature)) != 0)
/* Main Interpreter Loop */
ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type, int path_flag)
ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ccl_addr_space PathState *state, ShaderType type, int path_flag)
{
float stack[SVM_STACK_SIZE];
int offset = ccl_fetch(sd, shader) & SHADER_MASK;
@ -335,7 +335,7 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, Shade
svm_node_brightness(sd, stack, node.y, node.z, node.w);
break;
case NODE_LIGHT_PATH:
svm_node_light_path(sd, stack, node.y, node.z, path_flag);
svm_node_light_path(sd, state, stack, node.y, node.z, path_flag);
break;
case NODE_OBJECT_INFO:
svm_node_object_info(kg, sd, stack, node.y, node.z);

7
intern/cycles/kernel/svm/svm_light_path.h
View File

@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN
/* Light Path Node */
ccl_device void svm_node_light_path(ShaderData *sd, float *stack, uint type, uint out_offset, int path_flag)
ccl_device void svm_node_light_path(ShaderData *sd, ccl_addr_space PathState *state, float *stack, uint type, uint out_offset, int path_flag)
{
float info = 0.0f;
@ -33,8 +33,9 @@ ccl_device void svm_node_light_path(ShaderData *sd, float *stack, uint type, uin
case NODE_LP_volume_scatter: info = (path_flag & PATH_RAY_VOLUME_SCATTER)? 1.0f: 0.0f; break;
case NODE_LP_backfacing: info = (ccl_fetch(sd, flag) & SD_BACKFACING)? 1.0f: 0.0f; break;
case NODE_LP_ray_length: info = ccl_fetch(sd, ray_length); break;
case NODE_LP_ray_depth: info = (float)ccl_fetch(sd, ray_depth); break;
case NODE_LP_ray_transparent: info = (float)ccl_fetch(sd, transparent_depth); break;
case NODE_LP_ray_depth: info = (float)state->bounce; break;
case NODE_LP_ray_transparent: info = (float)state->transparent_bounce; break;
case NODE_LP_ray_transmission: info = (float)state->transmission_bounce; break;
}
stack_store_float(stack, out_offset, info);

3
intern/cycles/kernel/svm/svm_types.h
View File

@ -183,7 +183,8 @@ typedef enum NodeLightPath {
NODE_LP_backfacing,
NODE_LP_ray_length,
NODE_LP_ray_depth,
NODE_LP_ray_transparent
NODE_LP_ray_transparent,
NODE_LP_ray_transmission,
} NodeLightPath;
typedef enum NodeLightFalloff {

7
intern/cycles/render/nodes.cpp
View File

@ -2893,6 +2893,7 @@ LightPathNode::LightPathNode()
add_output("Ray Length", SHADER_SOCKET_FLOAT);
add_output("Ray Depth", SHADER_SOCKET_FLOAT);
add_output("Transparent Depth", SHADER_SOCKET_FLOAT);
add_output("Transmission Depth", SHADER_SOCKET_FLOAT);
}
void LightPathNode::compile(SVMCompiler& compiler)
@ -2965,6 +2966,12 @@ void LightPathNode::compile(SVMCompiler& compiler)
compiler.stack_assign(out);
compiler.add_node(NODE_LIGHT_PATH, NODE_LP_ray_transparent, out->stack_offset);
}
out = output("Transmission Depth");
if(!out->links.empty()) {
compiler.stack_assign(out);
compiler.add_node(NODE_LIGHT_PATH, NODE_LP_ray_transmission, out->stack_offset);
}
}
void LightPathNode::compile(OSLCompiler& compiler)

4
source/blender/gpu/shaders/gpu_shader_material.glsl
View File

@ -2560,7 +2560,8 @@ void node_light_path(
out float is_transmission_ray,
out float ray_length,
out float ray_depth,
out float transparent_depth)
out float transparent_depth,
out float transmission_depth)
{
is_camera_ray = 1.0;
is_shadow_ray = 0.0;
@ -2572,6 +2573,7 @@ void node_light_path(
ray_length = 1.0;
ray_depth = 1.0;
transparent_depth = 1.0;
transmission_depth = 1.0;
}
void node_light_falloff(float strength, float tsmooth, out float quadratic, out float linear, out float constant)

1
source/blender/nodes/shader/nodes/node_shader_light_path.c
View File

@ -40,6 +40,7 @@ static bNodeSocketTemplate sh_node_light_path_out[] = {
{ SOCK_FLOAT, 0, N_("Ray Length"), 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
{ SOCK_FLOAT, 0, N_("Ray Depth"), 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
{ SOCK_FLOAT, 0, N_("Transparent Depth"), 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
{ SOCK_FLOAT, 0, N_("Transmission Depth"), 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
{ -1, 0, "" }
};