diff options
52 files changed, 631 insertions, 188 deletions
diff --git a/release/datafiles/clkernelstoh.py b/release/datafiles/clkernelstoh.py new file mode 100755 index 00000000000..8fb5d4e6bae --- /dev/null +++ b/release/datafiles/clkernelstoh.py @@ -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() diff --git a/source/blender/compositor/COM_defines.h b/source/blender/compositor/COM_defines.h index df807091cb8..f87265c50f5 100644 --- a/source/blender/compositor/COM_defines.h +++ b/source/blender/compositor/COM_defines.h @@ -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. diff --git a/source/blender/compositor/intern/COM_ExecutionGroup.cpp b/source/blender/compositor/intern/COM_ExecutionGroup.cpp index 0ebf9af207b..44b3c8dafbb 100644 --- a/source/blender/compositor/intern/COM_ExecutionGroup.cpp +++ b/source/blender/compositor/intern/COM_ExecutionGroup.cpp @@ -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; diff --git a/source/blender/compositor/intern/COM_ExecutionGroup.h b/source/blender/compositor/intern/COM_ExecutionGroup.h index cbdc9bb1787..416a78eb8b8 100644 --- a/source/blender/compositor/intern/COM_ExecutionGroup.h +++ b/source/blender/compositor/intern/COM_ExecutionGroup.h @@ -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 diff --git a/source/blender/compositor/intern/COM_ExecutionSystem.cpp b/source/blender/compositor/intern/COM_ExecutionSystem.cpp index 96d2a6f4434..8c0b37a0685 100644 --- a/source/blender/compositor/intern/COM_ExecutionSystem.cpp +++ b/source/blender/compositor/intern/COM_ExecutionSystem.cpp @@ -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; diff --git a/source/blender/compositor/intern/COM_ExecutionSystem.h b/source/blender/compositor/intern/COM_ExecutionSystem.h index 85fec8b6145..510e58ba1bb 100644 --- a/source/blender/compositor/intern/COM_ExecutionSystem.h +++ b/source/blender/compositor/intern/COM_ExecutionSystem.h @@ -141,6 +141,11 @@ private: //methods /** * find all execution group with output nodes */ + void findOutputExecutionGroup(vector<ExecutionGroup*> *result, CompositorPriority priority) const; + + /** + * find all execution group with output nodes + */ void findOutputExecutionGroup(vector<ExecutionGroup*> *result) const; public: @@ -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 diff --git a/source/blender/compositor/intern/COM_Node.cpp b/source/blender/compositor/intern/COM_Node.cpp index ba5e21d53ae..264725b4b2c 100644 --- a/source/blender/compositor/intern/COM_Node.cpp +++ b/source/blender/compositor/intern/COM_Node.cpp @@ -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); } } diff --git a/source/blender/compositor/intern/COM_Node.h b/source/blender/compositor/intern/COM_Node.h index 2666d0a6980..23744adf642 100644 --- a/source/blender/compositor/intern/COM_Node.h +++ b/source/blender/compositor/intern/COM_Node.h @@ -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); diff --git a/source/blender/compositor/intern/COM_NodeOperation.cpp b/source/blender/compositor/intern/COM_NodeOperation.cpp index fae652e39d7..650e4af5ae0 100644 --- a/source/blender/compositor/intern/COM_NodeOperation.cpp +++ b/source/blender/compositor/intern/COM_NodeOperation.cpp @@ -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; + +} diff --git a/source/blender/compositor/intern/COM_NodeOperation.h b/source/blender/compositor/intern/COM_NodeOperation.h index 73ba5b472d7..2219907b0c8 100644 --- a/source/blender/compositor/intern/COM_NodeOperation.h +++ b/source/blender/compositor/intern/COM_NodeOperation.h @@ -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); }; diff --git a/source/blender/compositor/intern/COM_OpenCLDevice.cpp b/source/blender/compositor/intern/COM_OpenCLDevice.cpp index 692b96f40b3..e6d3789b06d 100644 --- a/source/blender/compositor/intern/COM_OpenCLDevice.cpp +++ b/source/blender/compositor/intern/COM_OpenCLDevice.cpp @@ -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(); - } } diff --git a/source/blender/compositor/intern/COM_WorkScheduler.cpp b/source/blender/compositor/intern/COM_WorkScheduler.cpp index 80b91b2364c..172107f720b 100644 --- a/source/blender/compositor/intern/COM_WorkScheduler.cpp +++ b/source/blender/compositor/intern/COM_WorkScheduler.cpp @@ -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; diff --git a/source/blender/compositor/nodes/COM_BlurNode.cpp b/source/blender/compositor/nodes/COM_BlurNode.cpp index b209e36dd48..d9cf2c2fef0 100644 --- a/source/blender/compositor/nodes/COM_BlurNode.cpp +++ b/source/blender/compositor/nodes/COM_BlurNode.cpp @@ -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); diff --git a/source/blender/compositor/nodes/COM_BokehBlurNode.cpp b/source/blender/compositor/nodes/COM_BokehBlurNode.cpp index d6f4f58fe70..abae1b88890 100644 --- a/source/blender/compositor/nodes/COM_BokehBlurNode.cpp +++ b/source/blender/compositor/nodes/COM_BokehBlurNode.cpp @@ -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); diff --git a/source/blender/compositor/nodes/COM_BokehImageNode.cpp b/source/blender/compositor/nodes/COM_BokehImageNode.cpp index 35511d213f5..f498fa11e30 100644 --- a/source/blender/compositor/nodes/COM_BokehImageNode.cpp +++ b/source/blender/compositor/nodes/COM_BokehImageNode.cpp @@ -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)); } diff --git a/source/blender/compositor/nodes/COM_ChannelMatteNode.cpp b/source/blender/compositor/nodes/COM_ChannelMatteNode.cpp index dbe5b9936dc..f1d5b8d39cc 100644 --- a/source/blender/compositor/nodes/COM_ChannelMatteNode.cpp +++ b/source/blender/compositor/nodes/COM_ChannelMatteNode.cpp @@ -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()); diff --git a/source/blender/compositor/nodes/COM_ChromaMatteNode.cpp b/source/blender/compositor/nodes/COM_ChromaMatteNode.cpp index dd3b3855e3f..82059ed8493 100644 --- a/source/blender/compositor/nodes/COM_ChromaMatteNode.cpp +++ b/source/blender/compositor/nodes/COM_ChromaMatteNode.cpp @@ -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()); diff --git a/source/blender/compositor/nodes/COM_ColorCurveNode.cpp b/source/blender/compositor/nodes/COM_ColorCurveNode.cpp index d7cde21a984..0d331ed9b05 100644 --- a/source/blender/compositor/nodes/COM_ColorCurveNode.cpp +++ b/source/blender/compositor/nodes/COM_ColorCurveNode.cpp @@ -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); + } } diff --git a/source/blender/compositor/nodes/COM_ColorMatteNode.cpp b/source/blender/compositor/nodes/COM_ColorMatteNode.cpp index 860d1a01194..ad117e1ca2c 100644 --- a/source/blender/compositor/nodes/COM_ColorMatteNode.cpp +++ b/source/blender/compositor/nodes/COM_ColorMatteNode.cpp @@ -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()); diff --git a/source/blender/compositor/nodes/COM_CompositorNode.cpp b/source/blender/compositor/nodes/COM_CompositorNode.cpp index 57821e7fe27..e2cc34bb6ce 100644 --- a/source/blender/compositor/nodes/COM_CompositorNode.cpp +++ b/source/blender/compositor/nodes/COM_CompositorNode.cpp @@ -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)); } } diff --git a/source/blender/compositor/nodes/COM_DifferenceMatteNode.cpp b/source/blender/compositor/nodes/COM_DifferenceMatteNode.cpp index c26fb4e5c5d..596fefff77c 100644 --- a/source/blender/compositor/nodes/COM_DifferenceMatteNode.cpp +++ b/source/blender/compositor/nodes/COM_DifferenceMatteNode.cpp @@ -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()); } diff --git a/source/blender/compositor/nodes/COM_DistanceMatteNode.cpp b/source/blender/compositor/nodes/COM_DistanceMatteNode.cpp index d7b4e481ec2..20a55ae195c 100644 --- a/source/blender/compositor/nodes/COM_DistanceMatteNode.cpp +++ b/source/blender/compositor/nodes/COM_DistanceMatteNode.cpp @@ -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()); diff --git a/source/blender/compositor/nodes/COM_FilterNode.cpp b/source/blender/compositor/nodes/COM_FilterNode.cpp index bdba69dc47d..7700bceb4ab 100644 --- a/source/blender/compositor/nodes/COM_FilterNode.cpp +++ b/source/blender/compositor/nodes/COM_FilterNode.cpp @@ -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); } diff --git a/source/blender/compositor/nodes/COM_ImageNode.cpp b/source/blender/compositor/nodes/COM_ImageNode.cpp index 7f14b068136..cfd530173a9 100644 --- a/source/blender/compositor/nodes/COM_ImageNode.cpp +++ b/source/blender/compositor/nodes/COM_ImageNode.cpp @@ -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) { diff --git a/source/blender/compositor/nodes/COM_LuminanceMatteNode.cpp b/source/blender/compositor/nodes/COM_LuminanceMatteNode.cpp index eb78657f3c4..37976216106 100644 --- a/source/blender/compositor/nodes/COM_LuminanceMatteNode.cpp +++ b/source/blender/compositor/nodes/COM_LuminanceMatteNode.cpp @@ -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()); diff --git a/source/blender/compositor/nodes/COM_MixNode.cpp b/source/blender/compositor/nodes/COM_MixNode.cpp index 86ca5ebc237..42e32a4e55e 100644 --- a/source/blender/compositor/nodes/COM_MixNode.cpp +++ b/source/blender/compositor/nodes/COM_MixNode.cpp @@ -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()); diff --git a/source/blender/compositor/nodes/COM_MovieClipNode.cpp b/source/blender/compositor/nodes/COM_MovieClipNode.cpp index 75831130936..eac581dc903 100644 --- a/source/blender/compositor/nodes/COM_MovieClipNode.cpp +++ b/source/blender/compositor/nodes/COM_MovieClipNode.cpp @@ -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()); } diff --git a/source/blender/compositor/nodes/COM_OutputFileNode.cpp b/source/blender/compositor/nodes/COM_OutputFileNode.cpp index cc060e9f7cd..ca18ea5fbf7 100644 --- a/source/blender/compositor/nodes/COM_OutputFileNode.cpp +++ b/source/blender/compositor/nodes/COM_OutputFileNode.cpp @@ -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; } } diff --git a/source/blender/compositor/nodes/COM_RenderLayersNode.cpp b/source/blender/compositor/nodes/COM_RenderLayersNode.cpp index 4e99db090e1..8216205b925 100644 --- a/source/blender/compositor/nodes/COM_RenderLayersNode.cpp +++ b/source/blender/compositor/nodes/COM_RenderLayersNode.cpp @@ -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; diff --git a/source/blender/compositor/nodes/COM_SplitViewerNode.cpp b/source/blender/compositor/nodes/COM_SplitViewerNode.cpp index 9f9efbd8fe5..bf434c164c0 100644 --- a/source/blender/compositor/nodes/COM_SplitViewerNode.cpp +++ b/source/blender/compositor/nodes/COM_SplitViewerNode.cpp @@ -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); } } diff --git a/source/blender/compositor/nodes/COM_TextureNode.cpp b/source/blender/compositor/nodes/COM_TextureNode.cpp index be8bb623f4c..fe8a8e2250e 100644 --- a/source/blender/compositor/nodes/COM_TextureNode.cpp +++ b/source/blender/compositor/nodes/COM_TextureNode.cpp @@ -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(); diff --git a/source/blender/compositor/nodes/COM_ViewerNode.cpp b/source/blender/compositor/nodes/COM_ViewerNode.cpp index 679589a7ce1..f5dab52d021 100644 --- a/source/blender/compositor/nodes/COM_ViewerNode.cpp +++ b/source/blender/compositor/nodes/COM_ViewerNode.cpp @@ -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)); } } diff --git a/source/blender/compositor/operations/COM_BokehBlurOperation.cpp b/source/blender/compositor/operations/COM_BokehBlurOperation.cpp index 1050fc57194..c48f3b0a291 100644 --- a/source/blender/compositor/operations/COM_BokehBlurOperation.cpp +++ b/source/blender/compositor/operations/COM_BokehBlurOperation.cpp @@ -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); +} diff --git a/source/blender/compositor/operations/COM_BokehBlurOperation.h b/source/blender/compositor/operations/COM_BokehBlurOperation.h index ce14faa8596..3cdd995b1df 100644 --- a/source/blender/compositor/operations/COM_BokehBlurOperation.h +++ b/source/blender/compositor/operations/COM_BokehBlurOperation.h @@ -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 diff --git a/source/blender/compositor/operations/COM_ColorCurveOperation.cpp b/source/blender/compositor/operations/COM_ColorCurveOperation.cpp index 8aee54013b1..069bbde5e73 100644 --- a/source/blender/compositor/operations/COM_ColorCurveOperation.cpp +++ b/source/blender/compositor/operations/COM_ColorCurveOperation.cpp @@ -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); +} diff --git a/source/blender/compositor/operations/COM_ColorCurveOperation.h b/source/blender/compositor/operations/COM_ColorCurveOperation.h index 15f9fd25e81..6ce5befb14a 100644 --- a/source/blender/compositor/operations/COM_ColorCurveOperation.h +++ b/source/blender/compositor/operations/COM_ColorCurveOperation.h @@ -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 diff --git a/source/blender/compositor/operations/COM_CompositorOperation.h b/source/blender/compositor/operations/COM_CompositorOperation.h index 41d43f896bb..13cb4f28324 100644 --- a/source/blender/compositor/operations/COM_CompositorOperation.h +++ b/source/blender/compositor/operations/COM_CompositorOperation.h @@ -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 diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl b/source/blender/compositor/operations/COM_OpenCLKernels.cl index 40932e54bc7..aeccfcab8b5 100644 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl +++ b/source/blender/compositor/operations/COM_OpenCLKernels.cl @@ -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 index 1024d460044..00000000000 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl.cpp +++ /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 index 00000000000..3cf33c75272 --- /dev/null +++ b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h @@ -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"; diff --git a/source/blender/compositor/operations/COM_OutputFileOperation.h b/source/blender/compositor/operations/COM_OutputFileOperation.h index 0e37432ca5b..9b9fb023467 100644 --- a/source/blender/compositor/operations/COM_OutputFileOperation.h +++ b/source/blender/compositor/operations/COM_OutputFileOperation.h @@ -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 diff --git a/source/blender/compositor/operations/COM_PreviewOperation.cpp b/source/blender/compositor/operations/COM_PreviewOperation.cpp index a7b6fc93b25..4975f13a285 100644 --- a/source/blender/compositor/operations/COM_PreviewOperation.cpp +++ b/source/blender/compositor/operations/COM_PreviewOperation.cpp @@ -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; } diff --git a/source/blender/compositor/operations/COM_PreviewOperation.h b/source/blender/compositor/operations/COM_PreviewOperation.h index 8450b7fc556..3d1cd38a5ea 100644 --- a/source/blender/compositor/operations/COM_PreviewOperation.h +++ b/source/blender/compositor/operations/COM_PreviewOperation.h @@ -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 diff --git a/source/blender/compositor/operations/COM_ViewerBaseOperation.cpp b/source/blender/compositor/operations/COM_ViewerBaseOperation.cpp index 71fc53e8f8e..809c688195f 100644 --- a/source/blender/compositor/operations/COM_ViewerBaseOperation.cpp +++ b/source/blender/compositor/operations/COM_ViewerBaseOperation.cpp @@ -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; } } diff --git a/source/blender/compositor/operations/COM_ViewerBaseOperation.h b/source/blender/compositor/operations/COM_ViewerBaseOperation.h index f5f30809f65..51fa8cecc0d 100644 --- a/source/blender/compositor/operations/COM_ViewerBaseOperation.h +++ b/source/blender/compositor/operations/COM_ViewerBaseOperation.h @@ -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;} diff --git a/source/blender/compositor/operations/COM_WriteBufferOperation.cpp b/source/blender/compositor/operations/COM_WriteBufferOperation.cpp index 498add2fc87..222b879645c 100644 --- a/source/blender/compositor/operations/COM_WriteBufferOperation.cpp +++ b/source/blender/compositor/operations/COM_WriteBufferOperation.cpp @@ -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); diff --git a/source/blender/compositor/operations/COM_WriteBufferOperation.h b/source/blender/compositor/operations/COM_WriteBufferOperation.h index b17122d68f0..6f2c49c49bd 100644 --- a/source/blender/compositor/operations/COM_WriteBufferOperation.h +++ b/source/blender/compositor/operations/COM_WriteBufferOperation.h @@ -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 diff --git a/source/blender/editors/space_node/drawnode.c b/source/blender/editors/space_node/drawnode.c index 1e8541214ce..d37a2739fc3 100644 --- a/source/blender/editors/space_node/drawnode.c +++ b/source/blender/editors/space_node/drawnode.c @@ -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; diff --git a/source/blender/makesdna/DNA_node_types.h b/source/blender/makesdna/DNA_node_types.h index 59506f31b60..5be7688d714 100644 --- a/source/blender/makesdna/DNA_node_types.h +++ b/source/blender/makesdna/DNA_node_types.h @@ -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 { diff --git a/source/blender/makesrna/intern/rna_nodetree.c b/source/blender/makesrna/intern/rna_nodetree.c index 562684799b4..65d572c6c11 100644 --- a/source/blender/makesrna/intern/rna_nodetree.c +++ b/source/blender/makesrna/intern/rna_nodetree.c @@ -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); diff --git a/source/blender/makesrna/intern/rna_nodetree_types.h b/source/blender/makesrna/intern/rna_nodetree_types.h index d5b33f0d01b..2b8050adca7 100644 --- a/source/blender/makesrna/intern/rna_nodetree_types.h +++ b/source/blender/makesrna/intern/rna_nodetree_types.h @@ -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", "" ) diff --git a/source/blender/nodes/composite/nodes/node_composite_bokehblur.c b/source/blender/nodes/composite/nodes/node_composite_bokehblur.c index ab679b9242f..06b6e79444a 100644 --- a/source/blender/nodes/composite/nodes/node_composite_bokehblur.c +++ b/source/blender/nodes/composite/nodes/node_composite_bokehblur.c @@ -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); } |