* Added OpenCL kernel for bokeh blur
authorJeroen Bakker <j.bakker@atmind.nl>
Fri, 8 Jun 2012 09:17:07 +0000 (09:17 +0000)
committerJeroen Bakker <j.bakker@atmind.nl>
Fri, 8 Jun 2012 09:17:07 +0000 (09:17 +0000)
 * Uncomment COM_OPENCL_ENABLED from COM_defines.h to test

52 files changed:
release/datafiles/clkernelstoh.py [new file with mode: 0755]
source/blender/compositor/COM_defines.h
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_ExecutionSystem.h
source/blender/compositor/intern/COM_Node.cpp
source/blender/compositor/intern/COM_Node.h
source/blender/compositor/intern/COM_NodeOperation.cpp
source/blender/compositor/intern/COM_NodeOperation.h
source/blender/compositor/intern/COM_OpenCLDevice.cpp
source/blender/compositor/intern/COM_WorkScheduler.cpp
source/blender/compositor/nodes/COM_BlurNode.cpp
source/blender/compositor/nodes/COM_BokehBlurNode.cpp
source/blender/compositor/nodes/COM_BokehImageNode.cpp
source/blender/compositor/nodes/COM_ChannelMatteNode.cpp
source/blender/compositor/nodes/COM_ChromaMatteNode.cpp
source/blender/compositor/nodes/COM_ColorCurveNode.cpp
source/blender/compositor/nodes/COM_ColorMatteNode.cpp
source/blender/compositor/nodes/COM_CompositorNode.cpp
source/blender/compositor/nodes/COM_DifferenceMatteNode.cpp
source/blender/compositor/nodes/COM_DistanceMatteNode.cpp
source/blender/compositor/nodes/COM_FilterNode.cpp
source/blender/compositor/nodes/COM_ImageNode.cpp
source/blender/compositor/nodes/COM_LuminanceMatteNode.cpp
source/blender/compositor/nodes/COM_MixNode.cpp
source/blender/compositor/nodes/COM_MovieClipNode.cpp
source/blender/compositor/nodes/COM_OutputFileNode.cpp
source/blender/compositor/nodes/COM_RenderLayersNode.cpp
source/blender/compositor/nodes/COM_SplitViewerNode.cpp
source/blender/compositor/nodes/COM_TextureNode.cpp
source/blender/compositor/nodes/COM_ViewerNode.cpp
source/blender/compositor/operations/COM_BokehBlurOperation.cpp
source/blender/compositor/operations/COM_BokehBlurOperation.h
source/blender/compositor/operations/COM_ColorCurveOperation.cpp
source/blender/compositor/operations/COM_ColorCurveOperation.h
source/blender/compositor/operations/COM_CompositorOperation.h
source/blender/compositor/operations/COM_OpenCLKernels.cl
source/blender/compositor/operations/COM_OpenCLKernels.cl.cpp [deleted file]
source/blender/compositor/operations/COM_OpenCLKernels.cl.h [new file with mode: 0644]
source/blender/compositor/operations/COM_OutputFileOperation.h
source/blender/compositor/operations/COM_PreviewOperation.cpp
source/blender/compositor/operations/COM_PreviewOperation.h
source/blender/compositor/operations/COM_ViewerBaseOperation.cpp
source/blender/compositor/operations/COM_ViewerBaseOperation.h
source/blender/compositor/operations/COM_WriteBufferOperation.cpp
source/blender/compositor/operations/COM_WriteBufferOperation.h
source/blender/editors/space_node/drawnode.c
source/blender/makesdna/DNA_node_types.h
source/blender/makesrna/intern/rna_nodetree.c
source/blender/makesrna/intern/rna_nodetree_types.h
source/blender/nodes/composite/nodes/node_composite_bokehblur.c

diff --git a/release/datafiles/clkernelstoh.py b/release/datafiles/clkernelstoh.py
new file mode 100755 (executable)
index 0000000..8fb5d4e
--- /dev/null
@@ -0,0 +1,70 @@
+#!/usr/bin/python
+# -*- coding: utf-8 -*-
+
+# ***** BEGIN GPL LICENSE BLOCK *****
+#
+# 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.
+#
+# The Original Code is Copyright (C) 2012 Blender Foundation.
+# All rights reserved.
+#
+# Contributor(s): Jeroen Bakker
+#
+# ***** END GPL LICENCE BLOCK *****
+
+# <pep8 compliant>
+
+import sys
+import os
+
+if len(sys.argv) < 2:
+    sys.stdout.write("Usage: clkernelstoh <cl_file>\n")
+    sys.exit(1)
+
+filename = sys.argv[1]
+
+try:
+    fpin = open(filename, "r")
+except:
+    sys.stdout.write("Unable to open input %s\n" % sys.argv[1])
+    sys.exit(1)
+
+if filename[0:2] == "." + os.sep:
+    filename = filename[2:]
+
+cname = filename + ".h"
+sys.stdout.write("Making H file <%s>\n" % cname)
+
+filename = filename.split("/")[-1].split("\\")[-1]
+filename = filename.replace(".", "_")
+
+try:
+    fpout = open(cname, "w")
+except:
+    sys.stdout.write("Unable to open output %s\n" % cname)
+    sys.exit(1)
+
+fpout.write("/* clkernelstoh output of file <%s> */\n\n" % filename)
+fpout.write("const char * clkernelstoh_%s = " % filename)
+
+lines = fpin.readlines()
+for line in lines:
+       fpout.write("\"")
+       fpout.write(line.rstrip())
+       fpout.write("\\n\" \\\n")
+fpout.write("\"\\0\";\n")
+
+fpin.close()
+fpout.close()
index df807091cb85463ecd6374e9c1c595747dc1fa9a..f87265c50f5ab202e4f5b83ff68e3bcce38f77c2 100644 (file)
@@ -52,12 +52,25 @@ typedef enum CompositorQuality {
        COM_QUALITY_LOW    = 2
 } CompositorQuality;
 
+/**
+  * @brief Possible priority settings
+  * @ingroup Execution
+  */
+typedef enum CompositorPriority {
+       /** @brief High quality setting */
+       COM_PRIORITY_HIGH   = 2 ,
+       /** @brief Medium quality setting */
+       COM_PRIORITY_MEDIUM = 1,
+       /** @brief Low quality setting */
+       COM_PRIORITY_LOW    = 0
+} CompositorPriority;
+
 // configurable items
 
 // chunk size determination
 #define COM_PREVIEW_SIZE 140.0f
-#define COM_OPENCL_ENABLED
-#define COM_PREVIEW_ENABLED
+//#define COM_OPENCL_ENABLED
+
 // workscheduler threading models
 /**
   * COM_TM_QUEUE is a multithreaded model, which uses the BLI_thread_queue pattern. This is the default option.
index 0ebf9af207ba3b09892d48d34999248b3d414f2f..44b3c8dafbb368681fec916294da957cc6a3d3ec 100644 (file)
@@ -57,7 +57,7 @@ ExecutionGroup::ExecutionGroup()
        this->chunksFinished = 0;
 }
 
-int ExecutionGroup::getRenderPriotrity()
+CompositorPriority ExecutionGroup::getRenderPriotrity()
 {
        return this->getOutputNodeOperation()->getRenderPriority();
 }
@@ -401,47 +401,11 @@ MemoryBuffer** ExecutionGroup::getInputBuffersOpenCL(int chunkNumber)
        return memoryBuffers;
 }
 
-// @todo: for opencl the memory buffers size needs to be same as the needed size
-// currently this method is not called, but will be when opencl nodes are created
 MemoryBuffer *ExecutionGroup::constructConsolidatedMemoryBuffer(MemoryProxy *memoryProxy, rcti *rect)
 {
-       // find all chunks inside the rect
-       // determine minxchunk, minychunk, maxxchunk, maxychunk where x and y are chunknumbers
-       float chunkSizef = this->chunkSize;
-
-       int indexx, indexy;
-
-       const int minxchunk = floor(rect->xmin/chunkSizef);
-       const int maxxchunk = ceil((rect->xmax-1)/chunkSizef);
-       const int minychunk = floor(rect->ymin/chunkSizef);
-       const int maxychunk = ceil((rect->ymax-1)/chunkSizef);
-
-       if (maxxchunk== minxchunk+1 && maxychunk == minychunk+1) {
-               MemoryBuffer *result =memoryProxy->getBuffer();
-               return result;
-       }
-
-       rcti chunkRect;
-       chunkRect.xmin = minxchunk*this->chunkSize;
-       chunkRect.xmax = maxxchunk*this->chunkSize;
-       chunkRect.ymin = minychunk*this->chunkSize;
-       chunkRect.ymax = maxychunk*this->chunkSize;
-
-       CLAMP(chunkRect.xmin, 0, (int)this->width);
-       CLAMP(chunkRect.xmax, 0, (int)this->width);
-       CLAMP(chunkRect.ymin, 0, (int)this->height);
-       CLAMP(chunkRect.ymax, 0, (int)this->height);
-
-       MemoryBuffer *result = new MemoryBuffer(memoryProxy, &chunkRect);
-
-       for (indexx = max(minxchunk, 0); indexx<min((int)this->numberOfXChunks, maxxchunk) ; indexx++) {
-               for (indexy = max(minychunk, 0); indexy<min((int)this->numberOfYChunks, maxychunk) ; indexy++) {
-                       /* int chunkNumber = indexx+indexy*this->numberOfXChunks; */ /* UNUSED */
-                       MemoryBuffer *chunkBuffer = memoryProxy->getBuffer();
-                       result->copyContentFrom(chunkBuffer);
-               }
-       }
-
+       MemoryBuffer* imageBuffer = memoryProxy->getBuffer();
+       MemoryBuffer* result = new MemoryBuffer(memoryProxy, rect);
+       result->copyContentFrom(imageBuffer);
        return result;
 }
 
