* optimized threading

 * break out with glare node
 * Added OpenCL kernels compatible with AMD still need some testing.

9 files changed, 191 insertions, 32 deletions

diff --git a/source/blender/compositor/intern/COM_ExecutionGroup.cpp b/source/blender/compositor/intern/COM_ExecutionGroup.cpp
index e46b4934217..7a53af7f58c 100644
--- a/source/blender/compositor/intern/COM_ExecutionGroup.cpp
+++ b/source/blender/compositor/intern/COM_ExecutionGroup.cpp
@@ -54,6 +54,7 @@ ExecutionGroup::ExecutionGroup()
 	this->numberOfChunks = 0;
 	this->initialized = false;
 	this->openCL = false;
+	this->singleThreaded = false;
 	this->chunksFinished = 0;
 }
 
@@ -100,6 +101,7 @@ void ExecutionGroup::addOperation(ExecutionSystem *system, NodeOperation *operat
 		if (!operation->isBufferOperation()) {
 			this->complex = operation->isComplex();
 			this->openCL = operation->isOpenCL();
+			this->singleThreaded = operation->isSingleThreaded();
 			this->initialized = true;
 		}
 		this->operations.push_back(operation);
@@ -191,10 +193,17 @@ void ExecutionGroup::determineResolution(unsigned int resolution[])
 
 void ExecutionGroup::determineNumberOfChunks()
 {
-	const float chunkSizef = this->chunkSize;
-	this->numberOfXChunks = ceil(this->width / chunkSizef);
-	this->numberOfYChunks = ceil(this->height / chunkSizef);
-	this->numberOfChunks = this->numberOfXChunks * this->numberOfYChunks;
+	if (singleThreaded) {
+		this->numberOfXChunks = 1;
+		this->numberOfYChunks = 1;
+		this->numberOfChunks = 1;
+	} 
+	else {
+		const float chunkSizef = this->chunkSize;
+		this->numberOfXChunks = ceil(this->width / chunkSizef);
+		this->numberOfYChunks = ceil(this->height / chunkSizef);
+		this->numberOfChunks = this->numberOfXChunks * this->numberOfYChunks;
+	}
 }
 
 /**
@@ -435,9 +444,14 @@ void ExecutionGroup::finalizeChunkExecution(int chunkNumber, MemoryBuffer** memo
 
 inline void ExecutionGroup::determineChunkRect(rcti *rect, const unsigned int xChunk, const unsigned int yChunk ) const
 {
-	const unsigned int minx = xChunk * chunkSize;
-	const unsigned int miny = yChunk * chunkSize;
-	BLI_init_rcti(rect, minx, min(minx + this->chunkSize, this->width), miny, min(miny + this->chunkSize, this->height));
+	if (singleThreaded) {
+		BLI_init_rcti(rect, 0, this->width, 0, this->height);
+	}
+	else {
+		const unsigned int minx = xChunk * chunkSize;
+		const unsigned int miny = yChunk * chunkSize;
+		BLI_init_rcti(rect, minx, min(minx + this->chunkSize, this->width), miny, min(miny + this->chunkSize, this->height));
+	}
 }
 
 void ExecutionGroup::determineChunkRect(rcti *rect, const unsigned int chunkNumber) const
@@ -462,6 +476,9 @@ MemoryBuffer *ExecutionGroup::allocateOutputBuffer(int chunkNumber, rcti *rect)
 
 bool ExecutionGroup::scheduleAreaWhenPossible(ExecutionSystem * graph, rcti *area)
 {
+	if (singleThreaded) {
+		return scheduleChunkWhenPossible(graph, 0, 0);
+	}
 	// find all chunks inside the rect
 	// determine minxchunk, minychunk, maxxchunk, maxychunk where x and y are chunknumbers
 
diff --git a/source/blender/compositor/intern/COM_ExecutionGroup.h b/source/blender/compositor/intern/COM_ExecutionGroup.h
index 416a78eb8b8..1698890cc34 100644
--- a/source/blender/compositor/intern/COM_ExecutionGroup.h
+++ b/source/blender/compositor/intern/COM_ExecutionGroup.h
@@ -63,10 +63,6 @@ class Device;
 class ExecutionGroup {
 private:
 	// fields
-	/**
-	  * @brief unique identifier of this node.
-	  */
-	string id;
 	
 	/**
 	  * @brief list of operations in this ExecutionGroup
@@ -121,6 +117,11 @@ private:
 	bool openCL;
 	
 	/**
+	 * @brief Is this Execution group SingleThreaded
+	 */
+	bool singleThreaded;
+	
+	/**
 	  * @brief what is the maximum number field of all ReadBufferOperation in this ExecutionGroup.
 	  * @note this is used to construct the MemoryBuffers that will be passed during execution.
 	  */
@@ -233,18 +234,7 @@ private:
 public:
 	// constructors
 	ExecutionGroup();
-	
-	/**
-	  * @brief set the id of this ExecutionGroup
-	  * @param id
-	  */
-	void setId(string id) {this->id = id;}
-	
-	/**
-	  * @brief return the id of this ExecutionGroup
-	  */
-	const string getId() const {return this->id;}
-	
+		
 	// methods
 	/**
 	  * @brief check to see if a NodeOperation is already inside this execution group
diff --git a/source/blender/compositor/intern/COM_ExecutionSystem.cpp b/source/blender/compositor/intern/COM_ExecutionSystem.cpp
index 1056c6d3f65..9681996c74d 100644
--- a/source/blender/compositor/intern/COM_ExecutionSystem.cpp
+++ b/source/blender/compositor/intern/COM_ExecutionSystem.cpp
@@ -124,6 +124,7 @@ void ExecutionSystem::execute()
 
 	for (index = 0 ; index < this->operations.size() ; index ++) {
 		NodeOperation * operation = this->operations[index];
+		operation->setbNodeTree(this->context.getbNodeTree());
 		operation->initExecution();
 	}
 	for (index = 0 ; index < this->groups.size() ; index ++) {
@@ -153,7 +154,7 @@ void ExecutionSystem::execute()
 
 void ExecutionSystem::executeGroups(CompositorPriority priority)
 {
-	int index;
+	unsigned int index;
 	vector<ExecutionGroup*> executionGroups;
 	this->findOutputExecutionGroup(&executionGroups, priority);
 
@@ -166,6 +167,7 @@ void ExecutionSystem::executeGroups(CompositorPriority priority)
 void ExecutionSystem::addOperation(NodeOperation *operation)
 {
 	ExecutionSystemHelper::addOperation(this->operations, operation);
+//	operation->setBTree
 }
 
 void ExecutionSystem::addReadWriteBufferOperations(NodeOperation *operation)
diff --git a/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp b/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp
index 75be8df74de..d5ca2ec619a 100644
--- a/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp
+++ b/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp
@@ -65,7 +65,7 @@ Node *ExecutionSystemHelper::addbNodeTree(ExecutionSystem &system, int nodes_sta
 	}
 
 	/* Expand group nodes */
-	for (int i=nodes_start; i < nodes.size(); ++i) {
+	for (unsigned int i=nodes_start; i < nodes.size(); ++i) {
 		Node *execnode = nodes[i];
 		if (execnode->isGroupNode()) {
 			GroupNode * groupNode = (GroupNode*)execnode;
diff --git a/source/blender/compositor/intern/COM_NodeOperation.cpp b/source/blender/compositor/intern/COM_NodeOperation.cpp
index 650e4af5ae0..148ad48ba3a 100644
--- a/source/blender/compositor/intern/COM_NodeOperation.cpp
+++ b/source/blender/compositor/intern/COM_NodeOperation.cpp
@@ -34,6 +34,7 @@ NodeOperation::NodeOperation()
 	this->width = 0;
 	this->height = 0;
 	this->openCL = false;
+	this->btree = NULL;
 }
 
 void NodeOperation::determineResolution(unsigned int resolution[], unsigned int preferredResolution[])
@@ -74,10 +75,22 @@ void NodeOperation::initMutex()
 {
 	BLI_mutex_init(&mutex);
 }
+
+void NodeOperation::lockMutex()
+{
+	BLI_mutex_lock(&mutex);
+}
+
+void NodeOperation::unlockMutex()
+{
+	BLI_mutex_unlock(&mutex);
+}
+
 void NodeOperation::deinitMutex()
 {
 	BLI_mutex_end(&mutex);
 }
+
 void NodeOperation::deinitExecution()
 {
 }
@@ -196,14 +209,15 @@ void NodeOperation::COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel,
 	size_t size[2];
 	cl_int2 offset;
 	
-	for (offsety = 0 ; offsety < height; offsety+=localSize) {
+	bool breaked = false;
+	for (offsety = 0 ; offsety < height && (!breaked); offsety+=localSize) {
 		offset[1] = offsety;
 		if (offsety+localSize < height) {
 			size[1] = localSize;
 		} else {
 			size[1] = height - offsety;
 		}
-		for (offsetx = 0 ; offsetx < width ; offsetx+=localSize) {
+		for (offsetx = 0 ; offsetx < width && (!breaked) ; offsetx+=localSize) {
 			if (offsetx+localSize < width) {
 				size[0] = localSize;
 			} else {
@@ -216,6 +230,9 @@ void NodeOperation::COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel,
 			error = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, size, 0, 0, 0, NULL);
 			if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));	}
 			clFlush(queue);
+			if (isBreaked()) {
+				breaked = false;
+			}
 		}
 	}
 }
diff --git a/source/blender/compositor/intern/COM_NodeOperation.h b/source/blender/compositor/intern/COM_NodeOperation.h
index 3f536fb3f2d..e56ac7bd51f 100644
--- a/source/blender/compositor/intern/COM_NodeOperation.h
+++ b/source/blender/compositor/intern/COM_NodeOperation.h
@@ -77,6 +77,11 @@ private:
 	  * @see NodeOperation.getMutex retrieve a pointer to this mutex.
 	  */
 	ThreadMutex mutex;
+	
+	/**
+	 * @brief reference to the editing bNodeTree only used for break callback
+	 */
+	const bNodeTree *btree;
 
 public:
 	/**
@@ -119,9 +124,10 @@ public:
 	  * for all other operations this will result in false.
 	  */
 	virtual int isBufferOperation() {return false;}
+	virtual int isSingleThreaded() {return false;}
 
+	void setbNodeTree(const bNodeTree * tree) {this->btree = tree;}
 	virtual void initExecution();
-	void initMutex();
 	
 	/**
 	  * @brief when a chunk is executed by a CPUDevice, this method is called
@@ -161,7 +167,6 @@ public:
 	  */
 	virtual void executeOpenCL(cl_context context,cl_program program, cl_command_queue queue, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer** inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, list<cl_kernel> *clKernelsToCleanUp) {}
 	virtual void deinitExecution();
-	void deinitMutex();
 
 	bool isResolutionSet() {
 		return this->width != 0 && height != 0;
@@ -236,6 +241,11 @@ public:
 	
 	virtual bool isViewerOperation() {return false;}
 	virtual bool isPreviewOperation() {return false;}
+	
+	inline bool isBreaked() {
+		return btree->test_break(btree->tbh);
+	}
+
 protected:
 	NodeOperation();
 
@@ -244,7 +254,11 @@ protected:
 	SocketReader *getInputSocketReader(unsigned int inputSocketindex);
 	NodeOperation *getInputOperation(unsigned int inputSocketindex);
 
-	inline ThreadMutex *getMutex() {return &this->mutex;}
+	void deinitMutex();
+	void initMutex();
+	void lockMutex();
+	void unlockMutex();
+	
 
 	/**
 	  * @brief set whether this operation is complex
@@ -264,7 +278,7 @@ protected:
 	static void COM_clAttachOutputMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, cl_mem clOutputMemoryBuffer);
 	void COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex);
 	static void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer* outputMemoryBuffer);
-	static void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex);
+	void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex);
 	cl_kernel COM_clCreateKernel(cl_program program, const char* kernelname, list<cl_kernel> *clKernelsToCleanUp);
 
 };
diff --git a/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.cpp b/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.cpp
new file mode 100644
index 00000000000..9ea90809de4
--- /dev/null
+++ b/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.cpp
@@ -0,0 +1,60 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ *
+ * Contributor: 
+ *		Jeroen Bakker 
+ *		Monique Dewanchand
+ */
+
+#include "COM_SingleThreadedNodeOperation.h"
+
+SingleThreadedNodeOperation::SingleThreadedNodeOperation(): NodeOperation()
+{
+	this->cachedInstance = NULL;
+	setComplex(true);
+}
+
+void SingleThreadedNodeOperation::initExecution()
+{
+	initMutex();
+}
+
+void SingleThreadedNodeOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data)
+{
+	this->cachedInstance->read(color, x, y);
+}
+
+void SingleThreadedNodeOperation::deinitExecution()
+{
+	deinitMutex();
+	if (this->cachedInstance) {
+		delete cachedInstance;
+		this->cachedInstance = NULL;
+	}
+}
+void *SingleThreadedNodeOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
+{
+	if (this->cachedInstance) return this->cachedInstance;
+	
+	lockMutex();
+	if (this->cachedInstance == NULL) {
+		//
+		this->cachedInstance = createMemoryBuffer(rect, memoryBuffers);
+	}
+	unlockMutex();
+	return this->cachedInstance;
+}
diff --git a/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.h b/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.h
new file mode 100644
index 00000000000..ace48365752
--- /dev/null
+++ b/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.h
@@ -0,0 +1,60 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ *
+ * Contributor: 
+ *		Jeroen Bakker 
+ *		Monique Dewanchand
+ */
+
+#ifndef _COM_SingleThreadedNodeOperation_h
+#define _COM_SingleThreadedNodeOperation_h
+#include "COM_NodeOperation.h"
+
+class SingleThreadedNodeOperation : public NodeOperation {
+private:
+	MemoryBuffer *cachedInstance;
+	
+protected:
+	inline bool isCached() {
+		return cachedInstance != NULL;
+	}
+
+public:
+	SingleThreadedNodeOperation();
+	
+	/**
+	  * the inner loop of this program
+	  */
+	void executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data);
+	
+	/**
+	  * Initialize the execution
+	  */
+	void initExecution();
+	
+	/**
+	  * Deinitialize the execution
+	  */
+	void deinitExecution();
+
+	void *initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers);
+
+	virtual MemoryBuffer* createMemoryBuffer(rcti *rect, MemoryBuffer **memoryBuffers) = 0;
+	
+	int isSingleThreaded() {return true;}
+};
+#endif
diff --git a/source/blender/compositor/intern/COM_compositor.cpp b/source/blender/compositor/intern/COM_compositor.cpp
index e27bff4401e..2bbfd18e7c5 100644
--- a/source/blender/compositor/intern/COM_compositor.cpp
+++ b/source/blender/compositor/intern/COM_compositor.cpp
@@ -51,7 +51,6 @@ void COM_execute(bNodeTree *editingtree, int rendering)
 
 	/* set progress bar to 0% and status to init compositing*/
 	editingtree->progress(editingtree->prh, 0.0);
-	editingtree->stats_draw(editingtree->sdh, (char*)"Compositing");
 
 	/* initialize execution system */
 	ExecutionSystem *system = new ExecutionSystem(editingtree, rendering);