Cycles:

* Add alpha pass output, to use set Transparent option in Film panel.
* Add Holdout closure (OSL terminology), this is like the Sky option in the
  internal renderer, objects with this closure show the background / zero
  alpha.
* Add option to use Gaussian instead of Box pixel filter in the UI.
* Remove camera response curves for now, they don't really belong here in
  the pipeline, should be moved to compositor.

* Output full float values for rendering now, previously was only byte precision.
* Add a patch from Thomas to get a preview passes option, but still disabled
  because it isn't quite working right yet.
* CUDA: don't compile shader graph evaluation inline.
* Convert tabs to spaces in python files.

22 files changed, 128 insertions, 67 deletions

diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp
index 656e90eb526..d9ba237d964 100644
--- a/intern/cycles/kernel/kernel.cpp
+++ b/intern/cycles/kernel/kernel.cpp
@@ -129,18 +129,6 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t
 		kg->__filter_table.data = (float*)mem;
 		kg->__filter_table.width = width;
 	}
-	else if(strcmp(name, "__response_curve_R") == 0) {
-		kg->__response_curve_R.data = (float*)mem;
-		kg->__response_curve_R.width = width;
-	}
-	else if(strcmp(name, "__response_curve_B") == 0) {
-		kg->__response_curve_B.data = (float*)mem;
-		kg->__response_curve_B.width = width;
-	}
-	else if(strcmp(name, "__response_curve_G") == 0) {
-		kg->__response_curve_G.data = (float*)mem;
-		kg->__response_curve_G.width = width;
-	}
 	else if(strcmp(name, "__sobol_directions") == 0) {
 		kg->__sobol_directions.data = (uint*)mem;
 		kg->__sobol_directions.width = width;
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index b7b29d46323..3a0eff5210c 100644
--- a/intern/cycles/kernel/kernel_compat_cuda.h
+++ b/intern/cycles/kernel/kernel_compat_cuda.h
@@ -33,6 +33,7 @@ CCL_NAMESPACE_BEGIN
 
 #define __device  __device__ __inline__
 #define __device_inline  __device__ __inline__
+#define __device_noinline  __device__ __noinline__
 #define __global
 #define __shared __shared__
 #define __constant
diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h
index 3d493c61fe4..287bf320881 100644
--- a/intern/cycles/kernel/kernel_compat_opencl.h
+++ b/intern/cycles/kernel/kernel_compat_opencl.h
@@ -30,6 +30,7 @@
 /* in opencl all functions are device functions, so leave this empty */
 #define __device
 #define __device_inline
+#define __device_noinline
 
 /* no assert in opencl */
 #define kernel_assert(cond)
diff --git a/intern/cycles/kernel/kernel_film.h b/intern/cycles/kernel/kernel_film.h
index 0bbd6c202a1..f6351a73295 100644
--- a/intern/cycles/kernel/kernel_film.h
+++ b/intern/cycles/kernel/kernel_film.h
@@ -20,21 +20,17 @@ CCL_NAMESPACE_BEGIN
 
 __device float4 film_map(KernelGlobals *kg, float4 irradiance, int pass)
 {
-	float scale = kernel_data.film.exposure*(1.0f/(pass+1));
+	float scale = 1.0f/(float)(pass+1);
+	float exposure = kernel_data.film.exposure;
 	float4 result = irradiance*scale;
 
-	if(kernel_data.film.use_response_curve) {
-		/* camera response curve */
-		result.x = kernel_tex_interp(__response_curve_R, result.x);
-		result.y = kernel_tex_interp(__response_curve_G, result.y);
-		result.z = kernel_tex_interp(__response_curve_B, result.z);
-	}
-	else {
-		/* conversion to srgb */
-		result.x = color_scene_linear_to_srgb(result.x);
-		result.y = color_scene_linear_to_srgb(result.y);
-		result.z = color_scene_linear_to_srgb(result.z);
-	}
+	/* conversion to srgb */
+	result.x = color_scene_linear_to_srgb(result.x*exposure);
+	result.y = color_scene_linear_to_srgb(result.y*exposure);
+	result.z = color_scene_linear_to_srgb(result.z*exposure);
+
+	/* clamp since alpha might be > 1.0 due to russian roulette */
+	result.w = clamp(result.w, 0.0f, 1.0f);
 
 	return result;
 }
@@ -47,7 +43,7 @@ __device uchar4 film_float_to_byte(float4 color)
 	result.x = (uchar)clamp(color.x*255.0f, 0.0f, 255.0f);
 	result.y = (uchar)clamp(color.y*255.0f, 0.0f, 255.0f);
 	result.z = (uchar)clamp(color.z*255.0f, 0.0f, 255.0f);
-	result.w = 255;
+	result.w = (uchar)clamp(color.w*255.0f, 0.0f, 255.0f);
 
 	return result;
 }
diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h
index 47aeea4f5c5..2cbb4ae306e 100644
--- a/intern/cycles/kernel/kernel_path.h
+++ b/intern/cycles/kernel/kernel_path.h
@@ -105,10 +105,11 @@ __device int path_flag_from_label(int path_flag, int label)
 	return path_flag;
 }
 
-__device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray ray, float3 throughput)
+__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray ray, float3 throughput)
 {
 	/* initialize */
 	float3 L = make_float3(0.0f, 0.0f, 0.0f);
+	float Ltransparent = 0.0f;
 
 #ifdef __EMISSION__
 	float ray_pdf = 0.0f;
@@ -123,14 +124,20 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 
 		if(!scene_intersect(kg, &ray, false, &isect)) {
 			/* eval background shader if nothing hit */
+			if(kernel_data.background.transparent && (path_flag & PATH_RAY_CAMERA)) {
+				Ltransparent += average(throughput);
+			}
+			else {
 #ifdef __BACKGROUND__
-			ShaderData sd;
-			shader_setup_from_background(kg, &sd, &ray);
-			L += throughput*shader_eval_background(kg, &sd, path_flag);
-			shader_release(kg, &sd);
+				ShaderData sd;
+				shader_setup_from_background(kg, &sd, &ray);
+				L += throughput*shader_eval_background(kg, &sd, path_flag);
+				shader_release(kg, &sd);
 #else
-			L += make_float3(0.8f, 0.8f, 0.8f);
+				L += make_float3(0.8f, 0.8f, 0.8f);
 #endif
+			}
+
 			break;
 		}
 
@@ -140,6 +147,22 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 		float rbsdf = path_rng(kg, rng, pass, rng_offset + PRNG_BSDF);
 		shader_eval_surface(kg, &sd, rbsdf, path_flag);
 
+#ifdef __HOLDOUT__
+		if((sd.flag & SD_HOLDOUT) && (path_flag & PATH_RAY_CAMERA)) {
+			float3 holdout_weight = shader_holdout_eval(kg, &sd);
+
+			if(kernel_data.background.transparent) {
+				Ltransparent += average(holdout_weight*throughput);
+			}
+			else {
+				ShaderData sd;
+				shader_setup_from_background(kg, &sd, &ray);
+				L += holdout_weight*throughput*shader_eval_background(kg, &sd, path_flag);
+				shader_release(kg, &sd);
+			}
+		}
+#endif
+
 #ifdef __EMISSION__
 		/* emission */
 		if(kernel_data.integrator.use_emission) {
@@ -166,6 +189,12 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 		}
 #endif
 
+		/* no BSDF? we can stop here */
+		if(!(sd.flag & SD_BSDF)) {
+			path_flag &= ~PATH_RAY_CAMERA;
+			break;
+		}
+
 		/* sample BSDF */
 		float bsdf_pdf;
 		float3 bsdf_eval;
@@ -180,8 +209,10 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 
 		shader_release(kg, &sd);
 
-		if(bsdf_pdf == 0.0f || is_zero(bsdf_eval))
+		if(bsdf_pdf == 0.0f || is_zero(bsdf_eval)) {
+			path_flag &= ~PATH_RAY_CAMERA;
 			break;
+		}
 
 		/* modify throughput */
 		throughput *= bsdf_eval/bsdf_pdf;
@@ -197,8 +228,10 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 		float probability = path_terminate_probability(kg, bounce, throughput);
 		float terminate = path_rng(kg, rng, pass, rng_offset + PRNG_TERMINATE);
 
-		if(terminate >= probability)
+		if(terminate >= probability) {
+			path_flag &= ~PATH_RAY_CAMERA;
 			break;
+		}
 
 		throughput /= probability;
 
@@ -212,7 +245,7 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 #endif
 	}
 
