* optimized threading
authorJeroen Bakker <j.bakker@atmind.nl>
Wed, 13 Jun 2012 12:34:56 +0000 (12:34 +0000)
committerJeroen Bakker <j.bakker@atmind.nl>
Wed, 13 Jun 2012 12:34:56 +0000 (12:34 +0000)
 * break out with glare node
 * Added OpenCL kernels compatible with AMD still need some testing.

40 files changed:
source/blender/compositor/CMakeLists.txt
source/blender/compositor/intern/COM_ExecutionGroup.cpp
source/blender/compositor/intern/COM_ExecutionGroup.h
source/blender/compositor/intern/COM_ExecutionSystem.cpp
source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp
source/blender/compositor/intern/COM_NodeOperation.cpp
source/blender/compositor/intern/COM_NodeOperation.h
source/blender/compositor/intern/COM_SingleThreadedNodeOperation.cpp [new file with mode: 0644]
source/blender/compositor/intern/COM_SingleThreadedNodeOperation.h [new file with mode: 0644]
source/blender/compositor/intern/COM_compositor.cpp
source/blender/compositor/nodes/COM_DilateErodeNode.cpp
source/blender/compositor/nodes/COM_MuteNode.cpp
source/blender/compositor/operations/COM_AntiAliasOperation.cpp
source/blender/compositor/operations/COM_CalculateMeanOperation.cpp
source/blender/compositor/operations/COM_CalculateStandardDeviationOperation.cpp
source/blender/compositor/operations/COM_CompositorOperation.cpp
source/blender/compositor/operations/COM_CompositorOperation.h
source/blender/compositor/operations/COM_DilateErodeOperation.cpp
source/blender/compositor/operations/COM_DilateErodeOperation.h
source/blender/compositor/operations/COM_DoubleEdgeMaskOperation.cpp
source/blender/compositor/operations/COM_FastGaussianBlurOperation.cpp
source/blender/compositor/operations/COM_GlareBaseOperation.cpp
source/blender/compositor/operations/COM_GlareBaseOperation.h
source/blender/compositor/operations/COM_GlareGhostOperation.cpp
source/blender/compositor/operations/COM_GlareSimpleStarOperation.cpp
source/blender/compositor/operations/COM_GlareStreaksOperation.cpp
source/blender/compositor/operations/COM_MaskOperation.cpp
source/blender/compositor/operations/COM_MovieDistortionOperation.cpp
source/blender/compositor/operations/COM_MultilayerImageOperation.cpp
source/blender/compositor/operations/COM_NormalizeOperation.cpp
source/blender/compositor/operations/COM_OpenCLKernels.cl
source/blender/compositor/operations/COM_OpenCLKernels.cl.h
source/blender/compositor/operations/COM_OutputFileOperation.cpp
source/blender/compositor/operations/COM_PreviewOperation.h
source/blender/compositor/operations/COM_TonemapOperation.cpp
source/blender/compositor/operations/COM_VectorBlurOperation.cpp
source/blender/compositor/operations/COM_ViewerBaseOperation.h
source/blender/compositor/operations/COM_ViewerOperation.cpp
source/blender/compositor/operations/COM_WriteBufferOperation.cpp
source/blender/compositor/operations/COM_WriteBufferOperation.h

index 22f72270734e7a87bc6b8f0562f42e56b4d9e0c4..bc2aeaefc834edbeff3cb12b0d71daa5e82c484b 100644 (file)
@@ -97,6 +97,9 @@ set(SRC
        intern/COM_CompositorContext.h
        intern/COM_ChannelInfo.cpp
        intern/COM_ChannelInfo.h
+       intern/COM_SingleThreadedNodeOperation.cpp
+       intern/COM_SingleThreadedNodeOperation.h
+
        operations/COM_QualityStepHelper.h
        operations/COM_QualityStepHelper.cpp
 
index e46b4934217deb4e4f7a3010f9fa8e3f72fb3304..7a53af7f58cb61eefc52f770ce76f8edec53caa8 100644 (file)
@@ -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
 
index 416a78eb8b872bb2b52da220e13feb21c18478d6..1698890cc34cce8f6c7e107e6d88a3ef2c1c2f0a 100644 (file)
@@ -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
@@ -120,6 +116,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
index 1056c6d3f65297854b0f51003328544ec6ae592a..9681996c74d04c616809141995fe9cd9583c7231 100644 (file)
@@ -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)
index 75be8df74de07abdaa0b2d06c0da009a0d993c3c..d5ca2ec619a926e30038d7b468ec84e66e6bb909 100644 (file)
@@ -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;
index 650e4af5ae091baf97c8092536ae4cb5abbfebea..148ad48ba3a1064e9848fb0886e30f9f1b5b61ed 100644 (file)
@@ -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;
+                       }
                }
        }
 }
index 3f536fb3f2d5f934192dc6590876cfb2da4e078f..e56ac7bd51f0b954085231cc5288231e985f3dbd 100644 (file)
@@ -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 (file)
index 0000000..9ea9080
--- /dev/null
@@ -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 (file)
index 0000000..ace4836
--- /dev/null
@@ -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
index e27bff4401e5bdbff17b04f8eaa68a0151e6ed72..2bbfd18e7c584fc1a644ddae95d4943f79cdbc44 100644 (file)
@@ -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);
index 477919568653c621743ac92a879f4b2c4630bf8f..0619bb5133e6940edc9bbe905d8b2b4886da76a1 100644 (file)
@@ -36,7 +36,7 @@ void DilateErodeNode::convertToOperations(ExecutionSystem *graph, CompositorCont
        
        bNode *editorNode = this->getbNode();
        if (editorNode->custom1 == CMP_NODE_DILATEERODE_DISTANCE_THRESH) {
-               DilateErodeDistanceOperation *operation = new DilateErodeDistanceOperation();
+               DilateErodeThresholdOperation *operation = new DilateErodeThresholdOperation();
                operation->setDistance(editorNode->custom2);
                operation->setInset(editorNode->custom3);
                
index d02eb2a0b980370119bb2fb927d0cc81d9b02974..93e352cfedeaf18272e393b4d6d7d278d62d5b00 100644 (file)
@@ -44,7 +44,7 @@ void MuteNode::reconnect(ExecutionSystem * graph, OutputSocket * output)
                }
        }
        
-       NodeOperation *operation;
+       NodeOperation *operation = NULL;
        switch (output->getDataType()) {
                case COM_DT_VALUE:
                {
@@ -74,7 +74,6 @@ void MuteNode::reconnect(ExecutionSystem * graph, OutputSocket * output)
                }
                        /* quiet warnings */
                case COM_DT_UNKNOWN:
-                       operation = NULL;
                        break;
        }
 
