Merge branch 'master' into blender2.8

intern/cycles/device/device_cuda.cpp

@@ -993,7 +993,7 @@ public:
 				cuda_assert(cuCtxSynchronize());
 
 				if(task.get_cancel()) {
-					canceled = false;
+					canceled = true;
 					break;
 				}
 			}

intern/cycles/device/device_opencl.cpp

@@ -2346,7 +2346,9 @@ public:
 		}
 	}
 
-	void path_trace(SplitRenderTile& rtile, int2 max_render_feasible_tile_size)
+	void path_trace(DeviceTask *task,
+	                SplitRenderTile& rtile,
+	                int2 max_render_feasible_tile_size)
 	{
 		/* cast arguments to cl types */
 		cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
@@ -2739,6 +2741,7 @@ public:
 		/* Record number of time host intervention has been made */
 		unsigned int numHostIntervention = 0;
 		unsigned int numNextPathIterTimes = PathIteration_times;
+		bool canceled = false;
 		while(activeRaysAvailable) {
 			/* Twice the global work size of other kernels for
 			 * ckPathTraceKernel_shadow_blocked_direct_lighting. */
@@ -2757,6 +2760,10 @@ public:
 				ENQUEUE_SPLIT_KERNEL(direct_lighting, global_size, local_size);
 				ENQUEUE_SPLIT_KERNEL(shadow_blocked, global_size_shadow_blocked, local_size);
 				ENQUEUE_SPLIT_KERNEL(next_iteration_setup, global_size, local_size);
+				if(task->get_cancel()) {
+					canceled = true;
+					break;
+				}
 			}
 
 			/* Read ray-state into Host memory to decide if we should exit
@@ -2794,22 +2801,28 @@ public:
 				 */
 				numNextPathIterTimes += PATH_ITER_INC_FACTOR;
 			}
+			if(task->get_cancel()) {
+				canceled = true;
+				break;
+			}
 		}
 
 		/* Execute SumALLRadiance kernel to accumulate radiance calculated in
 		 * per_sample_output_buffers into RenderTile's output buffer.
 		 */
-		size_t sum_all_radiance_local_size[2] = {16, 16};
-		size_t sum_all_radiance_global_size[2];
-		sum_all_radiance_global_size[0] =
-			(((d_w - 1) / sum_all_radiance_local_size[0]) + 1) *
-			sum_all_radiance_local_size[0];
-		sum_all_radiance_global_size[1] =
-			(((d_h - 1) / sum_all_radiance_local_size[1]) + 1) *
-			sum_all_radiance_local_size[1];
-		ENQUEUE_SPLIT_KERNEL(sum_all_radiance,
-		                     sum_all_radiance_global_size,
-		                     sum_all_radiance_local_size);
+		if (!canceled) {
+			size_t sum_all_radiance_local_size[2] = {16, 16};
+			size_t sum_all_radiance_global_size[2];
+			sum_all_radiance_global_size[0] =
+				(((d_w - 1) / sum_all_radiance_local_size[0]) + 1) *
+				sum_all_radiance_local_size[0];
+			sum_all_radiance_global_size[1] =
+				(((d_h - 1) / sum_all_radiance_local_size[1]) + 1) *
+				sum_all_radiance_local_size[1];
+			ENQUEUE_SPLIT_KERNEL(sum_all_radiance,
+			                     sum_all_radiance_global_size,
+			                     sum_all_radiance_local_size);
+		}
 
 #undef ENQUEUE_SPLIT_KERNEL
 #undef GLUE
@@ -3182,7 +3195,8 @@ public:
 					    tile_iter < to_path_trace_render_tiles.size();
 					    ++tile_iter)
 					{
-						path_trace(to_path_trace_render_tiles[tile_iter],
+						path_trace(task,
+						           to_path_trace_render_tiles[tile_iter],
 						           max_render_feasible_tile_size);
 					}
 				}
@@ -3198,7 +3212,7 @@ public:
 					/* buffer_rng_state_stride is stride itself. */
 					SplitRenderTile split_tile(tile);
 					split_tile.buffer_rng_state_stride = tile.stride;