-	return L;
+	return make_float4(L.x, L.y, L.z, 1.0f - Ltransparent);
 }
 
 __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int pass, int x, int y)
@@ -236,25 +269,19 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl
 	/* integrate */
 #ifdef __MODIFY_TP__
 	float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, pass);
-	float3 L = kernel_path_integrate(kg, &rng, pass, ray, throughput)/throughput;
+	float4 L = kernel_path_integrate(kg, &rng, pass, ray, throughput)/throughput;
 #else
 	float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
-	float3 L = kernel_path_integrate(kg, &rng, pass, ray, throughput);
+	float4 L = kernel_path_integrate(kg, &rng, pass, ray, throughput);
 #endif
 
 	/* accumulate result in output buffer */
 	int index = x + y*kernel_data.cam.width;
 
-	float4 result;
-	result.x = L.x;
-	result.y = L.y;
-	result.z = L.z;
-	result.w = 1.0f;
-
 	if(pass == 0)
-		buffer[index] = result;
+		buffer[index] = L;
 	else
-		buffer[index] += result;
+		buffer[index] += L;
 
 	path_rng_end(kg, rng_state, rng, x, y);
 }
diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h
index c4fc65596b2..ed3aec21207 100644
--- a/intern/cycles/kernel/kernel_shader.h
+++ b/intern/cycles/kernel/kernel_shader.h
@@ -342,6 +342,26 @@ __device void shader_emissive_sample(KernelGlobals *kg, ShaderData *sd,
 	}
 }
 
+/* Holdout */
+
+__device float3 shader_holdout_eval(KernelGlobals *kg, ShaderData *sd)
+{
+#ifdef WITH_OSL
+	if(kg->osl.use) {
+		return OSLShader::holdout_eval(sd);
+	}
+	else
+#endif
+	{
+#ifdef __SVM__
+		if(sd->svm_closure == CLOSURE_HOLDOUT_ID)
+			return make_float3(1.0f, 1.0f, 1.0f);
+		else
+			return make_float3(0.0f, 0.0f, 0.0f);
+#endif
+	}
+}
+
 /* Surface Evaluation */
 
 __device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd,
diff --git a/intern/cycles/kernel/kernel_textures.h b/intern/cycles/kernel/kernel_textures.h
index bd44ed7eee5..dd4566ec184 100644
--- a/intern/cycles/kernel/kernel_textures.h
+++ b/intern/cycles/kernel/kernel_textures.h
@@ -38,9 +38,6 @@ KERNEL_TEX(uint4, texture_uint4, __svm_nodes)
 
 /* camera/film */
 KERNEL_TEX(float, texture_float, __filter_table)
-KERNEL_TEX(float, texture_float, __response_curve_R)
-KERNEL_TEX(float, texture_float, __response_curve_G)
-KERNEL_TEX(float, texture_float, __response_curve_B)
 
 /* sobol */
 KERNEL_TEX(uint, texture_uint, __sobol_directions)
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index c0cb3fc8a09..06581d593e4 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -37,6 +37,7 @@ CCL_NAMESPACE_BEGIN
 #ifndef __KERNEL_OPENCL__
 #define __SVM__
 #define __TEXTURES__
+#define __HOLDOUT__
 #endif
 #define __RAY_DIFFERENTIALS__
 #define __CAMERA_CLIPPING__