@@ -487,14 +451,14 @@ void ExecutionGroup::determineChunkRect(rcti *rect, const unsigned int chunkNumb
 
 MemoryBuffer *ExecutionGroup::allocateOutputBuffer(int chunkNumber, rcti *rect)
 {
-       MemoryBuffer *outputBuffer = NULL;
-       // output allocation is only valid when our outputoperation is a memorywriter
+       // we asume that this method is only called from complex execution groups.
        NodeOperation * operation = this->getOutputNodeOperation();
        if (operation->isWriteBufferOperation()) {
-/*             WriteBufferOperation *writeOperation = (WriteBufferOperation*)operation; */ /* UNUSED */
-// @todo               outputBuffer = MemoryManager::allocateMemoryBuffer(writeOperation->getMemoryProxy(), chunkNumber, rect);
+               WriteBufferOperation *writeOperation = (WriteBufferOperation*)operation;
+               MemoryBuffer *buffer = new MemoryBuffer(writeOperation->getMemoryProxy(), rect);
+               return buffer;
        }
-       return outputBuffer;
+       return NULL;
 }
 
 
@@ -600,11 +564,6 @@ void ExecutionGroup::determineDependingMemoryProxies(vector<MemoryProxy*> *memor
        }
 }
 
-bool ExecutionGroup::operator ==(const ExecutionGroup & executionGroup) const
-{
-       return this->getOutputNodeOperation() == executionGroup.getOutputNodeOperation();
-}
-
 bool ExecutionGroup::isOpenCL()
 {
        return this->openCL;
index cbdc9bb1787a3b38191a589db69e66186e349274..416a78eb8b872bb2b52da220e13feb21c18478d6 100644 (file)
@@ -167,13 +167,7 @@ private:
          * @param operation the operation to be added
          */
        bool canContainOperation(NodeOperation *operation);
-       
-       /**
-         * @brief get the Render priority of this ExecutionGroup
-         * @see ExecutionSystem.execute
-         */
-       int getRenderPriotrity();
-       
+               
        /**
          * @brief calculate the actual chunk size of this execution group.
          * @note A chunk size is an unsigned int that is both the height and width of a chunk.
@@ -396,17 +390,21 @@ public:
          * @see determineChunkSize()
          */
        void determineChunkRect(rcti *rect, const unsigned int chunkNumber) const;
-       
-       
-               bool operator ==(const ExecutionGroup &executionGroup) const;
-       
-               /**
+
+       /**
          * @brief can this ExecutionGroup be scheduled on an OpenCLDevice
          * @see WorkScheduler.schedule
          */
        bool isOpenCL();
-       
+
        void setChunksize(int chunksize) {this->chunkSize = chunksize;}
+
+       /**
+         * @brief get the Render priority of this ExecutionGroup
+         * @see ExecutionSystem.execute
+         */
+       CompositorPriority getRenderPriotrity();
+       
 };
 
 #endif
index 96d2a6f44342b12becd5576f891d1d2fe98fb63b..8c0b37a0685f9273cfa59b8bde84bf1afb991e6b 100644 (file)
@@ -127,20 +127,9 @@ void ExecutionSystem::execute()
 
        WorkScheduler::start(this->context);
 
-
-       vector<ExecutionGroup*> executionGroups;
-       this->findOutputExecutionGroup(&executionGroups);
-
-       /* start execution of the ExecutionGroups based on priority of their output node */
-       for (int priority = 9 ; priority>=0 ; priority--) {
-               for (index = 0 ; index < executionGroups.size(); index ++) {
-                       ExecutionGroup *group = executionGroups[index];
-                       NodeOperation *output = group->getOutputNodeOperation();
-                       if (output->getRenderPriority() == priority) {
-                               group->execute(this);
-                       }
-               }
-       }
+       executeGroups(COM_PRIORITY_HIGH);
+       executeGroups(COM_PRIORITY_MEDIUM);
+       executeGroups(COM_PRIORITY_LOW);
 
        WorkScheduler::finish();
        WorkScheduler::stop();
@@ -155,6 +144,18 @@ void ExecutionSystem::execute()
        }
 }
 
+void ExecutionSystem::executeGroups(CompositorPriority priority)
+{
+       int index;
+       vector<ExecutionGroup*> executionGroups;
+       this->findOutputExecutionGroup(&executionGroups, priority);
+
+       for (index = 0 ; index < executionGroups.size(); index ++) {
+               ExecutionGroup *group = executionGroups[index];
+               group->execute(this);
+       }
+}
+
 void ExecutionSystem::addOperation(NodeOperation *operation)
 {
        ExecutionSystemHelper::addOperation(this->operations, operation);
@@ -304,6 +305,17 @@ void ExecutionSystem::determineActualSocketDataTypes(vector<NodeBase*> &nodes)
        }
 }
 
+void ExecutionSystem::findOutputExecutionGroup(vector<ExecutionGroup*> *result, CompositorPriority priority) const
+{
+       unsigned int index;
+       for (index = 0 ; index < this->groups.size() ; index ++) {
+               ExecutionGroup *group = this->groups[index];
+               if (group->isOutputExecutionGroup() && group->getRenderPriotrity() == priority) {
+                       result->push_back(group);
+               }
+       }
+}
+
 void ExecutionSystem::findOutputExecutionGroup(vector<ExecutionGroup*> *result) const
 {
        unsigned int index;
index 85fec8b6145ae4290b589993e17d073f96c6467e..510e58ba1bbe5feeeec418901c9288551725c514 100644 (file)
@@ -138,6 +138,11 @@ private: //methods
        void addReadWriteBufferOperations(NodeOperation *operation);
 
 
+       /**
+         * find all execution group with output nodes
+         */
+       void findOutputExecutionGroup(vector<ExecutionGroup*> *result, CompositorPriority priority) const;
+       
        /**
          * find all execution group with output nodes
          */
@@ -224,6 +229,8 @@ private:
          * @param nodes list of nodes or operations to do the data type determination
          */
        void determineActualSocketDataTypes(vector<NodeBase*> &nodes);
+       
+       void executeGroups(CompositorPriority priority);
 
 };
 #endif
index ba5e21d53ae58f44901f10fe4edfa09e3f0fa315..264725b4b2c2718581619777413607c8c32ab33e 100644 (file)
@@ -83,23 +83,20 @@ void Node::addSetValueOperation(ExecutionSystem *graph, InputSocket *inputsocket
        graph->addOperation(operation);
 }
 
-void Node::addPreviewOperation(ExecutionSystem *system, OutputSocket *outputSocket, int priority)
+void Node::addPreviewOperation(ExecutionSystem *system, OutputSocket *outputSocket)
 {
-#ifdef COM_PREVIEW_ENABLED
        PreviewOperation *operation = new PreviewOperation();
        system->addOperation(operation);
        operation->setbNode(this->getbNode());
        operation->setbNodeTree(system->getContext().getbNodeTree());
-       operation->setPriority(priority);
        this->addLink(system, outputSocket, operation->getInputSocket(0));
-#endif
 }
 
-void Node::addPreviewOperation(ExecutionSystem *system, InputSocket *inputSocket, int priority)
+void Node::addPreviewOperation(ExecutionSystem *system, InputSocket *inputSocket)
 {
        if (inputSocket->isConnected()) {
                OutputSocket *outputsocket = inputSocket->getConnection()->getFromSocket();
-               this->addPreviewOperation(system, outputsocket, priority);
+               this->addPreviewOperation(system, outputsocket);
        }
 }
 
index 2666d0a698058a9d21b3de40402b17d1c4aaa691..23744adf64280b40694868ec3e0f980fdc717db9 100644 (file)
@@ -120,8 +120,8 @@ protected:
        
        Node();
        
-       void addPreviewOperation(ExecutionSystem *system, InputSocket *inputSocket, int priority);
-       void addPreviewOperation(ExecutionSystem *system, OutputSocket *inputSocket, int priority);
+       void addPreviewOperation(ExecutionSystem *system, InputSocket *inputSocket);
+       void addPreviewOperation(ExecutionSystem *system, OutputSocket *outputSocket);
        
        bNodeSocket *getEditorInputSocket(int editorNodeInputSocketIndex);
        bNodeSocket *getEditorOutputSocket(int editorNodeOutputSocketIndex);
index fae652e39d706b76e5e860a4eaa948b1387d806e..650e4af5ae091baf97c8092536ae4cb5abbfebea 100644 (file)
@@ -124,3 +124,111 @@ bool NodeOperation::determineDependingAreaOfInterest(rcti * input, ReadBufferOpe
                return false;
        }
 }