-					path_trace(split_tile, max_render_feasible_tile_size);
+					path_trace(task, split_tile, max_render_feasible_tile_size);
 				}
 				tile.sample = tile.start_sample + tile.num_samples;
 

intern/cycles/kernel/bvh/bvh.h

@@ -1,6 +1,5 @@
 /*
- * Adapted from code Copyright 2009-2010 NVIDIA Corporation
- * Modifications Copyright 2011, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

intern/cycles/kernel/bvh/bvh_types.h

@@ -1,6 +1,5 @@
 /*
- * Adapted from code Copyright 2009-2010 NVIDIA Corporation
- * Modifications Copyright 2011, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

intern/cycles/kernel/bvh/qbvh_nodes.h

@@ -12,6 +12,8 @@
  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  * See the License for the specific language governing permissions and
  * limitations under the License.
+ *
+ * Aligned nodes intersection SSE code is adopted from Embree,
  */
 
 struct QBVHStackItem {

intern/cycles/kernel/bvh/qbvh_shadow_all.h

@@ -1,8 +1,5 @@
 /*
- * Adapted from code Copyright 2009-2010 NVIDIA Corporation,
- * and code copyright 2009-2012 Intel Corporation
- *
- * Modifications Copyright 2011-2014, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

intern/cycles/kernel/bvh/qbvh_subsurface.h

@@ -1,8 +1,5 @@
 /*
- * Adapted from code Copyright 2009-2010 NVIDIA Corporation,
- * and code copyright 2009-2012 Intel Corporation
- *
- * Modifications Copyright 2011-2014, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

intern/cycles/kernel/bvh/qbvh_traversal.h

@@ -1,8 +1,5 @@
 /*
- * Adapted from code Copyright 2009-2010 NVIDIA Corporation,
- * and code copyright 2009-2012 Intel Corporation
- *
- * Modifications Copyright 2011-2014, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

intern/cycles/kernel/bvh/qbvh_volume.h

@@ -1,8 +1,5 @@
 /*
- * Adapted from code Copyright 2009-2010 NVIDIA Corporation,
- * and code copyright 2009-2012 Intel Corporation
- *
- * Modifications Copyright 2011-2014, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

intern/cycles/kernel/bvh/qbvh_volume_all.h

@@ -1,8 +1,5 @@
 /*
- * Adapted from code Copyright 2009-2010 NVIDIA Corporation,
- * and code copyright 2009-2012 Intel Corporation
- *
- * Modifications Copyright 2011-2014, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

intern/cycles/kernel/geom/geom.h

@@ -1,6 +1,5 @@
 /*
- * Adapted from code Copyright 2009-2010 NVIDIA Corporation
- * Modifications Copyright 2011, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

intern/cycles/kernel/geom/geom_motion_triangle.h

@@ -1,6 +1,5 @@
 /*
- * Adapted from code Copyright 2009-2010 NVIDIA Corporation
- * Modifications Copyright 2011, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

intern/cycles/kernel/geom/geom_triangle.h

@@ -1,6 +1,5 @@
 /*
- * Adapted from code Copyright 2009-2010 NVIDIA Corporation
- * Modifications Copyright 2011, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

intern/cycles/kernel/svm/svm_math_util.h

@@ -164,6 +164,9 @@ ccl_device float3 svm_math_blackbody_color(float t) {
 
 ccl_device_inline float3 svm_math_gamma_color(float3 color, float gamma)
 {
+	if(gamma == 0.0f)
+		return make_float3(1.0f, 1.0f, 1.0f);
+
 	if(color.x > 0.0f)
 		color.x = powf(color.x, gamma);
 	if(color.y > 0.0f)

intern/cycles/render/constant_fold.cpp

@@ -89,6 +89,19 @@ void ConstantFolder::make_zero() const
 	}
 }
 
+void ConstantFolder::make_one() const
+{
+	if(output->type() == SocketType::FLOAT) {
+		make_constant(1.0f);
+	}
+	else if(SocketType::is_float3(output->type())) {
+		make_constant(make_float3(1.0f, 1.0f, 1.0f));
+	}
+	else {
+		assert(0);
+	}
+}
+
 void ConstantFolder::bypass(ShaderOutput *new_output) const
 {
 	assert(new_output);
@@ -321,6 +334,15 @@ void ConstantFolder::fold_math(NodeMath type, bool clamp) const
 				make_zero();
 			}
 			break;
+		case NODE_MATH_POWER:
+			/* 1 ^ X == X ^ 0 == 1 */
+			if(is_one(value1_in) || is_zero(value2_in)) {
+				make_one();
+			}
+			/* X ^ 1 == X */
+			else if(is_one(value2_in)) {
+				try_bypass_or_make_constant(value1_in, clamp);
+			}
 		default:
 			break;
 	}