@@ -194,8 +195,10 @@ struct FlatClosure {
 enum ShaderDataFlag {
 	SD_BACKFACING = 1,		/* backside of surface? */
 	SD_EMISSION = 2,		/* have emissive closure? */
-	SD_BSDF_HAS_EVAL = 4,	/* have non-singular bsdf closure? */
-	SD_BSDF_GLOSSY = 8		/* have glossy bsdf */
+	SD_BSDF = 4,			/* have bsdf closure? */
+	SD_BSDF_HAS_EVAL = 8,	/* have non-singular bsdf closure? */
+	SD_BSDF_GLOSSY = 16,	/* have glossy bsdf */
+	SD_HOLDOUT = 32			/* have holdout closure? */
 };
 
 typedef struct ShaderData {
@@ -257,6 +260,8 @@ typedef struct ShaderData {
 		float emissive_sample_sum;
 		float volume_sample_sum;
 
+		float3 holdout_weight;
+
 		float randb;
 	} osl_closure;
 
@@ -313,14 +318,14 @@ typedef struct KernelCamera {
 
 typedef struct KernelFilm {
 	float exposure;
-	int use_response_curve;
-	int pad1, pad2;
+	int pad1, pad2, pad3;
 } KernelFilm;
 
 typedef struct KernelBackground {
 	/* only shader index */
 	int shader;
-	int pad1, pad2, pad3;
+	int transparent;
+	int pad1, pad2;
 } KernelBackground;
 
 typedef struct KernelSunSky {
diff --git a/intern/cycles/kernel/osl/osl_shader.cpp b/intern/cycles/kernel/osl/osl_shader.cpp
index a86946a680e..007d14b526b 100644
--- a/intern/cycles/kernel/osl/osl_shader.cpp
+++ b/intern/cycles/kernel/osl/osl_shader.cpp
@@ -145,9 +145,11 @@ static void flatten_surface_closure_tree(ShaderData *sd, bool no_glossy,
 
 					/* scattering flags */
 					if(scattering == OSL::Labels::DIFFUSE)
-						sd->flag |= SD_BSDF_HAS_EVAL;
+						sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL;
 					else if(scattering == OSL::Labels::GLOSSY)
-						sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+						sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+					else
+						sd->flag |= SD_BSDF;
 
 					/* add */
 					sd->osl_closure.bsdf[sd->osl_closure.num_bsdf++] = flat;
@@ -170,8 +172,11 @@ static void flatten_surface_closure_tree(ShaderData *sd, bool no_glossy,
 					sd->osl_closure.emissive[sd->osl_closure.num_emissive++] = flat;
 					break;
 				}
-				case ClosurePrimitive::BSSRDF:
 				case ClosurePrimitive::Holdout:
+					sd->osl_closure.holdout_weight += weight;
+					sd->flag |= SD_HOLDOUT;
+					break;
+				case ClosurePrimitive::BSSRDF:
 				case ClosurePrimitive::Debug:
 					break; /* not implemented */
 				case ClosurePrimitive::Background:
@@ -212,6 +217,7 @@ void OSLShader::eval_surface(KernelGlobals *kg, ShaderData *sd, float randb, int
 	sd->osl_closure.emissive_sample_sum = 0.0f;
 	sd->osl_closure.num_bsdf = 0;
 	sd->osl_closure.num_emissive = 0;
+	sd->osl_closure.holdout_weight = make_float3(0.0f, 0.0f, 0.0f);
 	sd->osl_closure.randb = randb;
 
 	if(globals->Ci) {
@@ -555,5 +561,12 @@ float3 OSLShader::volume_eval_phase(const ShaderData *sd, const float3 omega_in,
 	return eval;
 }
 
+/* Holdout Closure */
+
+float3 OSLShader::holdout_eval(const ShaderData *sd)
+{
+	return sd->osl_closure.holdout_weight;
+}
+
 CCL_NAMESPACE_END
 
diff --git a/intern/cycles/kernel/osl/osl_shader.h b/intern/cycles/kernel/osl/osl_shader.h
index 2e5279c12ce..9b578b159ae 100644
--- a/intern/cycles/kernel/osl/osl_shader.h
+++ b/intern/cycles/kernel/osl/osl_shader.h
@@ -74,6 +74,7 @@ public:
 		float3 *eval, float3 *I, float *pdf);
 	static float3 volume_eval_phase(const ShaderData *sd, const float3 omega_in,
 		const float3 omega_out);
+	static float3 holdout_eval(const ShaderData *sd);
 
 	/* release */
 	static void release(KernelGlobals *kg, const ShaderData *sd);
diff --git a/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h b/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h
index 04e4ccb8313..a04f4e2b076 100644
--- a/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h
+++ b/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h
@@ -50,7 +50,7 @@ __device void bsdf_ashikhmin_velvet_setup(ShaderData *sd, float3 N, float sigma)
 	self->m_invsigma2 = 1.0f/(sigma * sigma);
 
 	sd->svm_closure = CLOSURE_BSDF_ASHIKHMIN_VELVET_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL;
 }
 
 __device void bsdf_ashikhmin_velvet_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_diffuse.h b/intern/cycles/kernel/svm/bsdf_diffuse.h
index 00493e72203..dcd29534109 100644
--- a/intern/cycles/kernel/svm/bsdf_diffuse.h
+++ b/intern/cycles/kernel/svm/bsdf_diffuse.h
@@ -47,7 +47,7 @@ __device void bsdf_diffuse_setup(ShaderData *sd, float3 N)
 	//self->m_N = N;
 
 	sd->svm_closure = CLOSURE_BSDF_DIFFUSE_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL;
 }
 
 __device void bsdf_diffuse_blur(ShaderData *sd, float roughness)
@@ -110,7 +110,7 @@ __device void bsdf_translucent_setup(ShaderData *sd, float3 N)
 	//self->m_N = N;
 
 	sd->svm_closure = CLOSURE_BSDF_TRANSLUCENT_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL;
 }
 
 __device void bsdf_translucent_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_microfacet.h b/intern/cycles/kernel/svm/bsdf_microfacet.h
index a948ba06871..c98092b2f8f 100644
--- a/intern/cycles/kernel/svm/bsdf_microfacet.h
+++ b/intern/cycles/kernel/svm/bsdf_microfacet.h
@@ -58,7 +58,7 @@ __device void bsdf_microfacet_ggx_setup(ShaderData *sd, float3 N, float ag, floa
 	else
 		sd->svm_closure = CLOSURE_BSDF_MICROFACET_GGX_ID;
 
-	sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
 }
 
 __device void bsdf_microfacet_ggx_blur(ShaderData *sd, float roughness)
@@ -278,7 +278,7 @@ __device void bsdf_microfacet_beckmann_setup(ShaderData *sd, float3 N, float ab,
 	else
 		sd->svm_closure = CLOSURE_BSDF_MICROFACET_BECKMANN_ID;
 
-	sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
 }
 
 __device void bsdf_microfacet_beckmann_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_reflection.h b/intern/cycles/kernel/svm/bsdf_reflection.h
index 287cc9c2506..2c49eb7ef24 100644
--- a/intern/cycles/kernel/svm/bsdf_reflection.h
+++ b/intern/cycles/kernel/svm/bsdf_reflection.h
@@ -47,6 +47,7 @@ __device void bsdf_reflection_setup(ShaderData *sd, float3 N)
 	//self->m_N = N;
 
 	sd->svm_closure = CLOSURE_BSDF_REFLECTION_ID;
+	sd->flag |= SD_BSDF;
 }
 
 __device void bsdf_reflection_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_refraction.h b/intern/cycles/kernel/svm/bsdf_refraction.h
index 55a914f8334..534945f4f0b 100644
--- a/intern/cycles/kernel/svm/bsdf_refraction.h
+++ b/intern/cycles/kernel/svm/bsdf_refraction.h
@@ -48,6 +48,7 @@ __device void bsdf_refraction_setup(ShaderData *sd, float3 N, float eta)
 	self->m_eta = eta;
 
 	sd->svm_closure = CLOSURE_BSDF_REFRACTION_ID;
+	sd->flag |= SD_BSDF;
 }
 
 __device void bsdf_refraction_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_transparent.h b/intern/cycles/kernel/svm/bsdf_transparent.h
index e689e3db357..30288bf251d 100644
--- a/intern/cycles/kernel/svm/bsdf_transparent.h
+++ b/intern/cycles/kernel/svm/bsdf_transparent.h
@@ -38,6 +38,7 @@ CCL_NAMESPACE_BEGIN
 __device void bsdf_transparent_setup(ShaderData *sd)
 {
 	sd->svm_closure = CLOSURE_BSDF_TRANSPARENT_ID;
+	sd->flag |= SD_BSDF;
 }
 
 __device void bsdf_transparent_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_ward.h b/intern/cycles/kernel/svm/bsdf_ward.h
index c54418afa77..6680644a1a7 100644
--- a/intern/cycles/kernel/svm/bsdf_ward.h
+++ b/intern/cycles/kernel/svm/bsdf_ward.h
@@ -54,7 +54,7 @@ __device void bsdf_ward_setup(ShaderData *sd, float3 N, float3 T, float ax, floa
 	self->m_ay = clamp(ay, 1e-5f, 1.0f);
 
 	sd->svm_closure = CLOSURE_BSDF_WARD_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
 }
 
 __device void bsdf_ward_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_westin.h b/intern/cycles/kernel/svm/bsdf_westin.h
index 6031012d0ca..55fc8286529 100644
--- a/intern/cycles/kernel/svm/bsdf_westin.h
+++ b/intern/cycles/kernel/svm/bsdf_westin.h
@@ -51,7 +51,7 @@ __device void bsdf_westin_backscatter_setup(ShaderData *sd, float3 N, float roug
 	self->m_invroughness = 1.0f/roughness;
 
 	sd->svm_closure = CLOSURE_BSDF_WESTIN_BACKSCATTER_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
 }
 
 __device void bsdf_westin_backscatter_blur(ShaderData *sd, float roughness)
@@ -146,7 +146,7 @@ __device void bsdf_westin_sheen_setup(ShaderData *sd, float3 N, float edginess)
 	self->m_edginess = edginess;
 
 	sd->svm_closure = CLOSURE_BSDF_WESTIN_SHEEN_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
 }
 
 __device void bsdf_westin_sheen_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h