index 62639eeb24a2a2513e35cb9a29d7440a200a12ca..4cd9552b1087e21b5c4cabac2e9372612c834040 100644 (file)
@@ -44,7 +44,7 @@ void AntiAliasOperation::initExecution()
 
 void AntiAliasOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void * data)
 {
-       if (y < 0 || y >= this->height || x < 0 || x >= this->width) {
+       if (y < 0 || (unsigned int)y >= this->height || x < 0 || (unsigned int)x >= this->width) {
                color[0] = 0.0f;
        }
        else {
@@ -85,7 +85,7 @@ bool AntiAliasOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe
 void *AntiAliasOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
 {
        if (this->buffer) {return buffer;}
-       BLI_mutex_lock(getMutex());
+       lockMutex();
        if (this->buffer == NULL) {
                MemoryBuffer *tile = (MemoryBuffer*)valueReader->initializeTileData(rect, memoryBuffers);
                int size = tile->getHeight()*tile->getWidth();
@@ -100,6 +100,6 @@ void *AntiAliasOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu
                antialias_tagbuf(tile->getWidth(), tile->getHeight(), valuebuffer);
                this->buffer = valuebuffer;
        }
-       BLI_mutex_unlock(getMutex());
+       unlockMutex();
        return this->buffer;
 }
index a3438cea27b75b9a54e6219ab0f0f07e1f187dbb..077d8473f0b6a96a2ef99ac281c31049df187b77 100644 (file)
@@ -72,13 +72,13 @@ bool CalculateMeanOperation::determineDependingAreaOfInterest(rcti *input, ReadB
 
 void *CalculateMeanOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
 {
-       BLI_mutex_lock(getMutex());
+       lockMutex();
        if (!this->iscalculated) {
                MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers);
                calculateMean(tile);
                this->iscalculated = true;
        }
-       BLI_mutex_unlock(getMutex());
+       unlockMutex();
        return NULL;
 }
 
index 651c6674fdbfa7abced50c5d6488c8a4975a0800..dfe1b6aa329c214b417872c188f7b8038d34f76a 100644 (file)
@@ -37,7 +37,7 @@ void CalculateStandardDeviationOperation::executePixel(float *color, int x, int
 
 void *CalculateStandardDeviationOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
 {
-       BLI_mutex_lock(getMutex());
+       lockMutex();
        if (!this->iscalculated) {
                MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers);
                CalculateMeanOperation::calculateMean(tile);
@@ -92,6 +92,6 @@ void *CalculateStandardDeviationOperation::initializeTileData(rcti *rect, Memory
                this->standardDeviation = sqrt(sum / (float)(pixels-1));
                this->iscalculated = true;
        }
-       BLI_mutex_unlock(getMutex());
+       unlockMutex();
        return NULL;
 }
index c6e8faaa6389cc96c0211be893e2d50bec41c702..4ce7ffac9ef8e76a74b3dcca7f9463dc98fbaa4d 100644 (file)
@@ -59,7 +59,7 @@ void CompositorOperation::initExecution()
 
 void CompositorOperation::deinitExecution()
 {
-       if (tree->test_break && !tree->test_break(tree->tbh)) {
+       if (isBreaked()) {
                const Scene * scene = this->scene;
                Render *re = RE_GetRender(scene->id.name);
                RenderResult *rr = RE_AcquireResultWrite(re);
@@ -118,7 +118,7 @@ void CompositorOperation::executeRegion(rcti *rect, unsigned int tileNumber, Mem
                        buffer[offset+2] = color[2];
                        buffer[offset+3] = color[3];
                        offset +=COM_NUMBER_OF_CHANNELS;
-                       if (tree->test_break && tree->test_break(tree->tbh)) {
+                       if (isBreaked()) {
                                breaked = true;
                        }
                }
index 13cb4f2832437f1d6dc9547342431eeec9f63b14..36099b3eb91a0b8eb6387ed1da852810947c2174 100644 (file)
@@ -36,11 +36,6 @@ private:
          */
        const Scene *scene;
        
-       /**
-         * @brief local reference to the node tree
-         */
-       const bNodeTree *tree;
-       
        /**
          * @brief reference to the output float buffer
          */
@@ -59,7 +54,6 @@ public:
        CompositorOperation();
        void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer** memoryBuffers);
        void setScene(const Scene *scene) {this->scene = scene;}
-       void setbNodeTree(const bNodeTree *tree) {this->tree = tree;}
        bool isOutputOperation(bool rendering) const {return true;}
        void initExecution();
        void deinitExecution();
index 7bc49fa695c3876c7c3c6b12bb459326e2620304..bdd7362952a295aac61cf6223d7c8fa92d4081f7 100644 (file)
@@ -24,7 +24,7 @@
 #include "BLI_math.h"
 
 // DilateErode Distance Threshold
-DilateErodeDistanceOperation::DilateErodeDistanceOperation(): NodeOperation()
+DilateErodeThresholdOperation::DilateErodeThresholdOperation(): NodeOperation()
 {
        this->addInputSocket(COM_DT_VALUE);
        this->addOutputSocket(COM_DT_VALUE);
@@ -34,7 +34,7 @@ DilateErodeDistanceOperation::DilateErodeDistanceOperation(): NodeOperation()
        this->_switch = 0.5f;
        this->distance = 0.0f;
 }
-void DilateErodeDistanceOperation::initExecution()
+void DilateErodeThresholdOperation::initExecution()
 {
        this->inputProgram = this->getInputSocketReader(0);
        if (this->distance < 0.0f) {
@@ -53,13 +53,13 @@ void DilateErodeDistanceOperation::initExecution()
        }
 }
 
-void *DilateErodeDistanceOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
+void *DilateErodeThresholdOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
 {
        void *buffer = inputProgram->initializeTileData(NULL, memoryBuffers);
        return buffer;
 }
 
-void DilateErodeDistanceOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data)
+void DilateErodeThresholdOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data)
 {
        float inputValue[4];
        const float sw = this->_switch;
@@ -142,12 +142,12 @@ void DilateErodeDistanceOperation::executePixel(float *color, int x, int y, Memo
        }
 }
 
-void DilateErodeDistanceOperation::deinitExecution()
+void DilateErodeThresholdOperation::deinitExecution()
 {
        this->inputProgram = NULL;
 }
 
-bool DilateErodeDistanceOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output)
+bool DilateErodeThresholdOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output)
 {
        rcti newInput;
 
@@ -167,6 +167,7 @@ DilateDistanceOperation::DilateDistanceOperation(): NodeOperation()
        this->setComplex(true);
        this->inputProgram = NULL;
        this->distance = 0.0f;
+       this->setOpenCL(true);
 }
 void DilateDistanceOperation::initExecution()
 {
@@ -231,6 +232,28 @@ bool DilateDistanceOperation::determineDependingAreaOfInterest(rcti *input, Read
 
        return NodeOperation::determineDependingAreaOfInterest(&newInput, readOperation, output);
 }
+
+static cl_kernel dilateKernel = 0;
+void DilateDistanceOperation::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) 
+{
+       if (!dilateKernel) {
+               dilateKernel = COM_clCreateKernel(program, "dilateKernel", NULL);
+       }
+       cl_int distanceSquared = this->distance*this->distance;
+       cl_int scope = this->scope;
+       
+       COM_clAttachMemoryBufferToKernelParameter(context, dilateKernel, 0,  2, clMemToCleanUp, inputMemoryBuffers, this->inputProgram);
+       COM_clAttachOutputMemoryBufferToKernelParameter(dilateKernel, 1, clOutputBuffer);
+       COM_clAttachMemoryBufferOffsetToKernelParameter(dilateKernel, 3, outputMemoryBuffer);
+       clSetKernelArg(dilateKernel, 4, sizeof(cl_int), &scope);
+       clSetKernelArg(dilateKernel, 5, sizeof(cl_int), &distanceSquared);
+       COM_clAttachSizeToKernelParameter(dilateKernel, 6);
+       COM_clEnqueueRange(queue, dilateKernel, outputMemoryBuffer, 7);
+}
+
 // Erode Distance
 ErodeDistanceOperation::ErodeDistanceOperation() : DilateDistanceOperation() 
 {
@@ -268,6 +291,27 @@ void ErodeDistanceOperation::executePixel(float *color, int x, int y, MemoryBuff
        color[0] = value;
 }
 
+static cl_kernel erodeKernel = 0;
+void ErodeDistanceOperation::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) 
+{
+       if (!erodeKernel) {
+               erodeKernel = COM_clCreateKernel(program, "erodeKernel", NULL);
+       }
+       cl_int distanceSquared = this->distance*this->distance;
+       cl_int scope = this->scope;
+       
+       COM_clAttachMemoryBufferToKernelParameter(context, erodeKernel, 0,  2, clMemToCleanUp, inputMemoryBuffers, this->inputProgram);
+       COM_clAttachOutputMemoryBufferToKernelParameter(erodeKernel, 1, clOutputBuffer);
+       COM_clAttachMemoryBufferOffsetToKernelParameter(erodeKernel, 3, outputMemoryBuffer);
+       clSetKernelArg(erodeKernel, 4, sizeof(cl_int), &scope);
+       clSetKernelArg(erodeKernel, 5, sizeof(cl_int), &distanceSquared);
+       COM_clAttachSizeToKernelParameter(erodeKernel, 6);
+       COM_clEnqueueRange(queue, erodeKernel, outputMemoryBuffer, 7);
+}
+
 // Dilate step
 DilateStepOperation::DilateStepOperation(): NodeOperation()
 {
@@ -288,7 +332,7 @@ void *DilateStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB
        if (this->cached_buffer != NULL) {
                return this->cached_buffer;
        }
-       BLI_mutex_lock(getMutex());
+       lockMutex();
        if (this->cached_buffer == NULL) {
                MemoryBuffer *buffer = (MemoryBuffer*)inputProgram->initializeTileData(NULL, memoryBuffers);
                float *rectf = buffer->convertToValueBuffer();
@@ -327,7 +371,7 @@ void *DilateStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB
                }
                this->cached_buffer = rectf;
        }
-       BLI_mutex_unlock(getMutex());
+       unlockMutex();
        return this->cached_buffer;
 }
 
@@ -374,7 +418,7 @@ void *ErodeStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu
        if (this->cached_buffer != NULL) {
                return this->cached_buffer;
        }
-       BLI_mutex_lock(getMutex());
+       lockMutex();
        if (this->cached_buffer == NULL) {
                MemoryBuffer *buffer = (MemoryBuffer*)inputProgram->initializeTileData(NULL, memoryBuffers);
                float *rectf = buffer->convertToValueBuffer();
@@ -413,6 +457,6 @@ void *ErodeStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu
                }
                this->cached_buffer = rectf;
        }
-       BLI_mutex_unlock(getMutex());
+       unlockMutex();
        return this->cached_buffer;
 }
