diff options
29 files changed, 197 insertions, 204 deletions
diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp index c49583b07f3..64a311f2c13 100644 --- a/intern/cycles/device/device_opencl.cpp +++ b/intern/cycles/device/device_opencl.cpp @@ -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 += diff --git a/intern/cycles/kernel/kernel_bake.h b/intern/cycles/kernel/kernel_bake.h index b54afbd21b8..2a2220ceb99 100644 --- a/intern/cycles/kernel/kernel_bake.h +++ b/intern/cycles/kernel/kernel_bake.h @@ -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 */ diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h index 3863cf221c5..47d357215cf 100644 --- a/intern/cycles/kernel/kernel_emission.h +++ b/intern/cycles/kernel/kernel_emission.h @@ -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__ diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h index 82cad788eb0..fc32201596f 100644 --- a/intern/cycles/kernel/kernel_path.h +++ b/intern/cycles/kernel/kernel_path.h @@ -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__ diff --git a/intern/cycles/kernel/kernel_path_branched.h b/intern/cycles/kernel/kernel_path_branched.h index cc997e5eb47..7c685e3b191 100644 --- a/intern/cycles/kernel/kernel_path_branched.h +++ b/intern/cycles/kernel/kernel_path_branched.h @@ -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 */ diff --git a/intern/cycles/kernel/kernel_path_state.h b/intern/cycles/kernel/kernel_path_state.h index 15efb2371de..540f289eae0 100644 --- a/intern/cycles/kernel/kernel_path_state.h +++ b/intern/cycles/kernel/kernel_path_state.h @@ -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 diff --git a/intern/cycles/kernel/kernel_path_surface.h b/intern/cycles/kernel/kernel_path_surface.h index 4485be8a5b3..10af2537f69 100644 --- a/intern/cycles/kernel/kernel_path_surface.h +++ b/intern/cycles/kernel/kernel_path_surface.h @@ -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; diff --git a/intern/cycles/kernel/kernel_path_volume.h b/intern/cycles/kernel/kernel_path_volume.h index 82dc0f97622..fc4cd151bd3 100644 --- a/intern/cycles/kernel/kernel_path_volume.h +++ b/intern/cycles/kernel/kernel_path_volume.h @@ -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; diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h index 25e5822a21b..63e5ace8f94 100644 --- a/intern/cycles/kernel/kernel_shader.h +++ b/intern/cycles/kernel/kernel_shader.h @@ -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 } diff --git a/intern/cycles/kernel/kernel_shaderdata_vars.h b/intern/cycles/kernel/kernel_shaderdata_vars.h index b157b82e023..3f63f94912e 100644 --- a/intern/cycles/kernel/kernel_shaderdata_vars.h +++ b/intern/cycles/kernel/kernel_shaderdata_vars.h @@ -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 diff --git a/intern/cycles/kernel/kernel_shadow.h b/intern/cycles/kernel/kernel_shadow.h index 2811a8348ca..cb5fcf813f3 100644 --- a/intern/cycles/kernel/kernel_shadow.h +++ b/intern/cycles/kernel/kernel_shadow.h @@ -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); } diff --git a/intern/cycles/kernel/kernel_subsurface.h b/intern/cycles/kernel/kernel_subsurface.h index 27a341aa37a..e5572be114a 100644 --- a/intern/cycles/kernel/kernel_subsurface.h +++ b/intern/cycles/kernel/kernel_subsurface.h @@ -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); diff --git a/intern/cycles/kernel/kernel_volume.h b/intern/cycles/kernel/kernel_volume.h index 0a74a9deba9..a8cd7b0b4c0 100644 --- a/intern/cycles/kernel/kernel_volume.h +++ b/intern/cycles/kernel/kernel_volume.h @@ -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. */ diff --git a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl index 8329aa98d35..b8138676acd 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl @@ -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, diff --git a/intern/cycles/kernel/osl/osl_services.cpp b/intern/cycles/kernel/osl/osl_services.cpp index 8838c28ca13..0595282af64 100644 --- a/intern/cycles/kernel/osl/osl_services.cpp +++ b/intern/cycles/kernel/osl/osl_services.cpp @@ -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; } diff --git a/intern/cycles/kernel/osl/osl_services.h b/intern/cycles/kernel/osl/osl_services.h index f35ddcaef91..1d2f1556c3b 100644 --- a/intern/cycles/kernel/osl/osl_services.h +++ b/intern/cycles/kernel/osl/osl_services.h @@ -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; diff --git a/intern/cycles/kernel/osl/osl_shader.cpp b/intern/cycles/kernel/osl/osl_shader.cpp index 2f234aa25ea..950786ab213 100644 --- a/intern/cycles/kernel/osl/osl_shader.cpp +++ b/intern/cycles/kernel/osl/osl_shader.cpp @@ -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; diff --git a/intern/cycles/kernel/osl/osl_shader.h b/intern/cycles/kernel/osl/osl_shader.h index 15dd74f9d38..7d26cd40da5 100644 --- a/intern/cycles/kernel/osl/osl_shader.h +++ b/intern/cycles/kernel/osl/osl_shader.h @@ -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 */ diff --git a/intern/cycles/kernel/shaders/node_light_path.osl b/intern/cycles/kernel/shaders/node_light_path.osl index 99a92c4f403..a021a40467d 100644 --- a/intern/cycles/kernel/shaders/node_light_path.osl +++ b/intern/cycles/kernel/shaders/node_light_path.osl @@ -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; } diff --git a/intern/cycles/kernel/split/kernel_data_init.h b/intern/cycles/kernel/split/kernel_data_init.h index a5db6a83283..27defe2fb9a 100644 --- a/intern/cycles/kernel/split/kernel_data_init.h +++ b/intern/cycles/kernel/split/kernel_data_init.h @@ -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; diff --git a/intern/cycles/kernel/split/kernel_direct_lighting.h b/intern/cycles/kernel/split/kernel_direct_lighting.h index 34000517006..20d1728f9de 100644 --- a/intern/cycles/kernel/split/kernel_direct_lighting.h +++ b/intern/cycles/kernel/split/kernel_direct_lighting.h @@ -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. diff --git a/intern/cycles/kernel/split/kernel_lamp_emission.h b/intern/cycles/kernel/split/kernel_lamp_emission.h index 53dd6621127..6329f3ae943 100644 --- a/intern/cycles/kernel/split/kernel_lamp_emission.h +++ b/intern/cycles/kernel/split/kernel_lamp_emission.h @@ -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; diff --git a/intern/cycles/kernel/split/kernel_shader_eval.h b/intern/cycles/kernel/split/kernel_shader_eval.h index dca7d67ee93..e816a818915 100644 --- a/intern/cycles/kernel/split/kernel_shader_eval.h +++ b/intern/cycles/kernel/split/kernel_shader_eval.h @@ -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); } } diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h index 69b98060c62..633e1edfb19 100644 --- a/intern/cycles/kernel/svm/svm.h +++ b/intern/cycles/kernel/svm/svm.h @@ -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); diff --git a/intern/cycles/kernel/svm/svm_light_path.h b/intern/cycles/kernel/svm/svm_light_path.h index a235dd35224..94c9fddfab8 100644 --- a/intern/cycles/kernel/svm/svm_light_path.h +++ b/intern/cycles/kernel/svm/svm_light_path.h @@ -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); diff --git a/intern/cycles/kernel/svm/svm_types.h b/intern/cycles/kernel/svm/svm_types.h index 7df139103b7..e5fb20adecc 100644 --- a/intern/cycles/kernel/svm/svm_types.h +++ b/intern/cycles/kernel/svm/svm_types.h @@ -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 { diff --git a/intern/cycles/render/nodes.cpp b/intern/cycles/render/nodes.cpp index e35c4a9dd46..e6b91629515 100644 --- a/intern/cycles/render/nodes.cpp +++ b/intern/cycles/render/nodes.cpp @@ -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) diff --git a/source/blender/gpu/shaders/gpu_shader_material.glsl b/source/blender/gpu/shaders/gpu_shader_material.glsl index 5b62739b0eb..dfb1af0e863 100644 --- a/source/blender/gpu/shaders/gpu_shader_material.glsl +++ b/source/blender/gpu/shaders/gpu_shader_material.glsl @@ -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) diff --git a/source/blender/nodes/shader/nodes/node_shader_light_path.c b/source/blender/nodes/shader/nodes/node_shader_light_path.c index 023c6ea02a9..b1001cd3937 100644 --- a/source/blender/nodes/shader/nodes/node_shader_light_path.c +++ b/source/blender/nodes/shader/nodes/node_shader_light_path.c @@ -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, "" } }; |