index 68e54d9752e..3a2b6bfb1ae 100644
--- a/intern/cycles/kernel/svm/svm.h
+++ b/intern/cycles/kernel/svm/svm.h
@@ -148,7 +148,7 @@ CCL_NAMESPACE_BEGIN
 
 /* Main Interpreter Loop */
 
-__device void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type, float randb, int path_flag)
+__device_noinline void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type, float randb, int path_flag)
 {
 	float stack[SVM_STACK_SIZE];
 	float closure_weight = 1.0f;
@@ -172,6 +172,8 @@ __device void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type,
 			svm_node_closure_emission(sd);
 		else if(node.x == NODE_CLOSURE_BACKGROUND)
 			svm_node_closure_background(sd);
+		else if(node.x == NODE_CLOSURE_HOLDOUT)
+			svm_node_closure_holdout(sd);
 		else if(node.x == NODE_CLOSURE_SET_WEIGHT)
 			svm_node_closure_set_weight(sd, node.y, node.z, node.w);
 		else if(node.x == NODE_CLOSURE_WEIGHT)
diff --git a/intern/cycles/kernel/svm/svm_closure.h b/intern/cycles/kernel/svm/svm_closure.h
index a7462e2c358..1efadcefbd5 100644
--- a/intern/cycles/kernel/svm/svm_closure.h
+++ b/intern/cycles/kernel/svm/svm_closure.h
@@ -136,6 +136,12 @@ __device void svm_node_closure_background(ShaderData *sd)
 	sd->svm_closure = CLOSURE_BACKGROUND_ID;
 }
 