index 71bbab74a4bda013723409d7e6e2de14de19222c..63cee4c333a9c380bd9fc7984aeacd3a251106e0 100644 (file)
@@ -25,7 +25,7 @@
 #include "COM_NodeOperation.h"
 
 
-class DilateErodeDistanceOperation : public NodeOperation {
+class DilateErodeThresholdOperation : public NodeOperation {
 private:
        /**
          * Cached reference to the inputProgram
@@ -42,7 +42,7 @@ private:
          */
        int scope;
 public:
-       DilateErodeDistanceOperation();
+       DilateErodeThresholdOperation();
        
        /**
          * the inner loop of this program
@@ -70,11 +70,11 @@ public:
 
 class DilateDistanceOperation : public NodeOperation {
 private:
+protected:
        /**
          * Cached reference to the inputProgram
          */
        SocketReader * inputProgram;
-protected:
        float distance;
        int scope;
 public:
@@ -98,6 +98,11 @@ public:
        
        void setDistance(float distance) {this->distance = distance;}
        bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output);
+       
+       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);
 };
 class ErodeDistanceOperation : public DilateDistanceOperation {
 public:
@@ -107,6 +112,11 @@ public:
          * the inner loop of this program
          */
        void executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data);
+
+       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);
 };
 
 class DilateStepOperation : public NodeOperation {
index 40f492b0f10394ccaeed4ea72b512b2781c4b037..df04b889200d4e8b60360f8523f0ffa1efa1195a 100644 (file)
@@ -23,6 +23,7 @@
 #include "COM_DoubleEdgeMaskOperation.h"
 #include "BLI_math.h"
 #include "DNA_node_types.h"
+#include "MEM_guardedalloc.h"
 
 // this part has been copied from the double edge mask
 // Contributor(s): Peter Larabell.
@@ -1215,12 +1216,12 @@ void DoubleEdgeMaskOperation::doDoubleEdgeMask(float *imask, float *omask, float
                gsz=rsize[2];                  // by the do_*EdgeDetection() function.
                
                fsz=gsz+isz+osz;                                   // calculate size of pixel index buffer needed
-               gbuf = new unsigned short[fsz*2];     // allocate edge/gradient pixel index buffer
+               gbuf = (unsigned short*)MEM_callocN(sizeof (unsigned short)*fsz*2, "DEM"); // allocate edge/gradient pixel index buffer
                
                do_createEdgeLocationBuffer(t,rw,lres,res,gbuf,&innerEdgeOffset,&outerEdgeOffset,isz,gsz);
                do_fillGradientBuffer(rw,res,gbuf,isz,osz,gsz,innerEdgeOffset,outerEdgeOffset);
                
-               delete [] gbuf;          // free the gradient index buffer
+               MEM_freeN(gbuf);          // free the gradient index buffer
        }
 }
 