intern/cycles/render/constant_fold.h

@@ -43,6 +43,7 @@ public:
 	void make_constant_clamp(float value, bool clamp) const;
 	void make_constant_clamp(float3 value, bool clamp) const;
 	void make_zero() const;
+	void make_one() const;
 
 	/* Bypass node, relinking to another output socket. */
 	void bypass(ShaderOutput *output) const;

intern/cycles/render/mesh_subdivision.cpp

@@ -109,7 +109,7 @@ namespace Far {
 
 	template<>
 	void TopologyRefinerFactory<ccl::Mesh>::reportInvalidTopology(TopologyError /*err_code*/,
-		char const */*msg*/, ccl::Mesh const& /*mesh*/)
+		char const * /*msg*/, ccl::Mesh const& /*mesh*/)
 	{
 	}
 } /* namespace Far */

intern/cycles/render/nodes.cpp

@@ -3896,6 +3896,19 @@ void GammaNode::constant_fold(const ConstantFolder& folder)
 	if(folder.all_inputs_constant()) {
 		folder.make_constant(svm_math_gamma_color(color, gamma));
 	}
+	else {
+		ShaderInput *color_in = input("Color");
+		ShaderInput *gamma_in = input("Gamma");
+
+		/* 1 ^ X == X ^ 0 == 1 */
+		if(folder.is_one(color_in) || folder.is_zero(gamma_in)) {
+			folder.make_one();
+		}
+		/* X ^ 1 == X */
+		else if(folder.is_one(gamma_in)) {
+			folder.try_bypass_or_make_constant(color_in, false);
+		}
+	}
 }
 
 void GammaNode::compile(SVMCompiler& compiler)

intern/cycles/test/render_graph_finalize_test.cpp

@@ -930,6 +930,72 @@ TEST(render_graph, constant_fold_gamma)
 	graph.finalize(&scene);
 }
 