+
+cl_mem NodeOperation::COM_clAttachMemoryBufferToKernelParameter(cl_context context, cl_kernel kernel, int parameterIndex, int offsetIndex, list<cl_mem> *cleanup, MemoryBuffer **inputMemoryBuffers, SocketReader* reader) 
+{
+       cl_int error;
+       MemoryBuffer* result = (MemoryBuffer*)reader->initializeTileData(NULL, inputMemoryBuffers);
+
+       const cl_image_format imageFormat = {
+               CL_RGBA,
+               CL_FLOAT
+       };
+
+       cl_mem clBuffer = clCreateImage2D(context, CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR, &imageFormat, result->getWidth(), 
+                                         result->getHeight(), 0, result->getBuffer(), &error);
+       
+       if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  }
+       if (error == CL_SUCCESS) cleanup->push_back(clBuffer);
+
+       error = clSetKernelArg(kernel, parameterIndex, sizeof(cl_mem), &clBuffer);
+       if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  }
+       
+       COM_clAttachMemoryBufferOffsetToKernelParameter(kernel, offsetIndex, result);
+       return clBuffer;
+}
+       
+void NodeOperation::COM_clAttachMemoryBufferOffsetToKernelParameter(cl_kernel kernel, int offsetIndex, MemoryBuffer *memoryBuffer) 
+{
+       if (offsetIndex != -1) {
+               cl_int error;
+               rcti* rect = memoryBuffer->getRect();
+               cl_int2 offset = {rect->xmin, rect->ymin};
+
+               error = clSetKernelArg(kernel, offsetIndex, sizeof(cl_int2), &offset);
+               if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  }
+       }
+}
+
+void NodeOperation::COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex) 
+{
+       if (offsetIndex != -1) {
+               cl_int error;
+               cl_int2 offset = {this->getWidth(), this->getHeight()};
+
+               error = clSetKernelArg(kernel, offsetIndex, sizeof(cl_int2), &offset);
+               if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  }
+       }
+}
+
+void NodeOperation::COM_clAttachOutputMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, cl_mem clOutputMemoryBuffer) 
+{
+       cl_int error;
+       error = clSetKernelArg(kernel, parameterIndex, sizeof(cl_mem), &clOutputMemoryBuffer);
+       if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); }
+}
+
+void NodeOperation::COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer) {
+       cl_int error;
+       const size_t size[] = {outputMemoryBuffer->getWidth(),outputMemoryBuffer->getHeight()};
+       
+       error = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, size, 0, 0, 0, NULL);
+       if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  }
+}
+
+void NodeOperation::COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex) {
+       cl_int error;
+       const int width = outputMemoryBuffer->getWidth();
+       const int height = outputMemoryBuffer->getHeight();
+       int offsetx;
+       int offsety;
+       const int localSize = 32;
+       size_t size[2];
+       cl_int2 offset;
+       
+       for (offsety = 0 ; offsety < height; offsety+=localSize) {
+               offset[1] = offsety;
+               if (offsety+localSize < height) {
+                       size[1] = localSize;
+               } else {
+                       size[1] = height - offsety;
+               }
+               for (offsetx = 0 ; offsetx < width ; offsetx+=localSize) {
+                       if (offsetx+localSize < width) {
+                               size[0] = localSize;
+                       } else {
+                               size[0] = width - offsetx;
+                       }
+                       offset[0] = offsetx;
+
+                       error = clSetKernelArg(kernel, offsetIndex, sizeof(cl_int2), &offset);
+                       if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); }
+                       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);
+               }
+       }
+}
+
+cl_kernel NodeOperation::COM_clCreateKernel(cl_program program, const char *kernelname, list<cl_kernel> *clKernelsToCleanUp) 
+{
+       cl_int error;
+       cl_kernel kernel = clCreateKernel(program, kernelname, &error)  ;
+       if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  
+       }
+       else {
+               if (clKernelsToCleanUp) clKernelsToCleanUp->push_back(kernel);
+       }
+       return kernel;
+       
+}
index 73ba5b472d736c806ef3206fe6b46a5d6ea2a4c5..2219907b0c807fdbde4f344888da1bebf6b1bb38 100644 (file)
@@ -139,8 +139,10 @@ public:
          * @param rect the rectangle of the chunk (location and size)
          * @param chunkNumber the chunkNumber to be calculated
          * @param memoryBuffers all input MemoryBuffer's needed
+         * @param outputBuffer the outputbuffer to write to
          */
-       virtual void executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer** memoryBuffers) {}
+       virtual void executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, 
+                                        unsigned int chunkNumber, MemoryBuffer** memoryBuffers, MemoryBuffer* outputBuffer) {}
 
        /**
          * @brief custom handle to add new tasks to the OpenCL command queue in order to execute a chunk on an GPUDevice
@@ -207,9 +209,9 @@ public:
        /**
          * @brief get the render priority of this node.
          * @note only applicable for output operations like ViewerOperation
-         * @return [0:9] 9 is highest priority
+         * @return CompositorPriority
          */
-       virtual const int getRenderPriority() const {return 0;}
+       virtual const CompositorPriority getRenderPriority() const {return COM_PRIORITY_LOW;}
 
        /**
          * @brief can this NodeOperation be scheduled on an OpenCLDevice
@@ -242,6 +244,13 @@ protected:
          */
        void setOpenCL(bool openCL) {this->openCL = openCL;}
 
+       static cl_mem COM_clAttachMemoryBufferToKernelParameter(cl_context context, cl_kernel kernel, int parameterIndex, int offsetIndex, list<cl_mem> *cleanup, MemoryBuffer **inputMemoryBuffers, SocketReader* reader);
+       static void COM_clAttachMemoryBufferOffsetToKernelParameter(cl_kernel kernel, int offsetIndex, MemoryBuffer *memoryBuffers);
+       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);
+       cl_kernel COM_clCreateKernel(cl_program program, const char* kernelname, list<cl_kernel> *clKernelsToCleanUp);
 
 };
 
index 692b96f40b309de237afddacc80be605767f6e1b..e6d3789b06d28f035cc79a76c00acbf0677992ed 100644 (file)
@@ -56,10 +56,10 @@ void OpenCLDevice::execute(WorkPackage *work)
        MemoryBuffer ** inputBuffers = executionGroup->getInputBuffersOpenCL(chunkNumber);
        MemoryBuffer * outputBuffer = executionGroup->allocateOutputBuffer(chunkNumber, &rect);
 
-       executionGroup->getOutputNodeOperation()->executeOpenCLRegion(this->context, this->program, this->queue, &rect, chunkNumber, inputBuffers);
+       executionGroup->getOutputNodeOperation()->executeOpenCLRegion(this->context, this->program, this->queue, &rect, 
+                                                                     chunkNumber, inputBuffers, outputBuffer);
+
+       delete outputBuffer;
        
        executionGroup->finalizeChunkExecution(chunkNumber, inputBuffers);
-       if (outputBuffer != NULL) {
-               outputBuffer->setCreatedState();
-       }
 }
index 80b91b2364c0ff6e341c634b87b0f509b6e35fc9..172107f720b94e220d5301f77e9d0a3bed4849db 100644 (file)
@@ -28,7 +28,7 @@
 #include "COM_OpenCLDevice.h"
 #include "OCL_opencl.h"
 #include "stdio.h"
-#include "COM_OpenCLKernels.cl.cpp"
+#include "COM_OpenCLKernels.cl.h"
 #include "BKE_global.h"
 
 #if COM_CURRENT_THREADING_MODEL == COM_TM_NOTHREAD
@@ -260,7 +260,7 @@ void WorkScheduler::initialize()
                if (totalNumberOfDevices > 0) {
                        context = clCreateContext(NULL, totalNumberOfDevices, cldevices, clContextError, NULL, &error);
                        if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  }