+__device void svm_node_closure_holdout(ShaderData *sd)
+{
+	sd->svm_closure = CLOSURE_HOLDOUT_ID;
+	sd->flag |= SD_HOLDOUT;
+}
+
 /* Closure Nodes */
 
 __device void svm_node_closure_set_weight(ShaderData *sd, uint r, uint g, uint b)
diff --git a/intern/cycles/kernel/svm/svm_convert.h b/intern/cycles/kernel/svm/svm_convert.h
index 14925569bfb..188b0489d9e 100644
--- a/intern/cycles/kernel/svm/svm_convert.h
+++ b/intern/cycles/kernel/svm/svm_convert.h
@@ -30,7 +30,7 @@ __device void svm_node_convert(ShaderData *sd, float *stack, uint type, uint fro
 		}
 		case NODE_CONVERT_CF: {
 			float3 f = stack_load_float3(stack, from);
-			float g = f.x*0.2126f + f.y*0.7152f + f.z*0.0722f;
+			float g = linear_rgb_to_gray(f);
 			stack_store_float(stack, to, g);
 			break;
 		}
diff --git a/intern/cycles/kernel/svm/svm_types.h b/intern/cycles/kernel/svm/svm_types.h
index 5cbd05e4400..786478c0c03 100644
--- a/intern/cycles/kernel/svm/svm_types.h
+++ b/intern/cycles/kernel/svm/svm_types.h
@@ -81,7 +81,8 @@ typedef enum NodeType {
 	NODE_EMISSION_SET_WEIGHT_TOTAL = 4300,
 	NODE_ATTR_BUMP_DX = 4400,
 	NODE_ATTR_BUMP_DY = 4500,
-	NODE_TEX_ENVIRONMENT = 4600
+	NODE_TEX_ENVIRONMENT = 4600,
+	NODE_CLOSURE_HOLDOUT = 4700
 } NodeType;
 
 typedef enum NodeAttributeType {