+/*
+ * Tests: Gamma with one constant 0 input.
+ */
+TEST(render_graph, constant_fold_gamma_part_0)
+{
+	DEFINE_COMMON_VARIABLES(builder, log);
+
+	EXPECT_ANY_MESSAGE(log);
+	INVALID_INFO_MESSAGE(log, "Folding Gamma_Cx::");
+	CORRECT_INFO_MESSAGE(log, "Folding Gamma_xC::Color to constant (1, 1, 1).");
+
+	builder
+		.add_attribute("Attribute")
+		/* constant on the left */
+		.add_node(ShaderNodeBuilder<GammaNode>("Gamma_Cx")
+		          .set("Color", make_float3(0.0f, 0.0f, 0.0f)))
+		.add_connection("Attribute::Fac", "Gamma_Cx::Gamma")
+		/* constant on the right */
+		.add_node(ShaderNodeBuilder<GammaNode>("Gamma_xC")
+		          .set("Gamma", 0.0f))
+		.add_connection("Attribute::Color", "Gamma_xC::Color")
+		/* output sum */
+		.add_node(ShaderNodeBuilder<MixNode>("Out")
+		          .set(&MixNode::type, NODE_MIX_ADD)
+		          .set(&MixNode::use_clamp, true)
+		          .set("Fac", 1.0f))
+		.add_connection("Gamma_Cx::Color", "Out::Color1")
+		.add_connection("Gamma_xC::Color", "Out::Color2")
+		.output_color("Out::Color");
+
+	graph.finalize(&scene);
+}
+
+/*
+ * Tests: Gamma with one constant 1 input.
+ */
+TEST(render_graph, constant_fold_gamma_part_1)
+{
+	DEFINE_COMMON_VARIABLES(builder, log);
+
+	EXPECT_ANY_MESSAGE(log);
+	CORRECT_INFO_MESSAGE(log, "Folding Gamma_Cx::Color to constant (1, 1, 1).");
+	CORRECT_INFO_MESSAGE(log, "Folding Gamma_xC::Color to socket Attribute::Color.");
+
+	builder
+		.add_attribute("Attribute")
+		/* constant on the left */
+		.add_node(ShaderNodeBuilder<GammaNode>("Gamma_Cx")
+		          .set("Color", make_float3(1.0f, 1.0f, 1.0f)))
+		.add_connection("Attribute::Fac", "Gamma_Cx::Gamma")
+		/* constant on the right */
+		.add_node(ShaderNodeBuilder<GammaNode>("Gamma_xC")
+		          .set("Gamma", 1.0f))
+		.add_connection("Attribute::Color", "Gamma_xC::Color")
+		/* output sum */
+		.add_node(ShaderNodeBuilder<MixNode>("Out")
+		          .set(&MixNode::type, NODE_MIX_ADD)
+		          .set(&MixNode::use_clamp, true)
+		          .set("Fac", 1.0f))
+		.add_connection("Gamma_Cx::Color", "Out::Color1")
+		.add_connection("Gamma_xC::Color", "Out::Color2")
+		.output_color("Out::Color");
+
+	graph.finalize(&scene);
+}
+
 /*
  * Tests: BrightnessContrast with all constant inputs.
  */
@@ -1142,6 +1208,40 @@ TEST(render_graph, constant_fold_part_math_div_0)
 	graph.finalize(&scene);
 }
 
+/*
+ * Tests: partial folding for Math Power with known 0.
+ */
+TEST(render_graph, constant_fold_part_math_pow_0)
+{
+	DEFINE_COMMON_VARIABLES(builder, log);
+
+	EXPECT_ANY_MESSAGE(log);
+	/* X ^ 0 == 1 */
+	INVALID_INFO_MESSAGE(log, "Folding Math_Cx::");
+	CORRECT_INFO_MESSAGE(log, "Folding Math_xC::Value to constant (1).");
+	INVALID_INFO_MESSAGE(log, "Folding Out::");
+
+	build_math_partial_test_graph(builder, NODE_MATH_POWER, 0.0f);
+	graph.finalize(&scene);
+}
+
+/*
+ * Tests: partial folding for Math Power with known 1.
+ */
+TEST(render_graph, constant_fold_part_math_pow_1)
+{
+	DEFINE_COMMON_VARIABLES(builder, log);
+
+	EXPECT_ANY_MESSAGE(log);
+	/* 1 ^ X == 1; X ^ 1 == X */
+	CORRECT_INFO_MESSAGE(log, "Folding Math_Cx::Value to constant (1)");
+	CORRECT_INFO_MESSAGE(log, "Folding Math_xC::Value to socket Attribute::Fac.");
+	INVALID_INFO_MESSAGE(log, "Folding Out::");
+
+	build_math_partial_test_graph(builder, NODE_MATH_POWER, 1.0f);
+	graph.finalize(&scene);
+}
+
 /*
  * Tests: Vector Math with all constant inputs.
  */

intern/cycles/util/util_path.cpp

@@ -778,7 +778,9 @@ static string line_directive(const string& path, int line)
 }
 
 
