Merged revision(s) 59035-59107 from trunk/blender into soc-2013-dingto.

Skipping 59034 (Non-Progressive on GPU). 

24 files changed, 67 insertions, 109 deletions

diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index 1d54915dd35..64fda3ef535 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -353,7 +353,9 @@ class CyclesRender_PT_layer_passes(CyclesButtonsPanel, Panel):
         col.prop(rl, "use_pass_uv")
         col.prop(rl, "use_pass_object_index")
         col.prop(rl, "use_pass_material_index")
+        col.separator()
         col.prop(rl, "use_pass_shadow")
+        col.prop(rl, "use_pass_ambient_occlusion")
 
         col = split.column()
         col.label(text="Diffuse:")
@@ -380,7 +382,6 @@ class CyclesRender_PT_layer_passes(CyclesButtonsPanel, Panel):
         col.separator()
         col.prop(rl, "use_pass_emit", text="Emission")
         col.prop(rl, "use_pass_environment")
-        col.prop(rl, "use_pass_ambient_occlusion")
 
 
 class Cycles_PT_post_processing(CyclesButtonsPanel, Panel):
diff --git a/intern/cycles/blender/blender_object.cpp b/intern/cycles/blender/blender_object.cpp
index a5a2f1529a2..3410c73fbde 100644
--- a/intern/cycles/blender/blender_object.cpp
+++ b/intern/cycles/blender/blender_object.cpp
@@ -373,7 +373,7 @@ static bool object_render_hide(BL::Object b_ob, bool top_level, bool parent_hide
 	
 	if(show_emitter) {
 		hide_triangles = false;
-		return (hide_as_dupli_parent || hide_as_dupli_child_original);
+		return false;
 	}
 	else if(hair_present) {
 		hide_triangles = true;
diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp
index 722d0380c22..9cf0879742e 100644
--- a/intern/cycles/device/device_cuda.cpp
+++ b/intern/cycles/device/device_cuda.cpp
@@ -572,8 +572,11 @@ public:
 		/* get kernel function */
 		if(progressive)
 			cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace_progressive"))
-		else
+		else {
 			cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace_non_progressive"))
+			if(have_error())
+				return;
+		}
 		
 		/* pass in parameters */
 		int offset = 0;
diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp
index 88a8b2aa569..589c9e1e073 100644
--- a/intern/cycles/kernel/kernel.cpp
+++ b/intern/cycles/kernel/kernel.cpp
@@ -90,10 +90,12 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t
 
 void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
 {
-	if(kernel_data.integrator.progressive)
-		kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
-	else
+#ifdef __NON_PROGRESSIVE__
+	if(!kernel_data.integrator.progressive)
 		kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
+	else
+#endif
+		kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */
diff --git a/intern/cycles/kernel/kernel.cu b/intern/cycles/kernel/kernel.cu
index e3fe23d3a37..df865aa1b7e 100644
--- a/intern/cycles/kernel/kernel.cu
+++ b/intern/cycles/kernel/kernel.cu
@@ -35,6 +35,7 @@ extern "C" __global__ void kernel_cuda_path_trace_progressive(float *buffer, uin
 		kernel_path_trace_progressive(NULL, buffer, rng_state, sample, x, y, offset, stride);
 }
 
+#ifdef __NON_PROGRESSIVE__
 extern "C" __global__ void kernel_cuda_path_trace_non_progressive(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
 {
 	int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
@@ -43,6 +44,7 @@ extern "C" __global__ void kernel_cuda_path_trace_non_progressive(float *buffer,
 	if(x < sx + sw && y < sy + sh)
 		kernel_path_trace_non_progressive(NULL, buffer, rng_state, sample, x, y, offset, stride);
 }
+#endif
 
 extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float *buffer, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
 {
diff --git a/intern/cycles/kernel/kernel_bvh.h b/intern/cycles/kernel/kernel_bvh.h
index 9f6d79e13fb..c9e14645d47 100644
--- a/intern/cycles/kernel/kernel_bvh.h
+++ b/intern/cycles/kernel/kernel_bvh.h
@@ -215,12 +215,12 @@ __device_inline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Intersectio
 	float3 P, float3 idir, uint visibility, int object, int curveAddr, int segment, uint *lcg_state, float difl, float extmax)
 {
 	float epsilon = 0.0f;
-	int depth = kernel_data.curve_kernel_data.subdivisions;
+	int depth = kernel_data.curve.subdivisions;
 
 	/* curve Intersection check */
 	float3 dir = 1.0f/idir;
 	
-	int flags = kernel_data.curve_kernel_data.curveflags;
+	int flags = kernel_data.curve.curveflags;
 
 	int prim = kernel_tex_fetch(__prim_index, curveAddr);
 
@@ -525,7 +525,7 @@ __device_inline bool bvh_curve_intersect(KernelGlobals *kg, Intersection *isect,
 	float3 P, float3 idir, uint visibility, int object, int curveAddr, int segment, uint *lcg_state, float difl, float extmax)
 {
 	/* curve Intersection check */
-	int flags = kernel_data.curve_kernel_data.curveflags;
+	int flags = kernel_data.curve.curveflags;
 
 	int prim = kernel_tex_fetch(__prim_index, curveAddr);
 	float4 v00 = kernel_tex_fetch(__curves, prim);
@@ -649,7 +649,7 @@ __device_inline bool bvh_curve_intersect(KernelGlobals *kg, Intersection *isect,
 
 			if (flags & CURVE_KN_ENCLOSEFILTER) {
 
-				float enc_ratio = kernel_data.curve_kernel_data.encasing_ratio;
+				float enc_ratio = kernel_data.curve.encasing_ratio;
 				if((dot(P - p1, tg) > -r1 * enc_ratio) && (dot(P - p2, tg) < r2 * enc_ratio)) {
 					float a2 = 1.0f - (dirz*dirz*(1 + gd*gd*enc_ratio*enc_ratio));
 					float c2 = dot(dif,dif) - difz * difz * (1 + gd*gd*enc_ratio*enc_ratio) - r1*r1*enc_ratio*enc_ratio - 2*r1*difz*gd*enc_ratio;
@@ -1009,7 +1009,7 @@ __device_inline float3 curvepoint(float t, float3 p0, float3 p1, float3 p2, floa
 
 __device_inline float3 bvh_curve_refine(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray, float t)
 {
-	int flag = kernel_data.curve_kernel_data.curveflags;
+	int flag = kernel_data.curve.curveflags;
 	float3 P = ray->P;
 	float3 D = ray->D;
 
@@ -1062,7 +1062,7 @@ __device_inline float3 bvh_curve_refine(KernelGlobals *kg, ShaderData *sd, const
 		sd->v = 0.0f;
 #endif
 
-		if(kernel_data.curve_kernel_data.curveflags & CURVE_KN_RIBBONS)
+		if(kernel_data.curve.curveflags & CURVE_KN_RIBBONS)
 			sd->Ng = normalize(-(D - tg * (dot(tg,D))));
 		else {
 			sd->Ng = normalize(P - p_curr);
@@ -1080,7 +1080,7 @@ __device_inline float3 bvh_curve_refine(KernelGlobals *kg, ShaderData *sd, const
 #endif
 
 		if (flag & CURVE_KN_TRUETANGENTGNORMAL) {
-			sd->Ng = -(D - tg * (dot(tg,D) * kernel_data.curve_kernel_data.normalmix));
+			sd->Ng = -(D - tg * (dot(tg,D) * kernel_data.curve.normalmix));
 			sd->Ng = normalize(sd->Ng);
 			if (flag & CURVE_KN_NORMALCORRECTION) {
 				sd->Ng = sd->Ng - gd * tg;
@@ -1098,7 +1098,7 @@ __device_inline float3 bvh_curve_refine(KernelGlobals *kg, ShaderData *sd, const
 		sd->N = sd->Ng;
 
 		if (flag & CURVE_KN_TANGENTGNORMAL && !(flag & CURVE_KN_TRUETANGENTGNORMAL)) {
-			sd->N = -(D - tg * (dot(tg,D) * kernel_data.curve_kernel_data.normalmix));
+			sd->N = -(D - tg * (dot(tg,D) * kernel_data.curve.normalmix));
 			sd->N = normalize(sd->N);
 			if (flag & CURVE_KN_NORMALCORRECTION) {
 				sd->N = sd->N - gd * tg;
diff --git a/intern/cycles/kernel/kernel_bvh_traversal.h b/intern/cycles/kernel/kernel_bvh_traversal.h
index 8d423b7e7ce..cfca405e7a5 100644
--- a/intern/cycles/kernel/kernel_bvh_traversal.h
+++ b/intern/cycles/kernel/kernel_bvh_traversal.h
@@ -258,7 +258,7 @@ __device bool BVH_FUNCTION_NAME
 #if !FEATURE(BVH_SUBSURFACE)
 							if(segment != ~0) {
 
-								if(kernel_data.curve_kernel_data.curveflags & CURVE_KN_INTERPOLATE) 
+								if(kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) 
 #if FEATURE(BVH_HAIR_MINIMUM_WIDTH)
 									hit = bvh_cardinal_curve_intersect(kg, isect, P, idir, visibility, object, primAddr, segment, lcg_state, difl, extmax);
 								else
diff --git a/intern/cycles/kernel/kernel_curve.h b/intern/cycles/kernel/kernel_curve.h
index e065717888c..b9637e7df8b 100644
--- a/intern/cycles/kernel/kernel_curve.h
+++ b/intern/cycles/kernel/kernel_curve.h
@@ -120,14 +120,14 @@ __device float3 curve_tangent_normal(KernelGlobals *kg, ShaderData *sd)
 	float3 tgN = make_float3(0.0f,0.0f,0.0f);
 
 	if(sd->segment != ~0) {
-		float normalmix = kernel_data.curve_kernel_data.normalmix;
+		float normalmix = kernel_data.curve.normalmix;
 
 		tgN = -(-sd->I - sd->dPdu * (dot(sd->dPdu,-sd->I) * normalmix / len_squared(sd->dPdu)));
 		tgN = normalize(tgN);
 
 		/* need to find suitable scaled gd for corrected normal */
 #if 0
-		if (kernel_data.curve_kernel_data.use_tangent_normal_correction)
+		if (kernel_data.curve.use_tangent_normal_correction)
 			tgN = normalize(tgN - gd * sd->dPdu);
 #endif
 	}
diff --git a/intern/cycles/kernel/kernel_passes.h b/intern/cycles/kernel/kernel_passes.h
index 0b45e243283..4b8d6f84faa 100644
--- a/intern/cycles/kernel/kernel_passes.h
+++ b/intern/cycles/kernel/kernel_passes.h
@@ -48,7 +48,7 @@ __device_inline void kernel_write_data_passes(KernelGlobals *kg, __global float
 	if(!(flag & PASS_ALL))
 		return;
 	
-	/* todo: add alpha treshold */
+	/* todo: add alpha threshold */
 	if(!(path_flag & PATH_RAY_TRANSPARENT)) {
 		if(sample == 0) {
 			if(flag & PASS_DEPTH) {
diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h
index 48df60162b1..bf06f8dd5f6 100644
--- a/intern/cycles/kernel/kernel_path.h
+++ b/intern/cycles/kernel/kernel_path.h
@@ -271,10 +271,10 @@ __device float4 kernel_path_progressive(KernelGlobals *kg, RNG *rng, int sample,
 			if((kernel_data.cam.resolution == 1) && (state.flag & PATH_RAY_CAMERA)) {	
 				float3 pixdiff = ray.dD.dx + ray.dD.dy;
 				/*pixdiff = pixdiff - dot(pixdiff, ray.D)*ray.D;*/
-				difl = kernel_data.curve_kernel_data.minimum_width * len(pixdiff) * 0.5f;
+				difl = kernel_data.curve.minimum_width * len(pixdiff) * 0.5f;
 			}
 
-			extmax = kernel_data.curve_kernel_data.maximum_width;
+			extmax = kernel_data.curve.maximum_width;
 			lcg_state = lcg_init(*rng + rng_offset + sample*0x51633e2d);
 		}
 
@@ -990,10 +990,10 @@ __device float4 kernel_path_non_progressive(KernelGlobals *kg, RNG *rng, int sam
 			if((kernel_data.cam.resolution == 1) && (state.flag & PATH_RAY_CAMERA)) {	
 				float3 pixdiff = ray.dD.dx + ray.dD.dy;
 				/*pixdiff = pixdiff - dot(pixdiff, ray.D)*ray.D;*/
-				difl = kernel_data.curve_kernel_data.minimum_width * len(pixdiff) * 0.5f;
+				difl = kernel_data.curve.minimum_width * len(pixdiff) * 0.5f;
 			}
 
-			extmax = kernel_data.curve_kernel_data.maximum_width;
+			extmax = kernel_data.curve.maximum_width;
 			lcg_state = lcg_init(*rng + rng_offset + sample*0x51633e2d);
 		}
 
@@ -1194,6 +1194,7 @@ __device void kernel_path_trace_progressive(KernelGlobals *kg,
 	path_rng_end(kg, rng_state, rng);
 }
 
+#ifdef __NON_PROGRESSIVE__
 __device void kernel_path_trace_non_progressive(KernelGlobals *kg,
 	__global float *buffer, __global uint *rng_state,
 	int sample, int x, int y, int offset, int stride)
@@ -1215,11 +1216,7 @@ __device void kernel_path_trace_non_progressive(KernelGlobals *kg,
 	float4 L;
 
 	if (ray.t != 0.0f)
-#ifdef __NON_PROGRESSIVE__
 		L = kernel_path_non_progressive(kg, &rng, sample, ray, buffer);
-#else
-		L = kernel_path_progressive(kg, &rng, sample, ray, buffer);
-#endif
 	else
 		L = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 
@@ -1228,6 +1225,7 @@ __device void kernel_path_trace_non_progressive(KernelGlobals *kg,
 
 	path_rng_end(kg, rng_state, rng);
 }
+#endif
 
 CCL_NAMESPACE_END
 
diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h
index 5fb349ff1ee..d8ff662c505 100644
--- a/intern/cycles/kernel/kernel_shader.h
+++ b/intern/cycles/kernel/kernel_shader.h
@@ -91,7 +91,7 @@ void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd,
 		sd->segment = isect->segment;
 
 		float tcorr = isect->t;
-		if(kernel_data.curve_kernel_data.curveflags & CURVE_KN_POSTINTERSECTCORRECTION) {
+		if(kernel_data.curve.curveflags & CURVE_KN_POSTINTERSECTCORRECTION) {
 			tcorr = (isect->u < 0)? tcorr + sqrtf(isect->v) : tcorr - sqrtf(isect->v);
 			sd->ray_length = tcorr;
 		}
@@ -193,7 +193,7 @@ __device_inline void shader_setup_from_subsurface(KernelGlobals *kg, ShaderData
 		sd->segment = isect->segment;
 
 		float tcorr = isect->t;
-		if(kernel_data.curve_kernel_data.curveflags & CURVE_KN_POSTINTERSECTCORRECTION)
+		if(kernel_data.curve.curveflags & CURVE_KN_POSTINTERSECTCORRECTION)
 			tcorr = (isect->u < 0)? tcorr + sqrtf(isect->v) : tcorr - sqrtf(isect->v);
 
 		sd->P = bvh_curve_refine(kg, sd, isect, ray, tcorr);
diff --git a/intern/cycles/kernel/kernel_sse2.cpp b/intern/cycles/kernel/kernel_sse2.cpp
index 8321c679f07..344d530396b 100644
--- a/intern/cycles/kernel/kernel_sse2.cpp
+++ b/intern/cycles/kernel/kernel_sse2.cpp
@@ -39,10 +39,12 @@ CCL_NAMESPACE_BEGIN
 
 void kernel_cpu_sse2_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
 {
-	if(kernel_data.integrator.progressive)
-		kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
-	else
+#ifdef __NON_PROGRESSIVE__
+	if(!kernel_data.integrator.progressive)
 		kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
+	else
+#endif
+		kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */
diff --git a/intern/cycles/kernel/kernel_sse3.cpp b/intern/cycles/kernel/kernel_sse3.cpp
index 35a816f7da9..a25c0bca93a 100644
--- a/intern/cycles/kernel/kernel_sse3.cpp
+++ b/intern/cycles/kernel/kernel_sse3.cpp
@@ -41,10 +41,12 @@ CCL_NAMESPACE_BEGIN
 
 void kernel_cpu_sse3_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
 {
-	if(kernel_data.integrator.progressive)
-		kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
-	else
+#ifdef __NON_PROGRESSIVE__
+	if(!kernel_data.integrator.progressive)
 		kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
+	else
+#endif
+		kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index 0608dad5c05..836eacf7cb6 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -222,7 +222,7 @@ enum PathRayFlag {
 	PATH_RAY_CURVE = 1024,
 
 	/* this gives collisions with localview bits
-	 * see: CYCLES_LOCAL_LAYER_HACK(), grr - Campbell */
+	 * see: blender_util.h, grr - Campbell */
 	PATH_RAY_LAYER_SHIFT = (32-20)
 };
 
@@ -279,8 +279,6 @@ typedef enum PassType {
 
 #ifdef __PASSES__
 
-typedef float3 PathThroughput;
-
 typedef struct PathRadiance {
 	int use_light_pass;
 
@@ -328,7 +326,6 @@ typedef struct BsdfEval {
 
 #else
 
-typedef float3 PathThroughput;
 typedef float3 PathRadiance;
 typedef float3 BsdfEval;
 
@@ -600,7 +597,7 @@ typedef struct ShaderData {
 #endif
 } ShaderData;
 
-/* Constrant Kernel Data
+/* Constant Kernel Data
  *
  * These structs are passed from CPU to various devices, and the struct layout
  * must match exactly. Structs are padded to ensure 16 byte alignment, and we
@@ -774,8 +771,10 @@ typedef struct KernelIntegrator {
 	int transmission_samples;
 	int ao_samples;
 	int mesh_light_samples;
-	int use_lamp_mis;
 	int subsurface_samples;
+	
+	/* mis */
+	int use_lamp_mis;
 
 	/* sampler */
 	int sampling_pattern;
@@ -821,7 +820,6 @@ typedef struct KernelCurves {
 	float maximum_width;
 	float curve_epsilon;
 	int pad1;
-
 } KernelCurves;
 
 typedef struct KernelBSSRDF {
@@ -843,7 +841,7 @@ typedef struct KernelData {
 	KernelSunSky sunsky;
 	KernelIntegrator integrator;
 	KernelBVH bvh;
-	KernelCurves curve_kernel_data;
+	KernelCurves curve;
 	KernelBSSRDF bssrdf;
 	KernelBlackbody blackbody;
 } KernelData;
diff --git a/intern/cycles/kernel/shaders/node_texture.h b/intern/cycles/kernel/shaders/node_texture.h
index 645be23c6a5..3997803ec08 100644
--- a/intern/cycles/kernel/shaders/node_texture.h
+++ b/intern/cycles/kernel/shaders/node_texture.h
@@ -18,8 +18,11 @@
 
 float voronoi_distance(string distance_metric, vector d, float e)
 {
+#if 0
 	if (distance_metric == "Distance Squared")
+#endif
 		return dot(d, d);
+#if 0
 	if (distance_metric == "Actual Distance")
 		return length(d);
 	if (distance_metric == "Manhattan")
@@ -34,6 +37,7 @@ float voronoi_distance(string distance_metric, vector d, float e)
 		return pow(pow(fabs(d[0]), e) + pow(fabs(d[1]), e) + pow(fabs(d[2]), e), 1.0 / e);
 	
 	return 0.0;
+#endif
 }
 
 /* Voronoi / Worley like */
@@ -198,31 +202,6 @@ float noise_basis_hard(point p, string basis, int hard)
 	return (hard) ? fabs(2.0 * t - 1.0) : t;
 }
 
-/* Waves */
-
-float noise_wave(string wave, float a)
-{
-	if (wave == "Sine") {
-		return 0.5 + 0.5 * sin(a);
-	}
-	if (wave == "Saw") {
-		float b = 2 * M_PI;
-		int n = (int)(a / b);
-		a -= n * b;
-		if (a < 0) a += b;
-
-		return a / b;
-	}
-	if (wave == "Tri") {
-		float b = 2 * M_PI;
-		float rmax = 1.0;
-
-		return rmax - 2.0 * fabs(floor((a * (1.0 / b)) + 0.5) - (a * (1.0 / b)));
-	}
-
-	return 0.0;
-}
-
 /* Turbulence */
 
 float noise_turbulence(point p, string basis, float details, int hard)
diff --git a/intern/cycles/kernel/shaders/node_wave_texture.osl b/intern/cycles/kernel/shaders/node_wave_texture.osl
index 12b837bc448..93d45ded9c3 100644
--- a/intern/cycles/kernel/shaders/node_wave_texture.osl
+++ b/intern/cycles/kernel/shaders/node_wave_texture.osl
@@ -39,7 +39,7 @@ float wave(point p, float scale, string type, float detail, float distortion, fl
 	if (distortion != 0.0) {
 		n = n + (distortion * noise_turbulence(p * dscale, "Perlin", detail, 0));
 	}
-	return noise_wave("Sine", n);
+	return 0.5 + 0.5 * sin(n);
 }
 
 shader node_wave_texture(
diff --git a/intern/cycles/kernel/svm/svm_texture.h b/intern/cycles/kernel/svm/svm_texture.h
index 3347b623506..02fb3e5b86e 100644
--- a/intern/cycles/kernel/svm/svm_texture.h
+++ b/intern/cycles/kernel/svm/svm_texture.h
@@ -22,8 +22,11 @@ CCL_NAMESPACE_BEGIN
 
 __device float voronoi_distance(NodeDistanceMetric distance_metric, float3 d, float e)
 {
+#if 0
 	if(distance_metric == NODE_VORONOI_DISTANCE_SQUARED)
+#endif
 		return dot(d, d);
+#if 0
 	if(distance_metric == NODE_VORONOI_ACTUAL_DISTANCE)
 		return len(d);
 	if(distance_metric == NODE_VORONOI_MANHATTAN)
@@ -38,6 +41,7 @@ __device float voronoi_distance(NodeDistanceMetric distance_metric, float3 d, fl
 		return powf(powf(fabsf(d.x), e) + powf(fabsf(d.y), e) + powf(fabsf(d.z), e), 1.0f/e);
 	
 	return 0.0f;
+#endif
 }
 
 /* Voronoi / Worley like */
@@ -177,31 +181,6 @@ __device float noise_basis_hard(float3 p, NodeNoiseBasis basis, int hard)
 	return (hard)? fabsf(2.0f*t - 1.0f): t;
 }
 
-/* Waves */
-
-__device float noise_wave(NodeWaveBasis wave, float a)
-{
-	if(wave == NODE_WAVE_SINE) {
-		return 0.5f + 0.5f * sinf(a);
-	}
-	else if(wave == NODE_WAVE_SAW) {
-		float b = M_2PI_F;
-		int n = float_to_int(a / b);
-		a -= n*b;
-		if(a < 0.0f) a += b;
-
-		return a / b;
-	}
-	else if(wave == NODE_WAVE_TRI) {
-		float b = M_2PI_F;
-		float rmax = 1.0f;
-
-		return rmax - 2.0f*fabsf(floorf((a*(1.0f/b))+0.5f) - (a*(1.0f/b)));
-	}
-
-	return 0.0f;
-}
-
 /* Turbulence */
 
 __device_noinline float noise_turbulence(float3 p, NodeNoiseBasis basis, float octaves, int hard)
diff --git a/intern/cycles/kernel/svm/svm_types.h b/intern/cycles/kernel/svm/svm_types.h
index dd95f19f175..939decf80a9 100644
--- a/intern/cycles/kernel/svm/svm_types.h
+++ b/intern/cycles/kernel/svm/svm_types.h
@@ -274,12 +274,6 @@ typedef enum NodeNoiseBasis {
 	NODE_NOISE_CELL_NOISE
 } NodeNoiseBasis;
 
-typedef enum NodeWaveBasis {
-	NODE_WAVE_SINE,
-	NODE_WAVE_SAW,
-	NODE_WAVE_TRI
-} NodeWaveBasis;
-
 typedef enum NodeMusgraveType {
 	NODE_MUSGRAVE_MULTIFRACTAL,
 	NODE_MUSGRAVE_FBM,
diff --git a/intern/cycles/kernel/svm/svm_wave.h b/intern/cycles/kernel/svm/svm_wave.h
index d2d808e4063..33c1694283d 100644
--- a/intern/cycles/kernel/svm/svm_wave.h
+++ b/intern/cycles/kernel/svm/svm_wave.h
@@ -22,21 +22,19 @@ CCL_NAMESPACE_BEGIN
 
 __device_noinline float svm_wave(NodeWaveType type, float3 p, float scale, float detail, float distortion, float dscale)
 {
-	float w, n;
+	float n;
 
 	p *= scale;
 
 	if(type == NODE_WAVE_BANDS)
 		n = (p.x + p.y + p.z) * 10.0f;
-	else /* if(type == NODE_WAVE_RINGS) */
+	else /* NODE_WAVE_RINGS */
 		n = len(p) * 20.0f;
 	
 	if(distortion != 0.0f)
 		n += distortion * noise_turbulence(p*dscale, NODE_NOISE_PERLIN, detail, 0);
 
-	w = noise_wave(NODE_WAVE_SINE, n);
-
-	return w;
+	return 0.5f + 0.5f * sinf(n);
 }
 
 __device void svm_node_tex_wave(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset)
diff --git a/intern/cycles/render/curves.cpp b/intern/cycles/render/curves.cpp
index 3fcd78b0152..502bed06930 100644
--- a/intern/cycles/render/curves.cpp
+++ b/intern/cycles/render/curves.cpp
@@ -119,7 +119,7 @@ void CurveSystemManager::device_update(Device *device, DeviceScene *dscene, Scen
 
 	progress.set_status("Updating Hair settings", "Copying Hair settings to device");
 
-	KernelCurves *kcurve= &dscene->data.curve_kernel_data;
+	KernelCurves *kcurve= &dscene->data.curve;
 
 	kcurve->curveflags = 0;
 
diff --git a/intern/cycles/render/graph.cpp b/intern/cycles/render/graph.cpp
index 2e8bc77b9c9..4c6e64a1fb5 100644
--- a/intern/cycles/render/graph.cpp
+++ b/intern/cycles/render/graph.cpp
@@ -255,7 +255,7 @@ void ShaderGraph::finalize(bool do_bump, bool do_osl, bool do_multi_transform)
 
 void ShaderGraph::find_dependencies(set<ShaderNode*>& dependencies, ShaderInput *input)
 {
-	/* find all nodes that this input dependes on directly and indirectly */
+	/* find all nodes that this input depends on directly and indirectly */
 	ShaderNode *node = (input->link)? input->link->parent: NULL;
 
 	if(node) {
diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp
index 13c199f879b..87f5722d144 100644
--- a/intern/cycles/render/session.cpp
+++ b/intern/cycles/render/session.cpp
@@ -135,7 +135,7 @@ void Session::reset_gpu(BufferParams& buffer_params, int samples)
 {
 	thread_scoped_lock pause_lock(pause_mutex);
 
-	/* block for buffer acces and reset immediately. we can't do this
+	/* block for buffer access and reset immediately. we can't do this
 	 * in the thread, because we need to allocate an OpenGL buffer, and
 	 * that only works in the main thread */
 	thread_scoped_lock display_lock(display_mutex);
diff --git a/intern/cycles/render/svm.cpp b/intern/cycles/render/svm.cpp
index 9e887aeadf2..4e617155465 100644
--- a/intern/cycles/render/svm.cpp
+++ b/intern/cycles/render/svm.cpp
@@ -504,7 +504,7 @@ void SVMCompiler::generate_closure(ShaderNode *node, set<ShaderNode*>& done)
 
 void SVMCompiler::generate_multi_closure(ShaderNode *node, set<ShaderNode*>& done, set<ShaderNode*>& closure_done)
 {
-	/* todo: the weaks point here is that unlike the single closure sampling 
+	/* todo: the weak point here is that unlike the single closure sampling 
 	 * we will evaluate all nodes even if they are used as input for closures
 	 * that are unused. it's not clear what would be the best way to skip such
 	 * nodes at runtime, especially if they are tangled up  */
diff --git a/intern/cycles/render/tile.cpp b/intern/cycles/render/tile.cpp
index 55bbfeb3f1b..ea938943f93 100644
--- a/intern/cycles/render/tile.cpp
+++ b/intern/cycles/render/tile.cpp
@@ -189,7 +189,7 @@ list<Tile>::iterator TileManager::next_background_tile(int device, int tile_orde
 
 	int64_t cordx = max(1, params.width/resolution);
 	int64_t cordy = max(1, params.height/resolution);
-	int64_t mindist = cordx * cordy;
+	int64_t mindist = INT_MAX;
 	
 	int64_t centx = cordx / 2, centy = cordy / 2;