-                       program = clCreateProgramWithSource(context, 1, &sourcecode, 0, &error);
+                       program = clCreateProgramWithSource(context, 1, &clkernelstoh_COM_OpenCLKernels_cl, 0, &error);
                        error = clBuildProgram(program, totalNumberOfDevices, cldevices, 0, 0, 0);
                        if (error != CL_SUCCESS) { 
                                cl_int error2;
index b209e36dd4815d557d2ecf9d08b0423c76d0b2ba..d9cf2c2fef0a4856691441bf6e04746c5beecac3 100644 (file)
@@ -55,7 +55,7 @@ void BlurNode::convertToOperations(ExecutionSystem *graph, CompositorContext * c
                this->getInputSocket(1)->relinkConnections(operationfgb->getInputSocket(1), 1, graph);
                this->getOutputSocket(0)->relinkConnections(operationfgb->getOutputSocket(0));
                graph->addOperation(operationfgb);
-               addPreviewOperation(graph, operationfgb->getOutputSocket(), 5);
+               addPreviewOperation(graph, operationfgb->getOutputSocket());
        }
        else if (!data->bokeh) {
                GaussianXBlurOperation *operationx = new GaussianXBlurOperation();
@@ -71,7 +71,7 @@ void BlurNode::convertToOperations(ExecutionSystem *graph, CompositorContext * c
                graph->addOperation(operationy);
                addLink(graph, operationx->getOutputSocket(), operationy->getInputSocket(0));
                addLink(graph, operationx->getInputSocket(1)->getConnection()->getFromSocket(), operationy->getInputSocket(1));
-               addPreviewOperation(graph, operationy->getOutputSocket(), 5);
+               addPreviewOperation(graph, operationy->getOutputSocket());
 
                if (!connectedSizeSocket) {
                        operationx->setSize(size);
@@ -86,7 +86,7 @@ void BlurNode::convertToOperations(ExecutionSystem *graph, CompositorContext * c
                operation->setQuality(quality);
                graph->addOperation(operation);
                this->getOutputSocket(0)->relinkConnections(operation->getOutputSocket());
-               addPreviewOperation(graph, operation->getOutputSocket(), 5);
+               addPreviewOperation(graph, operation->getOutputSocket());
 
                if (!connectedSizeSocket) {
                        operation->setSize(size);
index d6f4f58fe7000541bb054f64e5821a4ab2acde34..abae1b8889087cb36ed89f86eecbcc4e2f0d809e 100644 (file)
@@ -41,9 +41,9 @@ void BokehBlurNode::convertToOperations(ExecutionSystem *graph, CompositorContex
        if (this->getInputSocket(2)->isConnected()) {
                VariableSizeBokehBlurOperation *operation = new VariableSizeBokehBlurOperation();
                ConvertDepthToRadiusOperation *converter = new ConvertDepthToRadiusOperation();
-               converter->setfStop(4.0f);
+               converter->setfStop(this->getbNode()->custom3);
                converter->setCameraObject(camob);
-               operation->setMaxBlur(16);
+               operation->setMaxBlur((int)this->getbNode()->custom4);
                operation->setQuality(context->getQuality());
                this->getInputSocket(0)->relinkConnections(operation->getInputSocket(0), 0, graph);
                this->getInputSocket(1)->relinkConnections(operation->getInputSocket(1), 1, graph);
index 35511d213f51b1241af6d55dd35e11de63c38861..f498fa11e30052545f868c9704be78f09aca5411 100644 (file)
@@ -35,5 +35,5 @@ void BokehImageNode::convertToOperations(ExecutionSystem *graph, CompositorConte
        this->getOutputSocket(0)->relinkConnections(operation->getOutputSocket(0));
        graph->addOperation(operation);
        operation->setData((NodeBokehImage*)this->getbNode()->storage);
-       addPreviewOperation(graph, operation->getOutputSocket(0), 9);
+       addPreviewOperation(graph, operation->getOutputSocket(0));
 }
index dbe5b9936dc60dfa387e2cd8b12907c93d2cf31d..f1d5b8d39cc195acbfe005d89a449e912ce24772 100644 (file)
@@ -82,7 +82,7 @@ void ChannelMatteNode::convertToOperations(ExecutionSystem *graph, CompositorCon
        graph->addOperation(operationAlpha);
 
        addLink(graph, operation->getOutputSocket(), operationAlpha->getInputSocket(1));
-       addPreviewOperation(graph, operationAlpha->getOutputSocket(), 9);
+       addPreviewOperation(graph, operationAlpha->getOutputSocket());
 
        if (outputSocketImage->isConnected()) {
                outputSocketImage->relinkConnections(operationAlpha->getOutputSocket());
index dd3b3855e3f22dcb4d710ed33e9327c37d5db804..82059ed84935968b69b9cc84e4d85500b8d0b4e0 100644 (file)
@@ -63,7 +63,7 @@ void ChromaMatteNode::convertToOperations(ExecutionSystem *graph, CompositorCont
        addLink(graph, operation->getOutputSocket(), operationAlpha->getInputSocket(1));
 
        graph->addOperation(operationAlpha);
-       addPreviewOperation(graph, operationAlpha->getOutputSocket(), 9);
+       addPreviewOperation(graph, operationAlpha->getOutputSocket());
 
        if (outputSocketImage->isConnected()) {
                outputSocketImage->relinkConnections(operationAlpha->getOutputSocket());
index d7cde21a9843a0f9df9bd26cd4606ab4e741a243..0d331ed9b057837bcf9d9db8923759c4186d43cf 100644 (file)
@@ -31,16 +31,32 @@ ColorCurveNode::ColorCurveNode(bNode *editorNode): Node(editorNode)
 
 void ColorCurveNode::convertToOperations(ExecutionSystem *graph, CompositorContext * context)
 {
-       ColorCurveOperation *operation = new ColorCurveOperation();
-
-       this->getInputSocket(0)->relinkConnections(operation->getInputSocket(0), 0, graph);
-       this->getInputSocket(1)->relinkConnections(operation->getInputSocket(1), 1, graph);
-       this->getInputSocket(2)->relinkConnections(operation->getInputSocket(2), 2, graph);
-       this->getInputSocket(3)->relinkConnections(operation->getInputSocket(3), 3, graph);
-
-       this->getOutputSocket(0)->relinkConnections(operation->getOutputSocket());
-
-       operation->setCurveMapping((CurveMapping*)this->getbNode()->storage);
-
-       graph->addOperation(operation);
+       if (this->getInputSocket(2)->isConnected() || this->getInputSocket(3)->isConnected()) {
+               ColorCurveOperation *operation = new ColorCurveOperation();
+       
+               this->getInputSocket(0)->relinkConnections(operation->getInputSocket(0), 0, graph);
+               this->getInputSocket(1)->relinkConnections(operation->getInputSocket(1), 1, graph);
+               this->getInputSocket(2)->relinkConnections(operation->getInputSocket(2), 2, graph);
+               this->getInputSocket(3)->relinkConnections(operation->getInputSocket(3), 3, graph);
+       
+               this->getOutputSocket(0)->relinkConnections(operation->getOutputSocket());
+       
+               operation->setCurveMapping((CurveMapping*)this->getbNode()->storage);
+       
+               graph->addOperation(operation);
+       } else {
+               ConstantLevelColorCurveOperation *operation = new ConstantLevelColorCurveOperation();
+       
+               this->getInputSocket(0)->relinkConnections(operation->getInputSocket(0), 0, graph);
+               this->getInputSocket(1)->relinkConnections(operation->getInputSocket(1), 1, graph);
+               bNodeSocketValueRGBA *val = (bNodeSocketValueRGBA*)this->getInputSocket(2)->getbNodeSocket()->default_value;
+               operation->setBlackLevel(val->value);
+               val = (bNodeSocketValueRGBA*)this->getInputSocket(3)->getbNodeSocket()->default_value;
+               operation->setWhiteLevel(val->value);
+               this->getOutputSocket(0)->relinkConnections(operation->getOutputSocket());
+       
+               operation->setCurveMapping((CurveMapping*)this->getbNode()->storage);
+       
+               graph->addOperation(operation);
+       }
 }
index 860d1a011944432eaef5995566e76a8b30311167..ad117e1ca2c1717284437021661bd43d1b172987 100644 (file)
@@ -60,7 +60,7 @@ void ColorMatteNode::convertToOperations(ExecutionSystem *graph, CompositorConte
        addLink(graph, operationRGBToHSV_Image->getInputSocket(0)->getConnection()->getFromSocket(), operationAlpha->getInputSocket(0));
        addLink(graph, operation->getOutputSocket(), operationAlpha->getInputSocket(1));
        graph->addOperation(operationAlpha);
-       addPreviewOperation(graph, operationAlpha->getOutputSocket(), 9);
+       addPreviewOperation(graph, operationAlpha->getOutputSocket());
 
        if (outputSocketImage->isConnected()) {
                outputSocketImage->relinkConnections(operationAlpha->getOutputSocket());
index 57821e7fe27034bbfc8a6f7e125493e91452acf7..e2cc34bb6ce30d96e853295434e15a84437de55d 100644 (file)
@@ -39,6 +39,6 @@ void CompositorNode::convertToOperations(ExecutionSystem *graph, CompositorConte
                imageSocket->relinkConnections(colourAlphaProg->getInputSocket(0));
                alphaSocket->relinkConnections(colourAlphaProg->getInputSocket(1));
                graph->addOperation(colourAlphaProg);
-               addPreviewOperation(graph, colourAlphaProg->getInputSocket(0), 5);
+               addPreviewOperation(graph, colourAlphaProg->getInputSocket(0));
        }
 }
index c26fb4e5c5dffffe2ae683713418848fb7079272..596fefff77c8ac7a160ba06b3d41b22ea32a0741 100644 (file)
@@ -49,5 +49,5 @@ void DifferenceMatteNode::convertToOperations(ExecutionSystem *graph, Compositor
        addLink(graph, operationSet->getOutputSocket(), operation->getInputSocket(1));
        outputSocketImage->relinkConnections(operation->getOutputSocket());
        graph->addOperation(operation);
-       addPreviewOperation(graph, operation->getOutputSocket(), 5);
+       addPreviewOperation(graph, operation->getOutputSocket());
 }
index d7b4e481ec281a8644a91a087f70d7acaa04a624..20a55ae195ca0d5a618ba977e6b91bdc1292d066 100644 (file)
@@ -52,7 +52,7 @@ void DistanceMatteNode::convertToOperations(ExecutionSystem *graph, CompositorCo
        addLink(graph, operation->getOutputSocket(), operationAlpha->getInputSocket(1));
 
        graph->addOperation(operationAlpha);
-       addPreviewOperation(graph, operationAlpha->getOutputSocket(), 9);
+       addPreviewOperation(graph, operationAlpha->getOutputSocket());
 
        if (outputSocketImage->isConnected()) {
                outputSocketImage->relinkConnections(operationAlpha->getOutputSocket());
index bdba69dc47d24e412597961a351f2e76c69cc067..7700bceb4abbf50b68a8067042c4e28e3c7c5bc8 100644 (file)
@@ -76,7 +76,7 @@ void FilterNode::convertToOperations(ExecutionSystem *graph, CompositorContext *
        inputImageSocket->relinkConnections(operation->getInputSocket(0), 1, graph);
        inputSocket->relinkConnections(operation->getInputSocket(1), 0, graph);
        outputSocket->relinkConnections(operation->getOutputSocket());
-       addPreviewOperation(graph, operation->getOutputSocket(0), 5);
+       addPreviewOperation(graph, operation->getOutputSocket(0));
        
        graph->addOperation(operation);
 }
index 7f14b06813671e312209a32288bac991f6f0faa2..cfd530173a98409fbe0258b37a285d2315eafcaa 100644 (file)
@@ -105,7 +105,7 @@ void ImageNode::convertToOperations(ExecutionSystem *graph, CompositorContext *
                                                        break;
                                                        }
                                                        if (index == 0 && operation) {
-                                                               addPreviewOperation(graph, operation->getOutputSocket(), 9);
+                                                               addPreviewOperation(graph, operation->getOutputSocket());
                                                        }
                                                }
                                        }
@@ -123,7 +123,7 @@ void ImageNode::convertToOperations(ExecutionSystem *graph, CompositorContext *
                        operation->setImageUser(imageuser);
                        operation->setFramenumber(framenumber);
                        graph->addOperation(operation);
-                       addPreviewOperation(graph, operation->getOutputSocket(), 9);
+                       addPreviewOperation(graph, operation->getOutputSocket());
                }
                
                if (numberOfOutputs > 1) {
index eb78657f3c4ec99691389e57d8e37a554fe8be92..3797621610676a55aac362a5f9327904712d204c 100644 (file)
@@ -53,7 +53,7 @@ void LuminanceMatteNode::convertToOperations(ExecutionSystem *graph, CompositorC
        addLink(graph, rgbToYUV->getInputSocket(0)->getConnection()->getFromSocket(), operation->getInputSocket(0));
        addLink(graph, operationSet->getOutputSocket(), operation->getInputSocket(1));
        graph->addOperation(operation);
-       addPreviewOperation(graph, operation->getOutputSocket(), 9);
+       addPreviewOperation(graph, operation->getOutputSocket());
 
        if (outputSocketImage->isConnected()) {
                outputSocketImage->relinkConnections(operation->getOutputSocket());
index 86ca5ebc2378adffdd7c14bb4f56b72f5883a79f..42e32a4e55ed764ca1ed017673dac8f43cff8122 100644 (file)
@@ -125,7 +125,7 @@ void MixNode::convertToOperations(ExecutionSystem *graph, CompositorContext * co
        color1Socket->relinkConnections(convertProg->getInputSocket(1), 1, graph);
        color2Socket->relinkConnections(convertProg->getInputSocket(2), 2, graph);
        outputSocket->relinkConnections(convertProg->getOutputSocket(0));
-       addPreviewOperation(graph, convertProg->getOutputSocket(0), 5);
+       addPreviewOperation(graph, convertProg->getOutputSocket(0));
        
        convertProg->getInputSocket(2)->setResizeMode(color2Socket->getResizeMode());
        
index 75831130936b0dbb1da07d2303687b62c9681551..eac581dc903dc7b34ca9b06c958126db2bbae095 100644 (file)
@@ -62,7 +62,7 @@ void MovieClipNode::convertToOperations(ExecutionSystem *graph, CompositorContex
                converter->setFromColorProfile(IB_PROFILE_LINEAR_RGB);
                converter->setToColorProfile(IB_PROFILE_SRGB);
                addLink(graph, operation->getOutputSocket(), converter->getInputSocket(0));
-               addPreviewOperation(graph, converter->getOutputSocket(), 9);
+               addPreviewOperation(graph, converter->getOutputSocket());
                if (outputMovieClip->isConnected()) {
                        outputMovieClip->relinkConnections(converter->getOutputSocket());
                }
@@ -72,7 +72,7 @@ void MovieClipNode::convertToOperations(ExecutionSystem *graph, CompositorContex
                }
        }
        else {
-               addPreviewOperation(graph, operation->getOutputSocket(), 9);
+               addPreviewOperation(graph, operation->getOutputSocket());
                if (outputMovieClip->isConnected()) {
                        outputMovieClip->relinkConnections(operation->getOutputSocket());
                }
index cc060e9f7cde41a3bf0a6a2fd5372b848791d531..ca18ea5fbf75021b551b0d89c62a359ac6b1f539 100644 (file)
@@ -59,7 +59,7 @@ void OutputFileNode::convertToOperations(ExecutionSystem *graph, CompositorConte
                                input->relinkConnections(outputOperation->getInputSocket(i));
                        }
                }
-               if (hasConnections) addPreviewOperation(graph, outputOperation->getInputSocket(0), 5);
+               if (hasConnections) addPreviewOperation(graph, outputOperation->getInputSocket(0));
                
                graph->addOperation(outputOperation);
        }
@@ -81,7 +81,7 @@ void OutputFileNode::convertToOperations(ExecutionSystem *graph, CompositorConte
                                input->relinkConnections(outputOperation->getInputSocket(0));
                                graph->addOperation(outputOperation);
                                if (!previewAdded) {
-                                       addPreviewOperation(graph, outputOperation->getInputSocket(0), 5);
+                                       addPreviewOperation(graph, outputOperation->getInputSocket(0));
                                        previewAdded = true;
                                }
                        }
index 4e99db090e12258dbb7b3b232b2324e110b2f4a0..8216205b9252e976d36e19aea089b19a2a0d10a5 100644 (file)
@@ -63,7 +63,7 @@ void RenderLayersNode::testSocketConnection(ExecutionSystem *system, int outputS
                outputSocket->relinkConnections(operation->getOutputSocket());
                system->addOperation(operation);
                if (outputSocketNumber == 0) { // only do for image socket if connected
-                       addPreviewOperation(system, operation->getOutputSocket(), 9);
+                       addPreviewOperation(system, operation->getOutputSocket());
                }
        }
        else {
@@ -71,7 +71,7 @@ void RenderLayersNode::testSocketConnection(ExecutionSystem *system, int outputS
                        system->addOperation(operation);
                        operation->setScene(scene);
                        operation->setLayerId(layerId);
-                       addPreviewOperation(system, operation->getOutputSocket(), 9);
+                       addPreviewOperation(system, operation->getOutputSocket());
                }
                else {
                        delete operation;
index 9f9efbd8fe571368453c583209d4925b7bd7a487..bf434c164c0205223c8bc89a88012c46a0260f9a 100644 (file)
@@ -45,7 +45,7 @@ void SplitViewerNode::convertToOperations(ExecutionSystem *graph, CompositorCont
                splitViewerOperation->setXSplit(!this->getbNode()->custom2);
                image1Socket->relinkConnections(splitViewerOperation->getInputSocket(0), 0, graph);
                image2Socket->relinkConnections(splitViewerOperation->getInputSocket(1), 1, graph);
-               addPreviewOperation(graph, splitViewerOperation->getInputSocket(0), 0);
+               addPreviewOperation(graph, splitViewerOperation->getInputSocket(0));
                graph->addOperation(splitViewerOperation);
        }
 }
index be8bb623f4c3bc07270be0d9c9b00445d1ad19ca..fe8a8e2250ef75312d4ab14efcaf5950121df4cd 100644 (file)
@@ -39,7 +39,7 @@ void TextureNode::convertToOperations(ExecutionSystem *system, CompositorContext
        operation->setTexture(texture);
        operation->setScene(context->getScene());
        system->addOperation(operation);
-       addPreviewOperation(system, operation->getOutputSocket(), 9);
+       addPreviewOperation(system, operation->getOutputSocket());
 
        if (this->getOutputSocket(0)->isConnected()) {
                TextureAlphaOperation *alphaOperation = new TextureAlphaOperation();
index 679589a7ce12c86150e85fec2f8464cd0e798c66..f5dab52d02169934a66e1feeae8d59a13ec4b14d 100644 (file)
@@ -51,6 +51,6 @@ void ViewerNode::convertToOperations(ExecutionSystem *graph, CompositorContext *
                imageSocket->relinkConnections(viewerOperation->getInputSocket(0), 0, graph);
                alphaSocket->relinkConnections(viewerOperation->getInputSocket(1));
                graph->addOperation(viewerOperation);
-               addPreviewOperation(graph, viewerOperation->getInputSocket(0), 0);
+               addPreviewOperation(graph, viewerOperation->getInputSocket(0));
        }
 }
index 1050fc571940e41fc9b0b2ecc07f3a383b18696b..c48f3b0a291f76737a284acf810fe70e7f31ef8f 100644 (file)
@@ -34,8 +34,9 @@ BokehBlurOperation::BokehBlurOperation() : NodeOperation()
        this->addInputSocket(COM_DT_VALUE);
        this->addOutputSocket(COM_DT_COLOR);
        this->setComplex(true);
+       this->setOpenCL(true);
 
-       this->size = .01;
+       this->size = 1.0f;
 
        this->inputProgram = NULL;
        this->inputBokehProgram = NULL;
@@ -90,7 +91,7 @@ void BokehBlurOperation::executePixel(float *color, int x, int y, MemoryBuffer *
                int bufferwidth = inputBuffer->getWidth();
                int bufferstartx = inputBuffer->getRect()->xmin;
                int bufferstarty = inputBuffer->getRect()->ymin;
-               int pixelSize = this->size*this->getWidth();
+               int pixelSize = this->size*this->getWidth()/100.0f;
 
                int miny = y - pixelSize;
                int maxy = y + pixelSize;
@@ -142,10 +143,10 @@ bool BokehBlurOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe
        rcti newInput;
        rcti bokehInput;
 
-       newInput.xmax = input->xmax + (size*this->getWidth());
-       newInput.xmin = input->xmin - (size*this->getWidth());
-       newInput.ymax = input->ymax + (size*this->getWidth());
-       newInput.ymin = input->ymin - (size*this->getWidth());
+       newInput.xmax = input->xmax + (size*this->getWidth()/100.0f);
+       newInput.xmin = input->xmin - (size*this->getWidth()/100.0f);
+       newInput.ymax = input->ymax + (size*this->getWidth()/100.0f);
+       newInput.ymin = input->ymin - (size*this->getWidth()/100.0f);
 
        NodeOperation *operation = getInputOperation(1);
        bokehInput.xmax = operation->getWidth();
@@ -165,3 +166,27 @@ bool BokehBlurOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe
        }
        return false;
 }
+
+static cl_kernel kernel = 0;
+void BokehBlurOperation::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 (!kernel) {
+               kernel = COM_clCreateKernel(program, "bokehBlurKernel", NULL);
+       }
+       cl_int radius = this->getWidth()*this->size/100.0f;
+       cl_int step = this->getStep();
+       
+       COM_clAttachMemoryBufferToKernelParameter(context, kernel, 0, -1, clMemToCleanUp, inputMemoryBuffers, this->inputBoundingBoxReader);
+       COM_clAttachMemoryBufferToKernelParameter(context, kernel, 1,  4, clMemToCleanUp, inputMemoryBuffers, this->inputProgram);
+       COM_clAttachMemoryBufferToKernelParameter(context, kernel, 2,  -1, clMemToCleanUp, inputMemoryBuffers, this->inputBokehProgram);
+       COM_clAttachOutputMemoryBufferToKernelParameter(kernel, 3, clOutputBuffer);
+       COM_clAttachMemoryBufferOffsetToKernelParameter(kernel, 5, outputMemoryBuffer);
+       clSetKernelArg(kernel, 6, sizeof(cl_int), &radius);
+       clSetKernelArg(kernel, 7, sizeof(cl_int), &step);
+       COM_clAttachSizeToKernelParameter(kernel, 8);
+       
+       COM_clEnqueueRange(queue, kernel, outputMemoryBuffer, 9);       
+}
index ce14faa859626520f47cb4561e4619d2dd5c1abf..3cdd995b1df5fc38833dd4e829abe7fb5a2cba21 100644 (file)
@@ -56,5 +56,7 @@ public:
        bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output);
 
        void setSize(float size) {this->size = size;}
+       
+       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);
 };
 #endif
index 8aee54013b19c70eeda57baa2081612b50b6c1a4..069bbde5e735717ef91aca379f91ff039c2a5ca6 100644 (file)
@@ -28,6 +28,7 @@ extern "C" {
        #include "BKE_colortools.h"
 #ifdef __cplusplus
 }
+#include "MEM_guardedalloc.h"
 #endif
 
 ColorCurveOperation::ColorCurveOperation(): CurveBaseOperation()
@@ -59,6 +60,9 @@ void ColorCurveOperation::initExecution()
 
 void ColorCurveOperation::executePixel(float *color, float x, float y, PixelSampler sampler, MemoryBuffer *inputBuffers[])
 {
+       CurveMapping* cumap = this->curveMapping;
+       CurveMapping* workingCopy = (CurveMapping*)MEM_dupallocN(cumap);
+       
        float black[4];
        float white[4];
        float fac[4];
@@ -67,13 +71,13 @@ void ColorCurveOperation::executePixel(float *color, float x, float y, PixelSamp
        this->inputBlackProgram->read(black, x, y, sampler, inputBuffers);
        this->inputWhiteProgram->read(white, x, y, sampler, inputBuffers);
 
-       curvemapping_set_black_white(this->curveMapping, black, white);
+       curvemapping_set_black_white(workingCopy, black, white);
 
        this->inputFacProgram->read(fac, x, y, sampler, inputBuffers);
        this->inputImageProgram->read(image, x, y, sampler, inputBuffers);
 
        if (fac[0] >= 1.0)
-               curvemapping_evaluate_premulRGBF(this->curveMapping, color, image);
+               curvemapping_evaluate_premulRGBF(workingCopy, color, image);
        else if (*fac<=0.0) {
                color[0] = image[0];
                color[1] = image[1];
@@ -81,12 +85,13 @@ void ColorCurveOperation::executePixel(float *color, float x, float y, PixelSamp
        }
        else {
                float col[4], mfac = 1.0f-*fac;
-               curvemapping_evaluate_premulRGBF(this->curveMapping, col, image);
+               curvemapping_evaluate_premulRGBF(workingCopy, col, image);
                color[0] = mfac*image[0] + *fac*col[0];
                color[1] = mfac*image[1] + *fac*col[1];
                color[2] = mfac*image[2] + *fac*col[2];
        }
        color[3] = image[3];
+       MEM_freeN(workingCopy);
 }
 
 void ColorCurveOperation::deinitExecution()
@@ -97,3 +102,61 @@ void ColorCurveOperation::deinitExecution()
        this->inputWhiteProgram = NULL;
        curvemapping_premultiply(this->curveMapping, 1);
 }
+
+
+// Constant level curve mapping
+
+ConstantLevelColorCurveOperation::ConstantLevelColorCurveOperation(): CurveBaseOperation()
+{
+       this->addInputSocket(COM_DT_VALUE);
+       this->addInputSocket(COM_DT_COLOR);
+       this->addOutputSocket(COM_DT_COLOR);
+
+       this->inputFacProgram = NULL;
+       this->inputImageProgram = NULL;
+
+       this->setResolutionInputSocketIndex(1);
+}
+void ConstantLevelColorCurveOperation::initExecution()
+{
+       CurveBaseOperation::initExecution();
+       this->inputFacProgram = this->getInputSocketReader(0);
+       this->inputImageProgram = this->getInputSocketReader(1);
+
+       curvemapping_premultiply(this->curveMapping, 0);
+
+       curvemapping_set_black_white(this->curveMapping, this->black, this->white);
+}
+
+void ConstantLevelColorCurveOperation::executePixel(float *color, float x, float y, PixelSampler sampler, MemoryBuffer *inputBuffers[])
+{
+       float fac[4];
+       float image[4];
+
+
+       this->inputFacProgram->read(fac, x, y, sampler, inputBuffers);
+       this->inputImageProgram->read(image, x, y, sampler, inputBuffers);
+
+       if (fac[0] >= 1.0)
+               curvemapping_evaluate_premulRGBF(this->curveMapping, color, image);
+       else if (*fac<=0.0) {
+               color[0] = image[0];
+               color[1] = image[1];
+               color[2] = image[2];
+       }
+       else {
+               float col[4], mfac = 1.0f-*fac;
+               curvemapping_evaluate_premulRGBF(this->curveMapping, col, image);
+               color[0] = mfac*image[0] + *fac*col[0];
+               color[1] = mfac*image[1] + *fac*col[1];
+               color[2] = mfac*image[2] + *fac*col[2];
+       }
+       color[3] = image[3];
+}
+
+void ConstantLevelColorCurveOperation::deinitExecution()
+{
+       this->inputFacProgram = NULL;
+       this->inputImageProgram = NULL;
+       curvemapping_premultiply(this->curveMapping, 1);
+}
index 15f9fd25e81a0d5fdb2130ee6e2c7b24d50d93e5..6ce5befb14a098187bdc500891f75aec2dec65ba 100644 (file)
@@ -53,4 +53,37 @@ public:
          */
        void deinitExecution();
 };
+
+class ConstantLevelColorCurveOperation : public CurveBaseOperation {
+private:
+       /**
+         * Cached reference to the inputProgram
+         */
+       SocketReader * inputFacProgram;
+       SocketReader * inputImageProgram;
+       float black[3];
+       float white[3];
+       
+public:
+       ConstantLevelColorCurveOperation();
+       
+       /**
+         * the inner loop of this program
+         */
+       void executePixel(float *color, float x, float y, PixelSampler sampler, MemoryBuffer *inputBuffers[]);
+       
+       /**
+         * Initialize the execution
+         */
+       void initExecution();
+       
+       /**
+         * Deinitialize the execution
+         */
+       void deinitExecution();
+       
+       void setBlackLevel(float black[3]) {this->black[0] =black[0];this->black[1] =black[1];this->black[2] =black[2]; }
+       void setWhiteLevel(float white[3]) {this->white[0] =white[0];this->white[1] =white[1];this->white[2] =white[2]; }
+};
+
 #endif
index 41d43f896bb265dc752b67f508a82e3f93236858..13cb4f2832437f1d6dc9547342431eeec9f63b14 100644 (file)
@@ -63,7 +63,7 @@ public:
        bool isOutputOperation(bool rendering) const {return true;}
        void initExecution();
        void deinitExecution();
-       const int getRenderPriority() const {return 7;}
+       const CompositorPriority getRenderPriority() const {return COM_PRIORITY_MEDIUM;}
        void determineResolution(unsigned int resolution[], unsigned int preferredResolution[]);
 };
 #endif
index 40932e54bc7e20432b8db77c3e429014e44b2178..aeccfcab8b5e276bfcded7dacb5580df7fc20c70 100644 (file)
@@ -1,10 +1,52 @@
 /// This file contains all opencl kernels for node-operation implementations 
 
-__kernel void testKernel(__global __write_only image2d_t output)
+// Global SAMPLERS
+const sampler_t SAMPLER_NEAREST      = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
+
+__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, 
+                              int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset) 
 {
-       int x = get_global_id(0);
-       int y = get_global_id(1);
-       int2 coords = {x, y}; 
-       float4 color = {0.0f, 1.0f, 0.0f, 1.0f};
+       int2 coords = {get_global_id(0), get_global_id(1)}; 
+       coords += offset;
+       float tempBoundingBox;
+       float4 color = {0.0f,0.0f,0.0f,0.0f};
+       float4 multiplyer = {0.0f,0.0f,0.0f,0.0f};
+       float4 bokeh;
+       const float radius2 = radius*2.0f;
+       const int2 realCoordinate = coords + offsetOutput;
+
+       tempBoundingBox = read_imagef(boundingBox, SAMPLER_NEAREST, coords).s0;
+
+       if (tempBoundingBox > 0.0f) {
+               const int2 bokehImageDim = get_image_dim(bokehImage);
+               const int2 bokehImageCenter = bokehImageDim/2;
+               const int2 minXY = max(realCoordinate - radius, zero);
+               const int2 maxXY = min(realCoordinate + radius, dimension);
+               int nx, ny;
+               
+               float2 uv;
+               int2 inputXy;
+               
+               for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny +=step, inputXy.y+=step) {
+                       uv.y = ((realCoordinate.y-ny)/radius2)*bokehImageDim.y+bokehImageCenter.y;
+                       
+                       for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx +=step, inputXy.x+=step) {
+                               uv.x = ((realCoordinate.x-nx)/radius2)*bokehImageDim.x+bokehImageCenter.x;
+                               bokeh = read_imagef(bokehImage, SAMPLER_NEAREST, uv);
+                               color += bokeh * read_imagef(inputImage, SAMPLER_NEAREST, inputXy);
+                               multiplyer += bokeh;
+                       }
+               }
+               color /= multiplyer;
+               
+       } else {
+               int2 imageCoordinates = realCoordinate - offsetInput;
+               color = read_imagef(inputImage, SAMPLER_NEAREST, imageCoordinates);
+       }
+       
        write_imagef(output, coords, color);
 }
diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl.cpp b/source/blender/compositor/operations/COM_OpenCLKernels.cl.cpp
deleted file mode 100644 (file)
index 1024d46..0000000
+++ /dev/null
@@ -1,15 +0,0 @@
-/// @todo: this source needs to be generated from COM_OpenCLKernels.cl.
-/// not implemented yet. new data to h
-
-const char *sourcecode = "/// This file contains all opencl kernels for node-operation implementations \n" \
-"\n" \
-"__kernel void testKernel(__global __write_only image2d_t output)\n" \
-"{\n" \
-"      int x = get_global_id(0);\n" \
-"      int y = get_global_id(1);\n" \
-"      int2 coords = {x, y}; \n" \
-"      float4 color = {0.0f, 1.0f, 0.0f, 1.0f};\n" \
-"      write_imagef(output, coords, color);\n" \
-"}\n" \
-"\0\n";
-
diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl.h b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h
new file mode 100644 (file)
index 0000000..3cf33c7
--- /dev/null
@@ -0,0 +1,55 @@
+/* clkernelstoh output of file <COM_OpenCLKernels_cl> */
+
+const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all opencl kernels for node-operation implementations\n" \
+"\n" \
+"// Global SAMPLERS\n" \
+"const sampler_t SAMPLER_NEAREST      = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n" \
+"\n" \
+"__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" \
+"                              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" \
+"      coords += offset;\n" \
+"      float tempBoundingBox;\n" \
+"      float4 color = {0.0f,0.0f,0.0f,0.0f};\n" \
+"      float4 multiplyer = {0.0f,0.0f,0.0f,0.0f};\n" \
+"      float4 bokeh;\n" \
+"      const float radius2 = radius*2.0f;\n" \
+"      const int2 realCoordinate = coords + offsetOutput;\n" \
+"\n" \
+"      tempBoundingBox = read_imagef(boundingBox, SAMPLER_NEAREST, coords).s0;\n" \
+"\n" \
+"      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" \
+"              int nx, ny;\n" \
+"\n" \
+"              float2 uv;\n" \
+"              int2 inputXy;\n" \
+"\n" \
+"              for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny +=step, inputXy.y+=step) {\n" \
+"                      uv.y = ((realCoordinate.y-ny)/radius2)*bokehImageDim.y+bokehImageCenter.y;\n" \
+"\n" \
+"                      for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx +=step, inputXy.x+=step) {\n" \
+"                              uv.x = ((realCoordinate.x-nx)/radius2)*bokehImageDim.x+bokehImageCenter.x;\n" \
+"                              bokeh = read_imagef(bokehImage, SAMPLER_NEAREST, uv);\n" \
+"                              color += bokeh * read_imagef(inputImage, SAMPLER_NEAREST, inputXy);\n" \
+"                              multiplyer += bokeh;\n" \
+"                      }\n" \
+"              }\n" \
+"              color /= multiplyer;\n" \
+"\n" \
+"      } else {\n" \
+"              int2 imageCoordinates = realCoordinate - offsetInput;\n" \
+"              color = read_imagef(inputImage, SAMPLER_NEAREST, imageCoordinates);\n" \
+"      }\n" \
+"\n" \
+"      write_imagef(output, coords, color);\n" \
+"}\n" \
+"\0";
index 0e37432ca5ba0bdc5516bd1013419bc7512ae902..9b9fb02346757649c59a54f667f2bcc332acdc08 100644 (file)
@@ -49,7 +49,7 @@ public:
        bool isOutputOperation(bool rendering) const {return true;}
        void initExecution();
        void deinitExecution();
-       const int getRenderPriority() const {return 7;}
+       const CompositorPriority getRenderPriority() const {return COM_PRIORITY_LOW;}
 };
 
 /* extra info for OpenEXR layers */
@@ -83,7 +83,7 @@ public:
        bool isOutputOperation(bool rendering) const {return true;}
        void initExecution();
        void deinitExecution();
-       const int getRenderPriority() const {return 7;}
+       const CompositorPriority getRenderPriority() const {return COM_PRIORITY_LOW;}
 };
 
 #endif
index a7b6fc93b25f18390816cf89b4ef93a2ca895908..4975f13a285fdc72448dd7714333040d7f8cbe92 100644 (file)
@@ -46,7 +46,6 @@ PreviewOperation::PreviewOperation() : NodeOperation()
        this->input = NULL;
        this->divider = 1.0f;
        this->node = NULL;
-       this->priority = 0;
 }
 
 void PreviewOperation::initExecution()
@@ -129,7 +128,7 @@ void PreviewOperation::determineResolution(unsigned int resolution[], unsigned i
        resolution[1] = height;
 }
 
-const int PreviewOperation::getRenderPriority() const
+const CompositorPriority PreviewOperation::getRenderPriority() const
 {
-       return this->priority;
+       return COM_PRIORITY_LOW;
 }
index 8450b7fc556ca139e167df785e0701de8fed41f2..3d1cd38a5ea4dea6f8db0f707499d84ddd152933 100644 (file)
@@ -37,20 +37,18 @@ protected:
        const bNodeTree *tree;
        SocketReader *input;
        float divider;
-       int priority;
 
 public:
        PreviewOperation();
        bool isOutputOperation(bool rendering) const {return true;}
        void initExecution();
        void deinitExecution();
-       const int getRenderPriority() const;
+       const CompositorPriority getRenderPriority() const;
        
        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);
-       void setPriority(int priority) { this->priority = priority; }
 };
 #endif
index 71fc53e8f8e167efe48c8f548bd5bccb9fa30327..809c688195f93bd535cda4a88d8cfe611283e00a 100644 (file)
@@ -88,12 +88,12 @@ void ViewerBaseOperation::deinitExecution()
        this->outputBuffer = NULL;
 }
 
-const int ViewerBaseOperation::getRenderPriority() const
+const CompositorPriority ViewerBaseOperation::getRenderPriority() const
 {
        if (this->isActiveViewerOutput()) {
-               return 8;
+               return COM_PRIORITY_HIGH;
        }
        else {
-               return 0;
+               return COM_PRIORITY_LOW;
        }
 }
index f5f30809f6504969d0c79ed464fa3224e202f1a5..51fa8cecc0d752226dc04ca2509e8faca6b3318c 100644 (file)
@@ -56,7 +56,7 @@ public:
        float getCenterX() { return this->centerX; }
        float getCenterY() { return this->centerY; }
        OrderOfChunks getChunkOrder() { return this->chunkOrder; }
-       const int getRenderPriority() const;
+       const CompositorPriority getRenderPriority() const;
        void setColorManagement(bool doColorManagement) {this->doColorManagement = doColorManagement;}
        void setColorPredivide(bool doColorPredivide) {this->doColorPredivide = doColorPredivide;}
        bool isViewerOperation() {return true;}
index 498add2fc870396af21288c318216eed3bbfa259..222b879645cd0eb1f195e0b4fe0dfdb232fedd56 100644 (file)
@@ -111,10 +111,9 @@ void WriteBufferOperation::executeRegion(rcti *rect, unsigned int tileNumber, Me
        memoryBuffer->setCreatedState();
 }
 
-void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer** inputMemoryBuffers)
+void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer** inputMemoryBuffers, MemoryBuffer* outputBuffer)
 {
-       MemoryBuffer *outputMemoryBuffer = this->getMemoryProxy()->getBuffer();// @todo wrong implementation needs revision
-       float *outputFloatBuffer = outputMemoryBuffer->getBuffer();
+       float *outputFloatBuffer = outputBuffer->getBuffer();
        cl_int error;
        /*
         * 1. create cl_mem from outputbuffer
@@ -125,8 +124,8 @@ void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program pr
         * note: list of cl_mem will be filled by 2, and needs to be cleaned up by 4
         */
        // STEP 1
-       const unsigned int outputBufferWidth = outputMemoryBuffer->getWidth();
-       const unsigned int outputBufferHeight = outputMemoryBuffer->getHeight();
+       const unsigned int outputBufferWidth = outputBuffer->getWidth();
+       const unsigned int outputBufferHeight = outputBuffer->getHeight();
 
        const cl_image_format imageFormat = {
                CL_RGBA,
@@ -141,19 +140,26 @@ void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program pr
        clMemToCleanUp->push_back(clOutputBuffer);
        list<cl_kernel> *clKernelsToCleanUp = new list<cl_kernel>();
 
-       this->input->executeOpenCL(context, program, queue, outputMemoryBuffer, clOutputBuffer, inputMemoryBuffers, clMemToCleanUp, clKernelsToCleanUp);
+       this->input->executeOpenCL(context, program, queue, outputBuffer, clOutputBuffer, inputMemoryBuffers, clMemToCleanUp, clKernelsToCleanUp);
 
        // STEP 3
 
        size_t origin[3] = {0,0,0};
        size_t region[3] = {outputBufferWidth,outputBufferHeight,1};
 
+//     clFlush(queue);
+//     clFinish(queue);
+
        error = clEnqueueBarrier(queue);
        if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  }
        error = clEnqueueReadImage(queue, clOutputBuffer, CL_TRUE, origin, region, 0, 0, outputFloatBuffer, 0, NULL, NULL);
        if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error));  }