-string path_source_replace_includes(const string& source, const string& path)
+string path_source_replace_includes(const string& source,
+                                    const string& path,
+                                    const string& source_filename)
 {
 	/* Our own little c preprocessor that replaces #includes with the file
 	 * contents, to work around issue of opencl drivers not supporting
@@ -807,12 +809,12 @@ string path_source_replace_includes(const string& source, const string& path)
 						 * and avoids having list of include directories.x
 						 */
 						text = path_source_replace_includes(
-						        text, path_dirname(filepath));
-						text = path_source_replace_includes(text, path);
+						        text, path_dirname(filepath), filename);
+						text = path_source_replace_includes(text, path, filename);
 						/* Use line directives for better error messages. */
 						line = line_directive(filepath, 1)
 						     + token.replace(0, n_end + 1, "\n" + text + "\n")
-						     + line_directive(path, i);
+						     + line_directive(path_join(path, source_filename), i);
 					}
 				}
 			}

intern/cycles/util/util_path.h

@@ -66,7 +66,9 @@ bool path_read_text(const string& path, string& text);
 bool path_remove(const string& path);
 
 /* source code utility */
-string path_source_replace_includes(const string& source, const string& path);
+string path_source_replace_includes(const string& source,
+                                    const string& path,
+                                    const string& source_filename="");
 
 /* cache utility */
 void path_cache_clear_except(const string& name, const set<string>& except);

intern/elbeem/intern/solver_class.h

@@ -332,7 +332,7 @@ class LbmFsgrSolver :
 		void debugMarkCellCall(int level, int vi,int vj,int vk);
 		
 		// loop over grid, stream&collide update
-		void mainLoop(int lev);
+		void mainLoop(const int lev);
 		// change time step size
 		void adaptTimestep();
 		//! init mObjectSpeeds for current parametrization

intern/elbeem/intern/solver_main.cpp

@@ -355,7 +355,7 @@ void LbmFsgrSolver::fineAdvance()
 //! fine step function
 /*****************************************************************************/
 void 
-LbmFsgrSolver::mainLoop(int lev)
+LbmFsgrSolver::mainLoop(const int lev)
 {
 	// loops over _only inner_ cells  -----------------------------------------------------------------------------------
 	
@@ -376,13 +376,16 @@ LbmFsgrSolver::mainLoop(int lev)
   // main loop region
 	const bool doReduce = true;
 	const int gridLoopBound=1;
+	const int gDebugLevel = ::gDebugLevel;
+	int calcNumInvIfCells = 0;
+	LbmFloat calcInitialMass = 0;
 	GRID_REGION_INIT();
 #if PARALLEL==1
-#pragma omp parallel default(shared) num_threads(mNumOMPThreads) \
+#pragma omp parallel default(none) num_threads(mNumOMPThreads) \
   reduction(+: \
 	  calcCurrentMass,calcCurrentVolume, \
 		calcCellsFilled,calcCellsEmptied, \
-		calcNumUsedCells )
+		calcNumUsedCells,calcNumInvIfCells,calcInitialMass)
 	GRID_REGION_START();
 #else // PARALLEL==1
 	GRID_REGION_START();
@@ -468,7 +471,7 @@ LbmFsgrSolver::mainLoop(int lev)
 				calcCurrentMass += iniRho; 
 				calcCurrentVolume += 1.0; 
 				calcNumUsedCells++;
-				mInitialMass += iniRho;
+				calcInitialMass += iniRho;
 				// dont treat cell until next step
 				continue;
 			} 
