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

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. :)

25 files changed, 185 insertions, 182 deletions

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 {