+       
+       this->getMemoryProxy()->getBuffer()->copyContentFrom(outputBuffer);
+       
        // STEP 4
 
+       
        while (clMemToCleanUp->size()>0) {
                cl_mem mem = clMemToCleanUp->front();
                error = clReleaseMemObject(mem);
index b17122d68f05bfa5e2e653199f28f21e17f4c854..6f2c49c49bd616c6106346429bfe5308e3236b16 100644 (file)
@@ -46,7 +46,7 @@ public:
        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);
+       void executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer** memoryBuffers, MemoryBuffer* outputBuffer);
 
 };
 #endif
index 1e8541214cea423197d72f6186a234b183c19e3a..d37a2739fc38c17fe6c44e31f2fbea8c18a3970c 100644 (file)
@@ -2294,7 +2294,7 @@ static void node_composit_buts_bokehimage(uiLayout *layout, bContext *UNUSED(C),
 void node_composit_backdrop_viewer(SpaceNode *snode, ImBuf *backdrop, bNode *node, int x, int y)
 {
 //     node_composit_backdrop_canvas(snode, backdrop, node, x, y);
-       if (node->custom1 == 0) { /// @todo: why did we need this one?
+       if (node->custom1 == 0) {
                const float backdropWidth = backdrop->x;
                const float backdropHeight = backdrop->y;
                const float cx  = x + snode->zoom * backdropWidth * node->custom3;
index 59506f31b60f6cf25b8c61156f841e762b38e76c..5be7688d714a5f2778c8f5f95cf3c620528c32c2 100644 (file)
@@ -241,6 +241,14 @@ typedef struct bNodeLink {
 #define NTREE_QUALITY_MEDIUM  1
 #define NTREE_QUALITY_LOW     2
 
+/* tree->chunksize */
+#define NTREE_CHUNCKSIZE_32 32
+#define NTREE_CHUNCKSIZE_64 64
+#define NTREE_CHUNCKSIZE_128 128
+#define NTREE_CHUNCKSIZE_256 256
+#define NTREE_CHUNCKSIZE_512 512
+#define NTREE_CHUNCKSIZE_1024 1024
+
 /* the basis for a Node tree, all links and nodes reside internal here */
 /* only re-usable node trees are in the library though, materials and textures allocate own tree struct */
 typedef struct bNodeTree {
index 562684799b4fe7ba78e8de706aafb89d08d0d17d..65d572c6c118b380075771be65b1cfb8b39368ce 100644 (file)
@@ -71,6 +71,16 @@ EnumPropertyItem node_quality_items[] = {
        {0, NULL, 0, NULL, NULL}
 };
 
+EnumPropertyItem node_chunksize_items[] = {
+    {NTREE_CHUNCKSIZE_32,   "32",     0,    "32x32",     "Chunksize of 32x32"},
+    {NTREE_CHUNCKSIZE_64,   "64",     0,    "64x64",     "Chunksize of 64x64"},
+    {NTREE_CHUNCKSIZE_128,   "128",     0,    "128x128",     "Chunksize of 128x128"},
+    {NTREE_CHUNCKSIZE_256,   "256",     0,    "256x256",     "Chunksize of 256x256"},
+    {NTREE_CHUNCKSIZE_512,   "512",     0,    "512x512",     "Chunksize of 512x512"},
+    {NTREE_CHUNCKSIZE_1024,   "1024",     0,    "1024x1024",     "Chunksize of 1024x1024"},
+       {0, NULL, 0, NULL, NULL}
+};
+
 EnumPropertyItem node_socket_type_items[] = {
        {SOCK_FLOAT,   "VALUE",     0,    "Value",     ""},
        {SOCK_VECTOR,  "VECTOR",    0,    "Vector",    ""},
@@ -3184,6 +3194,25 @@ static void def_cmp_ellipsemask(StructRNA *srna)
        RNA_def_property_update(prop, NC_NODE | NA_EDITED, "rna_Node_update");
 }
 
+static void def_cmp_bokehblur(StructRNA *srna)
+{
+       PropertyRNA *prop;
+       prop = RNA_def_property(srna, "f_stop", PROP_FLOAT, PROP_NONE);
+       RNA_def_property_float_sdna(prop, NULL, "custom3");
+       RNA_def_property_range(prop, 0.0f, 128.0f);
+       RNA_def_property_ui_text(prop, "fStop",
+                                "Amount of focal blur, 128=infinity=perfect focus, half the value doubles "
+                                "the blur radius");
+       RNA_def_property_update(prop, NC_NODE | NA_EDITED, "rna_Node_update");
+       
+       prop = RNA_def_property(srna, "blur_max", PROP_FLOAT, PROP_NONE);
+       RNA_def_property_float_sdna(prop, NULL, "custom4");
+       RNA_def_property_range(prop, 0.0f, 10000.0f);
+       RNA_def_property_ui_text(prop, "Max Blur", "Blur limit, maximum CoC radius");
+       RNA_def_property_update(prop, NC_NODE | NA_EDITED, "rna_Node_update");
+       
+}
+
 static void def_cmp_bokehimage(StructRNA *srna)
 {
        PropertyRNA *prop;
@@ -4029,11 +4058,11 @@ static void rna_def_composite_nodetree(BlenderRNA *brna)
        RNA_def_property_enum_items(prop, node_quality_items);
        RNA_def_property_ui_text(prop, "Edit Quality", "Quality when editing");
 
-       prop = RNA_def_property(srna, "chunk_size", PROP_INT, PROP_NONE);
-       RNA_def_property_int_sdna(prop, NULL, "chunksize");
+       prop = RNA_def_property(srna, "chunk_size", PROP_ENUM, PROP_NONE);
+       RNA_def_property_enum_sdna(prop, NULL, "chunksize");
+       RNA_def_property_enum_items(prop, node_chunksize_items);
        RNA_def_property_ui_text(prop, "Chunksize", "Max size of a tile (smaller values gives better distribution "
                                                    "of multiple threads, but more overhead)");
-       RNA_def_property_range(prop, 32, 1024);
 
        prop = RNA_def_property(srna, "use_opencl", PROP_BOOLEAN, PROP_NONE);
        RNA_def_property_boolean_sdna(prop, NULL, "flag", NTREE_COM_OPENCL);
index d5b33f0d01b990e1b9aa4ca311cc7cd11729ecbf..2b8050adca7a6c2b6071c914446a30c02f729a44 100644 (file)
@@ -163,6 +163,7 @@ DefNode( CompositorNode, CMP_NODE_MOVIEDISTORTION,def_cmp_moviedistortion,"MOVIE
 DefNode( CompositorNode, CMP_NODE_MASK_BOX,       def_cmp_boxmask,        "BOXMASK"        ,BoxMask,          "Box mask",          ""              )
 DefNode( CompositorNode, CMP_NODE_MASK_ELLIPSE,   def_cmp_ellipsemask,    "ELLIPSEMASK"    ,EllipseMask,      "Ellipse mask",      ""              )
 DefNode( CompositorNode, CMP_NODE_BOKEHIMAGE,     def_cmp_bokehimage,     "BOKEHIMAGE"     ,BokehImage,       "Bokeh image",       ""              )
+DefNode( CompositorNode, CMP_NODE_BOKEHBLUR,      def_cmp_bokehblur,      "BOKEHBLUR"      ,BokehBlur,        "Bokeh Blur",        ""              )
 DefNode( CompositorNode, CMP_NODE_SWITCH,         def_cmp_switch,         "SWITCH"         ,Switch,           "Switch",            ""              )
 DefNode( CompositorNode, CMP_NODE_COLORCORRECTION,def_cmp_colorcorrection,"COLORCORRECTION",ColorCorrection,  "ColorCorrection",   ""              )
 DefNode( CompositorNode, CMP_NODE_MASK,           def_cmp_mask,           "MASK",           Mask,             "Mask",              ""              )
index ab679b9242fbcfbe43f6ff628786fa81fa7a5ace..06b6e79444a68b8f4f13d83c138b528bf8166243 100644 (file)
@@ -39,7 +39,7 @@
 static bNodeSocketTemplate cmp_node_bokehblur_in[]= {
        {       SOCK_RGBA, 1, N_("Image"),                      0.8f, 0.8f, 0.8f, 1.0f, 0.0f, 1.0f},
        {       SOCK_RGBA, 1, N_("Bokeh"),                      1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f},
-       {       SOCK_FLOAT, 1, N_("Size"),                      0.01f, 0.0f, 0.0f, 0.0f, 0.0f, 0.5f},
+       {       SOCK_FLOAT, 1, N_("Size"),                      1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 10.0f},
        {       SOCK_FLOAT, 1, N_("Bounding box"),      1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f},
        {       -1, 0, ""       }
 };
@@ -49,6 +49,12 @@ static bNodeSocketTemplate cmp_node_bokehblur_out[]= {
        {       -1, 0, ""       }
 };
 
+static void node_composit_init_bokehblur(bNodeTree *UNUSED(ntree), bNode* node, bNodeTemplate *UNUSED(ntemp))
+{
+       node->custom3 = 4.0f;
+       node->custom4 = 16.0f;
+}
+
 void register_node_type_cmp_bokehblur(bNodeTreeType *ttype)
 {
        static bNodeType ntype;
@@ -56,5 +62,7 @@ void register_node_type_cmp_bokehblur(bNodeTreeType *ttype)
        node_type_base(ttype, &ntype, CMP_NODE_BOKEHBLUR, "Bokeh Blur", NODE_CLASS_OP_FILTER, NODE_OPTIONS);
        node_type_socket_templates(&ntype, cmp_node_bokehblur_in, cmp_node_bokehblur_out);
        node_type_size(&ntype, 120, 80, 200);
+       node_type_init(&ntype, node_composit_init_bokehblur);
+       
        nodeRegisterType(ttype, &ntype);
 }