@@ -479,7 +482,7 @@ LbmFsgrSolver::mainLoop(int lev)
 			if(isnotValid) {
 				// remove fluid cells, shouldnt be here anyway
 				LbmFloat fluidRho = m[0]; FORDF1 { fluidRho += m[l]; }
-				mInitialMass -= fluidRho;
+				calcInitialMass -= fluidRho;
 				const LbmFloat iniRho = 0.0;
 				RAC(tcel, dMass) = RAC(tcel, dFfrac) = iniRho;
 				RAC(tcel, dFlux) = FLUX_INIT;
@@ -608,8 +611,8 @@ LbmFsgrSolver::mainLoop(int lev)
 		// read distribution funtions of adjacent cells = stream step
 		DEFAULT_STREAM;
 
-		if((nbored & CFFluid)==0) { newFlag |= CFNoNbFluid; mNumInvIfCells++; }
-		if((nbored & CFEmpty)==0) { newFlag |= CFNoNbEmpty; mNumInvIfCells++; }
+		if((nbored & CFFluid)==0) { newFlag |= CFNoNbFluid; calcNumInvIfCells++; }
+		if((nbored & CFEmpty)==0) { newFlag |= CFNoNbEmpty; calcNumInvIfCells++; }
 
 		// calculate mass exchange for interface cells 
 		LbmFloat myfrac = RAC(ccel,dFfrac);
@@ -809,7 +812,7 @@ LbmFsgrSolver::mainLoop(int lev)
 			// fill if cells in inflow region
 			if(myfrac<0.5) { 
 				mass += 0.25; 
-				mInitialMass += 0.25;
+				calcInitialMass += 0.25;
 			}
 			const int OId = oldFlag>>24;
 			const LbmVec vel(mObjectSpeeds[OId]);
@@ -1013,7 +1016,7 @@ LbmFsgrSolver::mainLoop(int lev)
 		if( (mass) <= (rho * (   -FSGR_MAGICNR)) ) { ifemptied = 1; }
 
 		if(oldFlag & (CFMbndOutflow)) {
-			mInitialMass -= mass;
+			calcInitialMass -= mass;
 			mass = myfrac = 0.0;
 			iffilled = 0; ifemptied = 1;
 		}
@@ -1105,6 +1108,8 @@ LbmFsgrSolver::mainLoop(int lev)
 	mNumFilledCells  = calcCellsFilled;
 	mNumEmptiedCells = calcCellsEmptied;
 	mNumUsedCells = calcNumUsedCells;
+	mNumInvIfCells += calcNumInvIfCells;
+	mInitialMass += calcInitialMass;
 }
 
 
@@ -1115,13 +1120,14 @@ LbmFsgrSolver::preinitGrids()
 	const int lev = mMaxRefine;
 	const bool doReduce = false;
 	const int gridLoopBound=0;