@@ -1263,7 +1264,7 @@ void *DoubleEdgeMaskOperation::initializeTileData(rcti *rect, MemoryBuffer **mem
 {
        if (this->cachedInstance) return this->cachedInstance;
        
-       BLI_mutex_lock(getMutex());
+       lockMutex();
        if (this->cachedInstance == NULL) {
                MemoryBuffer *innerMask = (MemoryBuffer*)inputInnerMask->initializeTileData(rect, memoryBuffers);
                MemoryBuffer *outerMask = (MemoryBuffer*)inputOuterMask->initializeTileData(rect, memoryBuffers);
@@ -1275,7 +1276,7 @@ void *DoubleEdgeMaskOperation::initializeTileData(rcti *rect, MemoryBuffer **mem
                delete omask;
                this->cachedInstance = data;
        }
-       BLI_mutex_unlock(getMutex());
+       unlockMutex();
        return this->cachedInstance;
 }
 void DoubleEdgeMaskOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data)
index ad58631f2c1b0d37aecab018bca9e0b43204839e..92e1aace9e2ba1d655ea1e346021bbc61409f5fb 100644 (file)
@@ -79,7 +79,7 @@ void FastGaussianBlurOperation::deinitExecution()
 
 void *FastGaussianBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
 {
-       BLI_mutex_lock(this->getMutex());
+       lockMutex();
        if (!iirgaus) {
                MemoryBuffer *newBuf = (MemoryBuffer*)this->inputProgram->initializeTileData(rect, memoryBuffers);
                MemoryBuffer *copy = newBuf->duplicate();
@@ -105,7 +105,7 @@ void *FastGaussianBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **m
                }
                this->iirgaus = copy;
        }
-       BLI_mutex_unlock(this->getMutex());
+       unlockMutex();
        return iirgaus;
 }
 
index fdfd19a10ae33f0be7590af11d35099e725f80da..83132963f0b3526d303e33c64ccf48f0f8785445 100644 (file)
 #include "COM_GlareBaseOperation.h"
 #include "BLI_math.h"
 
-GlareBaseOperation::GlareBaseOperation(): NodeOperation()
+GlareBaseOperation::GlareBaseOperation(): SingleThreadedNodeOperation()
 {
        this->addInputSocket(COM_DT_COLOR);
        this->addOutputSocket(COM_DT_COLOR);
        this->settings = NULL;
-       this->cachedInstance = NULL;
-       setComplex(true);
 }
 void GlareBaseOperation::initExecution()
 {
-       initMutex();
+       SingleThreadedNodeOperation::initExecution();
        this->inputProgram = getInputSocketReader(0);
-       this->cachedInstance = NULL;
-}
-
-void GlareBaseOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data)\
-{
-       float *buffer = (float*) data;
-       int index = (y*this->getWidth() + x) * COM_NUMBER_OF_CHANNELS;
-       color[0] = buffer[index];
-       color[1] = buffer[index+1];
-       color[2] = buffer[index+2];
-       color[3] = buffer[index+3];
 }
 
 void GlareBaseOperation::deinitExecution()
 {
-       deinitMutex();
        this->inputProgram = NULL;
-       if (this->cachedInstance) {
-               delete cachedInstance;
-               this->cachedInstance = NULL;
-       }
+       SingleThreadedNodeOperation::deinitExecution();
 }
-void *GlareBaseOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
+
+MemoryBuffer *GlareBaseOperation::createMemoryBuffer(rcti *rect2, MemoryBuffer **memoryBuffers)
 {
-       BLI_mutex_lock(getMutex());
-       if (this->cachedInstance == NULL) {
-               MemoryBuffer *tile = (MemoryBuffer*)inputProgram->initializeTileData(rect, memoryBuffers);
-               float *data = new float[this->getWidth()*this->getHeight()*COM_NUMBER_OF_CHANNELS];
-               this->generateGlare(data, tile, this->settings);
-               this->cachedInstance = data;
-       }
-       BLI_mutex_unlock(getMutex());
-       return this->cachedInstance;
+       MemoryBuffer *tile = (MemoryBuffer*)inputProgram->initializeTileData(rect2, memoryBuffers);
+       rcti rect;
+       rect.xmin = 0;
+       rect.ymin = 0;
+       rect.xmax = getWidth();
+       rect.ymax = getHeight();
+       MemoryBuffer *result = new MemoryBuffer(NULL, &rect);
+       float *data = result->getBuffer();
+       this->generateGlare(data, tile, this->settings);
+       return result;
 }
 
 bool GlareBaseOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output)
 {
-       if (this->cachedInstance != NULL) {
+       if (isCached()) {
                return false;
        }
        else {
index 2fa8afc9c4c5be6426491bd7a3f98de9af484ad6..d79f449f42614d1f1d7923520fffb9d7b07b3aad 100644 (file)
@@ -22,7 +22,8 @@
 
 #ifndef _COM_GlareBaseOperation_h
 #define _COM_GlareBaseOperation_h
-#include "COM_NodeOperation.h"
+
+#include "COM_SingleThreadedNodeOperation.h"
 #include "DNA_node_types.h"
 
 
@@ -55,7 +56,7 @@ typedef float fRGB[4];
                           } (void)0
 
 
-class GlareBaseOperation : public NodeOperation {
+class GlareBaseOperation : public SingleThreadedNodeOperation {
 private:
        /**
          * @brief Cached reference to the inputProgram
@@ -66,16 +67,7 @@ private:
          * @brief settings of the glare node.
          */
        NodeGlare * settings;
-       
-       float *cachedInstance;
-
 public:
-       
-       /**
-         * the inner loop of this program
-         */
-       void executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data);
-       
        /**
          * Initialize the execution
          */
@@ -86,15 +78,15 @@ public:
          */
        void deinitExecution();
 
-       void *initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers);
-
        void setGlareSettings(NodeGlare * settings) {this->settings = settings;}
        bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output);
+       
 protected:
        GlareBaseOperation();
        
        virtual void generateGlare(float *data, MemoryBuffer *inputTile, NodeGlare *settings) = 0;
        
+       MemoryBuffer *createMemoryBuffer(rcti *rect, MemoryBuffer **memoryBuffers);
        
 };
 #endif
index c5b1d6caa89fd21852af896ee72570a98e082c70..383a13c54def1ec566b8f68c91b16b884d7e3ca4 100644 (file)
@@ -45,15 +45,21 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No
        MemoryBuffer *gbuf = inputTile->duplicate();
        MemoryBuffer *tbuf1 = inputTile->duplicate();
 
+       bool breaked = false;
+       
        FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 0, 3);
-       FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 1, 3);
-       FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 2, 3);
+       if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 1, 3);
+       if (isBreaked()) breaked = true;
+       if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 2, 3);
        
        MemoryBuffer *tbuf2 = tbuf1->duplicate();
        
-       FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 0, 3);
-       FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 1, 3);
-       FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 2, 3);
+       if (isBreaked()) breaked = true;
+       if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 0, 3);
+       if (isBreaked()) breaked = true;
+       if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 1, 3);
+       if (isBreaked()) breaked = true;
+       if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 2, 3);
        
        if (settings->iter & 1) ofs = 0.5f; else ofs = 0.f;
        for (x=0; x<(settings->iter*4); x++) {
@@ -68,7 +74,7 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No
 
        sc = 2.13;
        isc = -0.97;
-       for (y=0; y<gbuf->getHeight(); y++) {
+       for (y=0; y<gbuf->getHeight() &(!breaked); y++) {
                v = (float)(y+0.5f) / (float)gbuf->getHeight();
                for (x=0; x<gbuf->getWidth(); x++) {
                        u = (float)(x+0.5f) / (float)gbuf->getWidth();
@@ -83,11 +89,13 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No
                        
                        gbuf->writePixel(x, y, c);
                }
+               if (isBreaked()) breaked = true;
+               
        }
 
        memset(tbuf1->getBuffer(), 0, tbuf1->getWidth()*tbuf1->getHeight()*COM_NUMBER_OF_CHANNELS*sizeof(float));
-       for (n=1; n<settings->iter; n++) {
-               for (y=0; y<gbuf->getHeight(); y++) {
+       for (n=1; n<settings->iter &(!breaked); n++) {
+               for (y=0; y<gbuf->getHeight()&(!breaked); y++) {
                        v = (float)(y+0.5f) / (float)gbuf->getHeight();
                        for (x=0; x<gbuf->getWidth(); x++) {
                                u = (float)(x+0.5f) / (float)gbuf->getWidth();
@@ -103,6 +111,7 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No
                                }
                                tbuf1->writePixel(x, y, tc);
                        }
+                       if (isBreaked()) breaked = true;
                }
                memcpy(gbuf->getBuffer(), tbuf1->getBuffer(), tbuf1->getWidth()*tbuf1->getHeight()*COM_NUMBER_OF_CHANNELS*sizeof(float));
        }
index fba3eca4af93ae373dc0e21bab353c29d43c4740..4a393a330733ad0ccaa9ad8b1a702590f22d432f 100644 (file)
@@ -32,10 +32,11 @@ void GlareSimpleStarOperation::generateGlare(float *data, MemoryBuffer *inputTil
        MemoryBuffer *tbuf1 = inputTile->duplicate();
        MemoryBuffer *tbuf2 = inputTile->duplicate();
 
-       for (i=0; i<settings->iter; i++) {
+       bool breaked = false;
+       for (i=0; i<settings->iter && (!breaked); i++) {
 //             // (x || x-1, y-1) to (x || x+1, y+1)
 //             // F
-               for (y=0; y<this->getHeight(); y++) {
+               for (y=0; y<this->getHeight() && (!breaked); y++) {
                        ym = y - i;
                        yp = y + i;
                        for (x=0; x<this->getWidth(); x++) {
@@ -58,11 +59,13 @@ void GlareSimpleStarOperation::generateGlare(float *data, MemoryBuffer *inputTil
                                madd_v3_v3fl(c, tc, f2);
                                c[3] = 1.0f;
                                tbuf2->writePixel(x, y, c);
-
+                       }
+                       if (isBreaked()) {
+                               breaked = true;
                        }
                }
 //             // B
-               for (y=tbuf1->getHeight()-1; y>=0; y--) {
+               for (y=tbuf1->getHeight()-1 && (!breaked); y>=0; y--) {
                        ym = y - i;
                        yp = y + i;
                        for (x=tbuf1->getWidth()-1; x>=0; x--) {
@@ -86,6 +89,9 @@ void GlareSimpleStarOperation::generateGlare(float *data, MemoryBuffer *inputTil
                                c[3] = 1.0f;
                                tbuf2->writePixel(x, y, c);
                        }
+                       if (isBreaked()) {
+                               breaked = true;
+                       }
                }
        }
 
index 42b6a2b5e50daa978df228be438a7b166e948423..e735893ed6d6daf2ad69ab7cfa2a34305d1219a4 100644 (file)
@@ -33,22 +33,23 @@ void GlareStreaksOperation::generateGlare(float *data, MemoryBuffer *inputTile,
        int size = inputTile->getWidth()*inputTile->getHeight();
        int size4 = size*4;
 
+       bool breaked = false;
        
        MemoryBuffer *tsrc = inputTile->duplicate();
        MemoryBuffer *tdst = new MemoryBuffer(NULL, inputTile->getRect());
        tdst->clear();
        memset(data, 0, size4*sizeof(float));
        
-       for (a=0.f; a<DEG2RADF(360.0f); a+=ang) {
+       for (a=0.f; a<DEG2RADF(360.0f) && (!breaked); a+=ang) {
                const float an = a + settings->angle_ofs;
                const float vx = cos((double)an), vy = sin((double)an);
-               for (n=0; n<settings->iter; ++n) {
+               for (n=0; n<settings->iter && (!breaked); ++n) {
                        const float p4 = pow(4.0, (double)n);
                        const float vxp = vx*p4, vyp = vy*p4;
                        const float wt = pow((double)settings->fade, (double)p4);
                        const float cmo = 1.f - (float)pow((double)settings->colmod, (double)n+1);      // colormodulation amount relative to current pass
                        float *tdstcol = tdst->getBuffer();
-                       for (y=0; y<tsrc->getHeight(); ++y) {
+                       for (y=0; y<tsrc->getHeight() && (!breaked); ++y) {
                                for (x=0; x<tsrc->getWidth(); ++x, tdstcol+=4) {
                                        // first pass no offset, always same for every pass, exact copy,
                                        // otherwise results in uneven brightness, only need once
@@ -71,11 +72,13 @@ void GlareStreaksOperation::generateGlare(float *data, MemoryBuffer *inputTile,
                                        tdstcol[2] = 0.5f*(tdstcol[2] + c1[2] + wt*(c2[2] + wt*(c3[2] + wt*c4[2])));
                                        tdstcol[3] = 1.0f;
                                }
+                               if (isBreaked()) {
+                                       breaked = true;
+                               }
                        }
                        memcpy(tsrc->getBuffer(), tdst->getBuffer(), sizeof(float)*size4);
                }
 
-//             addImage(sbuf, tsrc, 1.f/(float)(6 - ndg->iter)); // add result to data @todo
                float *sourcebuffer = tsrc->getBuffer();
                float factor = 1.f/(float)(6 - settings->iter);
                for (int i = 0 ; i < size4; i ++) {
index 35174349a63fc2b633f55abbf664242defe5314d..bfbf2b42e825fe7252f045f9503b2e295264eae6 100644 (file)
@@ -67,7 +67,7 @@ void *MaskOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers
        if (!this->mask)
                return NULL;
 
-       BLI_mutex_lock(getMutex());
+       lockMutex();
        if (this->rasterizedMask == NULL) {
                int width = this->getWidth();
                int height = this->getHeight();
@@ -78,8 +78,7 @@ void *MaskOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers
 
                this->rasterizedMask = buffer;
        }
-       BLI_mutex_unlock(getMutex());
-
+       unlockMutex();
        return this->rasterizedMask;
 }
 
index ebea9e8b4a27cdc861cba75b32243fdb97c08d53..d9e8977871f8aed0f42c50dd7e2d9490df21e4e3 100644 (file)
@@ -52,7 +52,7 @@ void MovieDistortionOperation::initExecution()
                BKE_movieclip_user_set_frame(&clipUser, this->framenumber);
                BKE_movieclip_get_size(this->movieClip, &clipUser, &calibration_width, &calibration_height);
 
-               for (int i = 0 ; i < s_cache.size() ; i ++) {
+               for (unsigned int i = 0 ; i < s_cache.size() ; i ++) {
                        DistortionCache *c = (DistortionCache*)s_cache[i];
                        if (c->isCacheFor(this->movieClip, this->width, this->height,
                                          calibration_width, calibration_height, this->distortion))
index f95dd12a81a21efd6f715e9ca193c02e539d2b60..5e81cd639dd0580b6d5832f08f1757b95d47871b 100644 (file)
@@ -47,7 +47,7 @@ void MultilayerColorOperation::executePixel(float *color, float x, float y, Pixe
 {
        int yi = y;
        int xi = x;
-       if (this->imageBuffer == NULL || xi < 0 || yi < 0 || xi >= this->getWidth() || yi >= this->getHeight() ) {
+       if (this->imageBuffer == NULL || xi < 0 || yi < 0 || (unsigned int)xi >= this->getWidth() || (unsigned int)yi >= this->getHeight() ) {
                color[0] = 0.0f;
                color[1] = 0.0f;
                color[2] = 0.0f;
@@ -80,7 +80,7 @@ void MultilayerValueOperation::executePixel(float *color, float x, float y, Pixe
 {
        int yi = y;
        int xi = x;
-       if (this->imageBuffer == NULL || xi < 0 || yi < 0 || xi >= this->getWidth() || yi >= this->getHeight() ) {
+       if (this->imageBuffer == NULL || xi < 0 || yi < 0 || (unsigned int)xi >= this->getWidth() || (unsigned int)yi >= this->getHeight() ) {
                color[0] = 0.0f;
        }
        else {
@@ -93,7 +93,7 @@ void MultilayerVectorOperation::executePixel(float *color, float x, float y, Pix
 {
        int yi = y;
        int xi = x;
-       if (this->imageBuffer == NULL || xi < 0 || yi < 0 || xi >= this->getWidth() || yi >= this->getHeight() ) {
+       if (this->imageBuffer == NULL || xi < 0 || yi < 0 || (unsigned int)xi >= this->getWidth() || (unsigned int)yi >= this->getHeight() ) {
                color[0] = 0.0f;
        }
        else {
index d144739f8456daefc6299e2cd2b0099b236f05bd..4dd94943f4eecc242f1d8904f3be7b6b84f3e374 100644 (file)
@@ -76,8 +76,7 @@ bool NormalizeOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe
 
 void *NormalizeOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
 {
-       BLI_mutex_lock(getMutex());
-
+       lockMutex();
        if (this->cachedInstance == NULL) {
                MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers);
                /* using generic two floats struct to store x: min  y: mult */
@@ -105,7 +104,7 @@ void *NormalizeOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu
                this->cachedInstance = minmult;
        }
 
-       BLI_mutex_unlock(getMutex());
+       unlockMutex();
        return this->cachedInstance;
 }
 
index aeccfcab8b5e276bfcded7dacb5580df7fc20c70..e1f175b318a217c8d0a50137e51d41d9cf422371 100644 (file)
@@ -6,8 +6,8 @@ const sampler_t SAMPLER_NEAREST      = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS
 __constant const int2 zero = {0,0};
 
 // KERNEL --- BOKEH BLUR ---
-__kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __global __read_only image2d_t inputImage, 
-                              __global __read_only image2d_t bokehImage, __global __write_only image2d_t output, 
+__kernel void bokehBlurKernel(__read_only image2d_t boundingBox, __read_only image2d_t inputImage, 
+                              __read_only image2d_t bokehImage, __write_only image2d_t output, 
                               int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset) 
 {
        int2 coords = {get_global_id(0), get_global_id(1)}; 
@@ -50,3 +50,65 @@ __kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __glob
        
        write_imagef(output, coords, color);
 }
+
+// KERNEL --- DILATE ---
+__kernel void dilateKernel(__read_only image2d_t inputImage,  __write_only image2d_t output,
+                           int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension, 
+                           int2 offset)
+{
+       int2 coords = {get_global_id(0), get_global_id(1)}; 
+       coords += offset;
+       const int2 realCoordinate = coords + offsetOutput;
+
+       const int2 minXY = max(realCoordinate - scope, zero);
+       const int2 maxXY = min(realCoordinate + scope, dimension);
+       
+       float value = 0.0f;
+       int nx, ny;
+       int2 inputXy;
+       
+       for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {
+               for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {
+                       const float deltaX = (realCoordinate.x - nx);
+                       const float deltaY = (realCoordinate.y - ny);
+                       const float measuredDistance = deltaX*deltaX+deltaY*deltaY;
+                       if (measuredDistance <= distanceSquared) {
+                               value = max(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);
+                       }
+               }
+       }
+
+       float4 color = {value,0.0f,0.0f,0.0f};
+       write_imagef(output, coords, color);
+}
+
+// KERNEL --- DILATE ---
+__kernel void erodeKernel(__read_only image2d_t inputImage,  __write_only image2d_t output,
+                           int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension, 
+                           int2 offset)
+{
+       int2 coords = {get_global_id(0), get_global_id(1)}; 
+       coords += offset;
+       const int2 realCoordinate = coords + offsetOutput;
+
+       const int2 minXY = max(realCoordinate - scope, zero);
+       const int2 maxXY = min(realCoordinate + scope, dimension);
+       
+       float value = 1.0f;
+       int nx, ny;
+       int2 inputXy;
+       
+       for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {
+               for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {
+                       const float deltaX = (realCoordinate.x - nx);
+                       const float deltaY = (realCoordinate.y - ny);
+                       const float measuredDistance = deltaX*deltaX+deltaY*deltaY;
+                       if (measuredDistance <= distanceSquared) {
+                               value = min(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);
+                       }
+               }
+       }
+
+       float4 color = {value,0.0f,0.0f,0.0f};
+       write_imagef(output, coords, color);
+}
index 3cf33c752726d8fc3a3e48ef73e7a1ae3267b1d4..ef8668f6f21a4321aaa8ecb889bf983ae065df00 100644 (file)
@@ -8,8 +8,8 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope
 "__constant const int2 zero = {0,0};\n" \
 "\n" \
 "// KERNEL --- BOKEH BLUR ---\n" \
-"__kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __global __read_only image2d_t inputImage,\n" \
-"                              __global __read_only image2d_t bokehImage, __global __write_only image2d_t output,\n" \
+"__kernel void bokehBlurKernel(__read_only image2d_t boundingBox, __read_only image2d_t inputImage,\n" \
+"                              __read_only image2d_t bokehImage, __write_only image2d_t output,\n" \
 "                              int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset)\n" \
 "{\n" \
 "      int2 coords = {get_global_id(0), get_global_id(1)};\n" \
@@ -26,8 +26,8 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope
 "      if (tempBoundingBox > 0.0f) {\n" \
 "              const int2 bokehImageDim = get_image_dim(bokehImage);\n" \
 "              const int2 bokehImageCenter = bokehImageDim/2;\n" \
-"              const int2 minXY = max(realCoordinate - radius, zero);;\n" \
-"              const int2 maxXY = min(realCoordinate + radius, dimension);;\n" \
+"              const int2 minXY = max(realCoordinate - radius, zero);\n" \
+"              const int2 maxXY = min(realCoordinate + radius, dimension);\n" \
 "              int nx, ny;\n" \
 "\n" \
 "              float2 uv;\n" \
@@ -52,4 +52,66 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope
 "\n" \
 "      write_imagef(output, coords, color);\n" \
 "}\n" \
+"\n" \
+"// KERNEL --- DILATE ---\n" \
+"__kernel void dilateKernel(__read_only image2d_t inputImage,  __write_only image2d_t output,\n" \
+"                           int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,\n" \
+"                           int2 offset)\n" \
+"{\n" \
+"      int2 coords = {get_global_id(0), get_global_id(1)};\n" \
+"      coords += offset;\n" \
+"      const int2 realCoordinate = coords + offsetOutput;\n" \
+"\n" \
+"      const int2 minXY = max(realCoordinate - scope, zero);\n" \
+"      const int2 maxXY = min(realCoordinate + scope, dimension);\n" \
+"\n" \
+"      float value = 0.0f;\n" \
+"      int nx, ny;\n" \
+"      int2 inputXy;\n" \
+"\n" \
+"      for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\n" \
+"              for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {\n" \
+"                      const float deltaX = (realCoordinate.x - nx);\n" \
+"                      const float deltaY = (realCoordinate.y - ny);\n" \
+"                      const float measuredDistance = deltaX*deltaX+deltaY*deltaY;\n" \
+"                      if (measuredDistance <= distanceSquared) {\n" \
+"                              value = max(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \
+"                      }\n" \
+"              }\n" \
+"      }\n" \
+"\n" \
+"      float4 color = {value,0.0f,0.0f,0.0f};\n" \
+"      write_imagef(output, coords, color);\n" \
+"}\n" \
+"\n" \
+"// KERNEL --- DILATE ---\n" \
+"__kernel void erodeKernel(__read_only image2d_t inputImage,  __write_only image2d_t output,\n" \
+"                           int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,\n" \
+"                           int2 offset)\n" \
+"{\n" \
+"      int2 coords = {get_global_id(0), get_global_id(1)};\n" \
+"      coords += offset;\n" \
+"      const int2 realCoordinate = coords + offsetOutput;\n" \
+"\n" \
+"      const int2 minXY = max(realCoordinate - scope, zero);\n" \
+"      const int2 maxXY = min(realCoordinate + scope, dimension);\n" \
+"\n" \
+"      float value = 1.0f;\n" \
+"      int nx, ny;\n" \
+"      int2 inputXy;\n" \
+"\n" \
+"      for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\n" \
+"              for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {\n" \
+"                      const float deltaX = (realCoordinate.x - nx);\n" \
+"                      const float deltaY = (realCoordinate.y - ny);\n" \
+"                      const float measuredDistance = deltaX*deltaX+deltaY*deltaY;\n" \
+"                      if (measuredDistance <= distanceSquared) {\n" \
+"                              value = min(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \
+"                      }\n" \
+"              }\n" \
+"      }\n" \
+"\n" \
+"      float4 color = {value,0.0f,0.0f,0.0f};\n" \
+"      write_imagef(output, coords, color);\n" \
+"}\n" \
 "\0";
index 8d39e987bd44c90c7827d3fbd82c14a20d3575dd..1438116f31300d6bbe1fb6cf1d932cd2f0b1be32 100644 (file)
@@ -177,7 +177,7 @@ void OutputOpenExrMultiLayerOperation::add_layer(const char *name, DataType data
 
 void OutputOpenExrMultiLayerOperation::initExecution()
 {
-       for (int i=0; i < layers.size(); ++i) {
+       for (unsigned int i=0; i < layers.size(); ++i) {
                layers[i].imageInput = getInputSocketReader(i);
                layers[i].outputBuffer = init_buffer(this->getWidth(), this->getHeight(), layers[i].datatype);
        }
@@ -185,7 +185,7 @@ void OutputOpenExrMultiLayerOperation::initExecution()
 
 void OutputOpenExrMultiLayerOperation::executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer** memoryBuffers)
 {
-       for (int i=0; i < layers.size(); ++i) {
+       for (unsigned int i=0; i < layers.size(); ++i) {
                write_buffer_rect(rect, memoryBuffers, this->tree, layers[i].imageInput, layers[i].outputBuffer, this->getWidth(), layers[i].datatype);
        }
 }
@@ -203,7 +203,7 @@ void OutputOpenExrMultiLayerOperation::deinitExecution()
                                  (this->scene->r.scemode & R_EXTENSION), true);
                BLI_make_existing_file(filename);
                
-               for (int i=0; i < layers.size(); ++i) {
+               for (unsigned int i=0; i < layers.size(); ++i) {
                        char channelname[EXR_TOT_MAXNAME];
                        BLI_strncpy(channelname, layers[i].name, sizeof(channelname)-2);
                        char *channelname_ext = channelname + strlen(channelname);
@@ -251,7 +251,7 @@ void OutputOpenExrMultiLayerOperation::deinitExecution()
                }
                
                IMB_exr_close(exrhandle);
-               for (int i=0; i < layers.size(); ++i) {
+               for (unsigned int i=0; i < layers.size(); ++i) {
                        if (layers[i].outputBuffer) {
                                MEM_freeN(layers[i].outputBuffer);
                                layers[i].outputBuffer = NULL;
index 2b81b9147461e83c0b6d0a39349ed9f88924f67e..32c0b6ecc77513e4739c793e6920ad170a1f0766 100644 (file)
@@ -34,7 +34,6 @@ protected:
          * @brief holds reference to the SDNA bNode, where this nodes will render the preview image for
          */
        bNode *node;
-       const bNodeTree *tree;
        SocketReader *input;
        float divider;
 
@@ -48,7 +47,6 @@ public:
        void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer **memoryBuffers);
        void determineResolution(unsigned int resolution[], unsigned int preferredResolution[]);
        void setbNode(bNode *node) { this->node = node;}
-       void setbNodeTree(const bNodeTree *tree) { this->tree = tree;}
        bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output);
        bool isPreviewOperation() {return true;}
        
index 75adcf524c4f711f4dcbff2988b619632e45bf9f..9a02d6a58ca854564197247afa1775cc3fef8f86 100644 (file)
@@ -118,7 +118,7 @@ bool TonemapOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferO
 
 void *TonemapOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
 {
-       BLI_mutex_lock(getMutex());
+       lockMutex();
        if (this->cachedInstance == NULL) {
                MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers);
                AvgLogLum *data = new AvgLogLum();
@@ -150,7 +150,7 @@ void *TonemapOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuff
                data->igm = (this->data->gamma==0.f) ? 1 : (1.f / this->data->gamma);
                this->cachedInstance = data;
        }
-       BLI_mutex_unlock(getMutex());
+       unlockMutex();
        return this->cachedInstance;
 }
 
index e6305dc26a2840b5faffb16d48ecd07b0c61ff67..3dac3547b8acfd6e6a7f4ca8183d49b2cc1920e6 100644 (file)
@@ -77,7 +77,7 @@ void *VectorBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB
 {
        if (this->cachedInstance) return this->cachedInstance;
        
-       BLI_mutex_lock(getMutex());
+       lockMutex();
        if (this->cachedInstance == NULL) {
                MemoryBuffer *tile = (MemoryBuffer*)inputImageProgram->initializeTileData(rect, memoryBuffers);
                MemoryBuffer *speed = (MemoryBuffer*)inputSpeedProgram->initializeTileData(rect, memoryBuffers);
@@ -87,7 +87,7 @@ void *VectorBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB
                this->generateVectorBlur(data, tile, speed, z);
                this->cachedInstance = data;
        }
-       BLI_mutex_unlock(getMutex());
+       unlockMutex();
        return this->cachedInstance;
 }
 
index 51fa8cecc0d752226dc04ca2509e8faca6b3318c..8864894f0ede108b1aa81bdd3097cd8b26173d4c 100644 (file)
@@ -34,7 +34,6 @@ protected:
        ImageUser * imageUser;
        void *lock;
        bool active;
-       const bNodeTree *tree;
        float centerX;
        float centerY;
        OrderOfChunks chunkOrder;
@@ -49,7 +48,6 @@ public:
        void setImageUser(ImageUser *imageUser) {this->imageUser = imageUser;}
        const bool isActiveViewerOutput() const {return active;}
        void setActive(bool active) {this->active = active;}
-       void setbNodeTree(const bNodeTree *tree) {this->tree = tree;}
        void setCenterX(float centerX) {this->centerX = centerX;}
        void setCenterY(float centerY) {this->centerY = centerY;}
        void setChunkOrder(OrderOfChunks tileOrder) {this->chunkOrder = tileOrder;}
index 22e6511fbe7f0d1c39e9705042c960c2f9fa4fba..3bec1ef48f52fc28f312bd59ea331c4139242df4 100644 (file)
@@ -104,7 +104,7 @@ void ViewerOperation::executeRegion(rcti *rect, unsigned int tileNumber, MemoryB
 
                        offset +=4;
                }
-               if (tree->test_break && tree->test_break(tree->tbh)) {
+               if (isBreaked()) {
                        breaked = true;
                }
 
index 8888d30ba2fda11403454c1934a68939b18424b2..087ab50cf4789eb3e6417034d922b72599f44390 100644 (file)
@@ -30,7 +30,6 @@ WriteBufferOperation::WriteBufferOperation() :NodeOperation()
        this->memoryProxy = new MemoryProxy();
        this->memoryProxy->setWriteBufferOperation(this);
        this->memoryProxy->setExecutor(NULL);
-       this->tree = NULL;
 }
 WriteBufferOperation::~WriteBufferOperation()
 {
@@ -78,7 +77,7 @@ void WriteBufferOperation::executeRegion(rcti *rect, unsigned int tileNumber, Me
                                offset4 +=COM_NUMBER_OF_CHANNELS;
 
                        }
-                       if (tree->test_break && tree->test_break(tree->tbh)) {
+                       if (isBreaked()) {
                                breaked = true;
                        }
 
@@ -103,7 +102,7 @@ void WriteBufferOperation::executeRegion(rcti *rect, unsigned int tileNumber, Me
                                input->read(&(buffer[offset4]), x, y, COM_PS_NEAREST, memoryBuffers);
                                offset4 +=COM_NUMBER_OF_CHANNELS;
                        }
-                       if (tree->test_break && tree->test_break(tree->tbh)) {
+                       if (isBreaked()) {
                                breaked = true;
                        }
                }
index 068adc032934384a8869a8033b2271fff346e47a..fff3212b4109f2cc0989c85a6821a8cd5638603d 100644 (file)
@@ -33,7 +33,6 @@
 class WriteBufferOperation: public NodeOperation {
        MemoryProxy *memoryProxy;
        NodeOperation *input;
-       const bNodeTree * tree;
 public:
        WriteBufferOperation();
        ~WriteBufferOperation();
@@ -45,7 +44,6 @@ public:
        void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer** memoryBuffers);
        void initExecution();
        void deinitExecution();
-       void setbNodeTree(const bNodeTree *tree) {this->tree = tree;}
        void executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer** memoryBuffers, MemoryBuffer* outputBuffer);
        void readResolutionFromInputSocket();