+	const int gDebugLevel = ::gDebugLevel;
 
 	// preinit both grids
 	for(int s=0; s<2; s++) {
 	
 		GRID_REGION_INIT();
 #if PARALLEL==1
-#pragma omp parallel default(shared) num_threads(mNumOMPThreads) \
+#pragma omp parallel default(none) num_threads(mNumOMPThreads) \
   reduction(+: \
 	  calcCurrentMass,calcCurrentVolume, \
 		calcCellsFilled,calcCellsEmptied, \
@@ -1155,10 +1161,11 @@ LbmFsgrSolver::standingFluidPreinit()
 	const int lev = mMaxRefine;
 	const bool doReduce = false;
 	const int gridLoopBound=1;
+	const int gDebugLevel = ::gDebugLevel;
 
 	GRID_REGION_INIT();
 #if PARALLEL==1
-#pragma omp parallel default(shared) num_threads(mNumOMPThreads) \
+#pragma omp parallel default(none) num_threads(mNumOMPThreads) \
   reduction(+: \
 	  calcCurrentMass,calcCurrentVolume, \
 		calcCellsFilled,calcCellsEmptied, \

intern/ghost/intern/GHOST_SystemWin32.cpp

@@ -712,18 +712,26 @@ GHOST_EventCursor *GHOST_SystemWin32::processCursorEvent(GHOST_TEventType type,
 }
 
 
-GHOST_EventWheel *GHOST_SystemWin32::processWheelEvent(GHOST_WindowWin32 *window, WPARAM wParam, LPARAM lParam)
+void GHOST_SystemWin32::processWheelEvent(GHOST_WindowWin32 *window, WPARAM wParam, LPARAM lParam)
 {
-	// short fwKeys = LOWORD(wParam);			// key flags
-	int zDelta = (short) HIWORD(wParam);    // wheel rotation
-	
-	// zDelta /= WHEEL_DELTA;
-	// temporary fix below: microsoft now has added more precision, making the above division not work
-	zDelta = (zDelta <= 0) ? -1 : 1;
+	GHOST_SystemWin32 *system = (GHOST_SystemWin32 *)getSystem();
 
-	// short xPos = (short) LOWORD(lParam);	// horizontal position of pointer
-	// short yPos = (short) HIWORD(lParam);	// vertical position of pointer
-	return new GHOST_EventWheel(getSystem()->getMilliSeconds(), window, zDelta);
+	int acc = system->m_wheelDeltaAccum;
+	int delta = GET_WHEEL_DELTA_WPARAM(wParam);
+	
+	if (acc * delta < 0) {
+		// scroll direction reversed.
+		acc = 0;
+	}
+	acc += delta;
+	int direction = (acc >= 0) ? 1 : -1;
+	acc = abs(acc);
+	
+	while (acc >= WHEEL_DELTA) {
+		system->pushEvent(new GHOST_EventWheel(system->getMilliSeconds(), window, direction));
+		acc -= WHEEL_DELTA;
+	}
+	system->m_wheelDeltaAccum = acc * direction;
 }
 
 
@@ -1137,14 +1145,9 @@ LRESULT WINAPI GHOST_SystemWin32::s_wndProc(HWND hwnd, UINT msg, WPARAM wParam,
 					POINT mouse_pos = {GET_X_LPARAM(lParam), GET_Y_LPARAM(lParam)};
 					HWND mouse_hwnd = ChildWindowFromPoint(HWND_DESKTOP, mouse_pos);
 					GHOST_WindowWin32 *mouse_window = (GHOST_WindowWin32 *)::GetWindowLongPtr(mouse_hwnd, GWLP_USERDATA);
-					if (mouse_window != NULL) {
-						event = processWheelEvent(mouse_window, wParam, lParam);
-					}
-					else {
-						/* Happens when mouse is not over any of blender's windows. */
-						event = processWheelEvent(window, wParam, lParam);
-					}
-
+					
+					processWheelEvent(mouse_window ? mouse_window : window , wParam, lParam);
+					eventHandled = true;
 #ifdef BROKEN_PEEK_TOUCHPAD
 					PostMessage(hwnd, WM_USER, 0, 0);
 #endif
@@ -1203,6 +1206,7 @@ LRESULT WINAPI GHOST_SystemWin32::s_wndProc(HWND hwnd, UINT msg, WPARAM wParam,
 					GHOST_ModifierKeys modifiers;
 					modifiers.clear();
 					system->storeModifierKeys(modifiers);
+					system->m_wheelDeltaAccum = 0;
 					event = processWindowEvent(LOWORD(wParam) ? GHOST_kEventWindowActivate : GHOST_kEventWindowDeactivate, window);
 					/* WARNING: Let DefWindowProc handle WM_ACTIVATE, otherwise WM_MOUSEWHEEL
 					 * will not be dispatched to OUR active window if we minimize one of OUR windows. */

intern/ghost/intern/GHOST_SystemWin32.h

@@ -264,12 +264,12 @@ protected:
 	static GHOST_EventCursor *processCursorEvent(GHOST_TEventType type, GHOST_WindowWin32 *window);
 
 	/**
-	 * Creates a mouse wheel event.
+	 * Handles a mouse wheel event.
 	 * \param window	The window receiving the event (the active window).
 	 * \param wParam	The wParam from the wndproc
 	 * \param lParam	The lParam from the wndproc
 	 */
-	static GHOST_EventWheel *processWheelEvent(GHOST_WindowWin32 *window, WPARAM wParam, LPARAM lParam);
+	static void processWheelEvent(GHOST_WindowWin32 *window, WPARAM wParam, LPARAM lParam);
 
 	/**
 	 * Creates a key event and updates the key data stored locally (m_modifierKeys).
@@ -376,6 +376,9 @@ protected:
 
 	/** Console status */
 	int m_consoleStatus;
+
+	/** Wheel delta accumulator **/
+	int m_wheelDeltaAccum;
 };
 
 inline void GHOST_SystemWin32::retrieveModifierKeys(GHOST_ModifierKeys& keys) const