diff options
Diffstat (limited to 'source/blender/compositor')
40 files changed, 483 insertions, 163 deletions
diff --git a/source/blender/compositor/CMakeLists.txt b/source/blender/compositor/CMakeLists.txt index 22f72270734..bc2aeaefc83 100644 --- a/source/blender/compositor/CMakeLists.txt +++ b/source/blender/compositor/CMakeLists.txt @@ -97,6 +97,9 @@ set(SRC intern/COM_CompositorContext.h intern/COM_ChannelInfo.cpp intern/COM_ChannelInfo.h + intern/COM_SingleThreadedNodeOperation.cpp + intern/COM_SingleThreadedNodeOperation.h + operations/COM_QualityStepHelper.h operations/COM_QualityStepHelper.cpp diff --git a/source/blender/compositor/intern/COM_ExecutionGroup.cpp b/source/blender/compositor/intern/COM_ExecutionGroup.cpp index e46b4934217..7a53af7f58c 100644 --- a/source/blender/compositor/intern/COM_ExecutionGroup.cpp +++ b/source/blender/compositor/intern/COM_ExecutionGroup.cpp @@ -54,6 +54,7 @@ ExecutionGroup::ExecutionGroup() this->numberOfChunks = 0; this->initialized = false; this->openCL = false; + this->singleThreaded = false; this->chunksFinished = 0; } @@ -100,6 +101,7 @@ void ExecutionGroup::addOperation(ExecutionSystem *system, NodeOperation *operat if (!operation->isBufferOperation()) { this->complex = operation->isComplex(); this->openCL = operation->isOpenCL(); + this->singleThreaded = operation->isSingleThreaded(); this->initialized = true; } this->operations.push_back(operation); @@ -191,10 +193,17 @@ void ExecutionGroup::determineResolution(unsigned int resolution[]) void ExecutionGroup::determineNumberOfChunks() { - const float chunkSizef = this->chunkSize; - this->numberOfXChunks = ceil(this->width / chunkSizef); - this->numberOfYChunks = ceil(this->height / chunkSizef); - this->numberOfChunks = this->numberOfXChunks * this->numberOfYChunks; + if (singleThreaded) { + this->numberOfXChunks = 1; + this->numberOfYChunks = 1; + this->numberOfChunks = 1; + } + else { + const float chunkSizef = this->chunkSize; + this->numberOfXChunks = ceil(this->width / chunkSizef); + this->numberOfYChunks = ceil(this->height / chunkSizef); + this->numberOfChunks = this->numberOfXChunks * this->numberOfYChunks; + } } /** @@ -435,9 +444,14 @@ void ExecutionGroup::finalizeChunkExecution(int chunkNumber, MemoryBuffer** memo inline void ExecutionGroup::determineChunkRect(rcti *rect, const unsigned int xChunk, const unsigned int yChunk ) const { - const unsigned int minx = xChunk * chunkSize; - const unsigned int miny = yChunk * chunkSize; - BLI_init_rcti(rect, minx, min(minx + this->chunkSize, this->width), miny, min(miny + this->chunkSize, this->height)); + if (singleThreaded) { + BLI_init_rcti(rect, 0, this->width, 0, this->height); + } + else { + const unsigned int minx = xChunk * chunkSize; + const unsigned int miny = yChunk * chunkSize; + BLI_init_rcti(rect, minx, min(minx + this->chunkSize, this->width), miny, min(miny + this->chunkSize, this->height)); + } } void ExecutionGroup::determineChunkRect(rcti *rect, const unsigned int chunkNumber) const @@ -462,6 +476,9 @@ MemoryBuffer *ExecutionGroup::allocateOutputBuffer(int chunkNumber, rcti *rect) bool ExecutionGroup::scheduleAreaWhenPossible(ExecutionSystem * graph, rcti *area) { + if (singleThreaded) { + return scheduleChunkWhenPossible(graph, 0, 0); + } // find all chunks inside the rect // determine minxchunk, minychunk, maxxchunk, maxychunk where x and y are chunknumbers diff --git a/source/blender/compositor/intern/COM_ExecutionGroup.h b/source/blender/compositor/intern/COM_ExecutionGroup.h index 416a78eb8b8..1698890cc34 100644 --- a/source/blender/compositor/intern/COM_ExecutionGroup.h +++ b/source/blender/compositor/intern/COM_ExecutionGroup.h @@ -63,10 +63,6 @@ class Device; class ExecutionGroup { private: // fields - /** - * @brief unique identifier of this node. - */ - string id; /** * @brief list of operations in this ExecutionGroup @@ -121,6 +117,11 @@ private: bool openCL; /** + * @brief Is this Execution group SingleThreaded + */ + bool singleThreaded; + + /** * @brief what is the maximum number field of all ReadBufferOperation in this ExecutionGroup. * @note this is used to construct the MemoryBuffers that will be passed during execution. */ @@ -233,18 +234,7 @@ private: public: // constructors ExecutionGroup(); - - /** - * @brief set the id of this ExecutionGroup - * @param id - */ - void setId(string id) {this->id = id;} - - /** - * @brief return the id of this ExecutionGroup - */ - const string getId() const {return this->id;} - + // methods /** * @brief check to see if a NodeOperation is already inside this execution group diff --git a/source/blender/compositor/intern/COM_ExecutionSystem.cpp b/source/blender/compositor/intern/COM_ExecutionSystem.cpp index 1056c6d3f65..9681996c74d 100644 --- a/source/blender/compositor/intern/COM_ExecutionSystem.cpp +++ b/source/blender/compositor/intern/COM_ExecutionSystem.cpp @@ -124,6 +124,7 @@ void ExecutionSystem::execute() for (index = 0 ; index < this->operations.size() ; index ++) { NodeOperation * operation = this->operations[index]; + operation->setbNodeTree(this->context.getbNodeTree()); operation->initExecution(); } for (index = 0 ; index < this->groups.size() ; index ++) { @@ -153,7 +154,7 @@ void ExecutionSystem::execute() void ExecutionSystem::executeGroups(CompositorPriority priority) { - int index; + unsigned int index; vector<ExecutionGroup*> executionGroups; this->findOutputExecutionGroup(&executionGroups, priority); @@ -166,6 +167,7 @@ void ExecutionSystem::executeGroups(CompositorPriority priority) void ExecutionSystem::addOperation(NodeOperation *operation) { ExecutionSystemHelper::addOperation(this->operations, operation); +// operation->setBTree } void ExecutionSystem::addReadWriteBufferOperations(NodeOperation *operation) diff --git a/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp b/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp index 75be8df74de..d5ca2ec619a 100644 --- a/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp +++ b/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp @@ -65,7 +65,7 @@ Node *ExecutionSystemHelper::addbNodeTree(ExecutionSystem &system, int nodes_sta } /* Expand group nodes */ - for (int i=nodes_start; i < nodes.size(); ++i) { + for (unsigned int i=nodes_start; i < nodes.size(); ++i) { Node *execnode = nodes[i]; if (execnode->isGroupNode()) { GroupNode * groupNode = (GroupNode*)execnode; diff --git a/source/blender/compositor/intern/COM_NodeOperation.cpp b/source/blender/compositor/intern/COM_NodeOperation.cpp index 650e4af5ae0..148ad48ba3a 100644 --- a/source/blender/compositor/intern/COM_NodeOperation.cpp +++ b/source/blender/compositor/intern/COM_NodeOperation.cpp @@ -34,6 +34,7 @@ NodeOperation::NodeOperation() this->width = 0; this->height = 0; this->openCL = false; + this->btree = NULL; } void NodeOperation::determineResolution(unsigned int resolution[], unsigned int preferredResolution[]) @@ -74,10 +75,22 @@ void NodeOperation::initMutex() { BLI_mutex_init(&mutex); } + +void NodeOperation::lockMutex() +{ + BLI_mutex_lock(&mutex); +} + +void NodeOperation::unlockMutex() +{ + BLI_mutex_unlock(&mutex); +} + void NodeOperation::deinitMutex() { BLI_mutex_end(&mutex); } + void NodeOperation::deinitExecution() { } @@ -196,14 +209,15 @@ void NodeOperation::COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, size_t size[2]; cl_int2 offset; - for (offsety = 0 ; offsety < height; offsety+=localSize) { + bool breaked = false; + for (offsety = 0 ; offsety < height && (!breaked); offsety+=localSize) { offset[1] = offsety; if (offsety+localSize < height) { size[1] = localSize; } else { size[1] = height - offsety; } - for (offsetx = 0 ; offsetx < width ; offsetx+=localSize) { + for (offsetx = 0 ; offsetx < width && (!breaked) ; offsetx+=localSize) { if (offsetx+localSize < width) { size[0] = localSize; } else { @@ -216,6 +230,9 @@ void NodeOperation::COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, error = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, size, 0, 0, 0, NULL); if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } clFlush(queue); + if (isBreaked()) { + breaked = false; + } } } } diff --git a/source/blender/compositor/intern/COM_NodeOperation.h b/source/blender/compositor/intern/COM_NodeOperation.h index 3f536fb3f2d..e56ac7bd51f 100644 --- a/source/blender/compositor/intern/COM_NodeOperation.h +++ b/source/blender/compositor/intern/COM_NodeOperation.h @@ -77,6 +77,11 @@ private: * @see NodeOperation.getMutex retrieve a pointer to this mutex. */ ThreadMutex mutex; + + /** + * @brief reference to the editing bNodeTree only used for break callback + */ + const bNodeTree *btree; public: /** @@ -119,9 +124,10 @@ public: * for all other operations this will result in false. */ virtual int isBufferOperation() {return false;} + virtual int isSingleThreaded() {return false;} + void setbNodeTree(const bNodeTree * tree) {this->btree = tree;} virtual void initExecution(); - void initMutex(); /** * @brief when a chunk is executed by a CPUDevice, this method is called @@ -161,7 +167,6 @@ public: */ virtual void executeOpenCL(cl_context context,cl_program program, cl_command_queue queue, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer** inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, list<cl_kernel> *clKernelsToCleanUp) {} virtual void deinitExecution(); - void deinitMutex(); bool isResolutionSet() { return this->width != 0 && height != 0; @@ -236,6 +241,11 @@ public: virtual bool isViewerOperation() {return false;} virtual bool isPreviewOperation() {return false;} + + inline bool isBreaked() { + return btree->test_break(btree->tbh); + } + protected: NodeOperation(); @@ -244,7 +254,11 @@ protected: SocketReader *getInputSocketReader(unsigned int inputSocketindex); NodeOperation *getInputOperation(unsigned int inputSocketindex); - inline ThreadMutex *getMutex() {return &this->mutex;} + void deinitMutex(); + void initMutex(); + void lockMutex(); + void unlockMutex(); + /** * @brief set whether this operation is complex @@ -264,7 +278,7 @@ protected: static void COM_clAttachOutputMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, cl_mem clOutputMemoryBuffer); void COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex); static void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer* outputMemoryBuffer); - static void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex); + void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex); cl_kernel COM_clCreateKernel(cl_program program, const char* kernelname, list<cl_kernel> *clKernelsToCleanUp); }; diff --git a/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.cpp b/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.cpp new file mode 100644 index 00000000000..9ea90809de4 --- /dev/null +++ b/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.cpp @@ -0,0 +1,60 @@ +/* + * Copyright 2011, Blender Foundation. + * + * This program is free software; you can redistribute it and/or + * modify it under the terms of the GNU General Public License + * as published by the Free Software Foundation; either version 2 + * of the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program; if not, write to the Free Software Foundation, + * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. + * + * Contributor: + * Jeroen Bakker + * Monique Dewanchand + */ + +#include "COM_SingleThreadedNodeOperation.h" + +SingleThreadedNodeOperation::SingleThreadedNodeOperation(): NodeOperation() +{ + this->cachedInstance = NULL; + setComplex(true); +} + +void SingleThreadedNodeOperation::initExecution() +{ + initMutex(); +} + +void SingleThreadedNodeOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data) +{ + this->cachedInstance->read(color, x, y); +} + +void SingleThreadedNodeOperation::deinitExecution() +{ + deinitMutex(); + if (this->cachedInstance) { + delete cachedInstance; + this->cachedInstance = NULL; + } +} +void *SingleThreadedNodeOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers) +{ + if (this->cachedInstance) return this->cachedInstance; + + lockMutex(); + if (this->cachedInstance == NULL) { + // + this->cachedInstance = createMemoryBuffer(rect, memoryBuffers); + } + unlockMutex(); + return this->cachedInstance; +} diff --git a/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.h b/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.h new file mode 100644 index 00000000000..ace48365752 --- /dev/null +++ b/source/blender/compositor/intern/COM_SingleThreadedNodeOperation.h @@ -0,0 +1,60 @@ +/* + * Copyright 2011, Blender Foundation. + * + * This program is free software; you can redistribute it and/or + * modify it under the terms of the GNU General Public License + * as published by the Free Software Foundation; either version 2 + * of the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program; if not, write to the Free Software Foundation, + * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. + * + * Contributor: + * Jeroen Bakker + * Monique Dewanchand + */ + +#ifndef _COM_SingleThreadedNodeOperation_h +#define _COM_SingleThreadedNodeOperation_h +#include "COM_NodeOperation.h" + +class SingleThreadedNodeOperation : public NodeOperation { +private: + MemoryBuffer *cachedInstance; + +protected: + inline bool isCached() { + return cachedInstance != NULL; + } + +public: + SingleThreadedNodeOperation(); + + /** + * the inner loop of this program + */ + void executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data); + + /** + * Initialize the execution + */ + void initExecution(); + + /** + * Deinitialize the execution + */ + void deinitExecution(); + + void *initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers); + + virtual MemoryBuffer* createMemoryBuffer(rcti *rect, MemoryBuffer **memoryBuffers) = 0; + + int isSingleThreaded() {return true;} +}; +#endif diff --git a/source/blender/compositor/intern/COM_compositor.cpp b/source/blender/compositor/intern/COM_compositor.cpp index e27bff4401e..2bbfd18e7c5 100644 --- a/source/blender/compositor/intern/COM_compositor.cpp +++ b/source/blender/compositor/intern/COM_compositor.cpp @@ -51,7 +51,6 @@ void COM_execute(bNodeTree *editingtree, int rendering) /* set progress bar to 0% and status to init compositing*/ editingtree->progress(editingtree->prh, 0.0); - editingtree->stats_draw(editingtree->sdh, (char*)"Compositing"); /* initialize execution system */ ExecutionSystem *system = new ExecutionSystem(editingtree, rendering); diff --git a/source/blender/compositor/nodes/COM_DilateErodeNode.cpp b/source/blender/compositor/nodes/COM_DilateErodeNode.cpp index 47791956865..0619bb5133e 100644 --- a/source/blender/compositor/nodes/COM_DilateErodeNode.cpp +++ b/source/blender/compositor/nodes/COM_DilateErodeNode.cpp @@ -36,7 +36,7 @@ void DilateErodeNode::convertToOperations(ExecutionSystem *graph, CompositorCont bNode *editorNode = this->getbNode(); if (editorNode->custom1 == CMP_NODE_DILATEERODE_DISTANCE_THRESH) { - DilateErodeDistanceOperation *operation = new DilateErodeDistanceOperation(); + DilateErodeThresholdOperation *operation = new DilateErodeThresholdOperation(); operation->setDistance(editorNode->custom2); operation->setInset(editorNode->custom3); diff --git a/source/blender/compositor/nodes/COM_MuteNode.cpp b/source/blender/compositor/nodes/COM_MuteNode.cpp index d02eb2a0b98..93e352cfede 100644 --- a/source/blender/compositor/nodes/COM_MuteNode.cpp +++ b/source/blender/compositor/nodes/COM_MuteNode.cpp @@ -44,7 +44,7 @@ void MuteNode::reconnect(ExecutionSystem * graph, OutputSocket * output) } } - NodeOperation *operation; + NodeOperation *operation = NULL; switch (output->getDataType()) { case COM_DT_VALUE: { @@ -74,7 +74,6 @@ void MuteNode::reconnect(ExecutionSystem * graph, OutputSocket * output) } /* quiet warnings */ case COM_DT_UNKNOWN: - operation = NULL; break; } diff --git a/source/blender/compositor/operations/COM_AntiAliasOperation.cpp b/source/blender/compositor/operations/COM_AntiAliasOperation.cpp index 62639eeb24a..4cd9552b108 100644 --- a/source/blender/compositor/operations/COM_AntiAliasOperation.cpp +++ b/source/blender/compositor/operations/COM_AntiAliasOperation.cpp @@ -44,7 +44,7 @@ void AntiAliasOperation::initExecution() void AntiAliasOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void * data) { - if (y < 0 || y >= this->height || x < 0 || x >= this->width) { + if (y < 0 || (unsigned int)y >= this->height || x < 0 || (unsigned int)x >= this->width) { color[0] = 0.0f; } else { @@ -85,7 +85,7 @@ bool AntiAliasOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe void *AntiAliasOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers) { if (this->buffer) {return buffer;} - BLI_mutex_lock(getMutex()); + lockMutex(); if (this->buffer == NULL) { MemoryBuffer *tile = (MemoryBuffer*)valueReader->initializeTileData(rect, memoryBuffers); int size = tile->getHeight()*tile->getWidth(); @@ -100,6 +100,6 @@ void *AntiAliasOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu antialias_tagbuf(tile->getWidth(), tile->getHeight(), valuebuffer); this->buffer = valuebuffer; } - BLI_mutex_unlock(getMutex()); + unlockMutex(); return this->buffer; } diff --git a/source/blender/compositor/operations/COM_CalculateMeanOperation.cpp b/source/blender/compositor/operations/COM_CalculateMeanOperation.cpp index a3438cea27b..077d8473f0b 100644 --- a/source/blender/compositor/operations/COM_CalculateMeanOperation.cpp +++ b/source/blender/compositor/operations/COM_CalculateMeanOperation.cpp @@ -72,13 +72,13 @@ bool CalculateMeanOperation::determineDependingAreaOfInterest(rcti *input, ReadB void *CalculateMeanOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers) { - BLI_mutex_lock(getMutex()); + lockMutex(); if (!this->iscalculated) { MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers); calculateMean(tile); this->iscalculated = true; } - BLI_mutex_unlock(getMutex()); + unlockMutex(); return NULL; } diff --git a/source/blender/compositor/operations/COM_CalculateStandardDeviationOperation.cpp b/source/blender/compositor/operations/COM_CalculateStandardDeviationOperation.cpp index 651c6674fdb..dfe1b6aa329 100644 --- a/source/blender/compositor/operations/COM_CalculateStandardDeviationOperation.cpp +++ b/source/blender/compositor/operations/COM_CalculateStandardDeviationOperation.cpp @@ -37,7 +37,7 @@ void CalculateStandardDeviationOperation::executePixel(float *color, int x, int void *CalculateStandardDeviationOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers) { - BLI_mutex_lock(getMutex()); + lockMutex(); if (!this->iscalculated) { MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers); CalculateMeanOperation::calculateMean(tile); @@ -92,6 +92,6 @@ void *CalculateStandardDeviationOperation::initializeTileData(rcti *rect, Memory this->standardDeviation = sqrt(sum / (float)(pixels-1)); this->iscalculated = true; } - BLI_mutex_unlock(getMutex()); + unlockMutex(); return NULL; } diff --git a/source/blender/compositor/operations/COM_CompositorOperation.cpp b/source/blender/compositor/operations/COM_CompositorOperation.cpp index c6e8faaa638..4ce7ffac9ef 100644 --- a/source/blender/compositor/operations/COM_CompositorOperation.cpp +++ b/source/blender/compositor/operations/COM_CompositorOperation.cpp @@ -59,7 +59,7 @@ void CompositorOperation::initExecution() void CompositorOperation::deinitExecution() { - if (tree->test_break && !tree->test_break(tree->tbh)) { + if (isBreaked()) { const Scene * scene = this->scene; Render *re = RE_GetRender(scene->id.name); RenderResult *rr = RE_AcquireResultWrite(re); @@ -118,7 +118,7 @@ void CompositorOperation::executeRegion(rcti *rect, unsigned int tileNumber, Mem buffer[offset+2] = color[2]; buffer[offset+3] = color[3]; offset +=COM_NUMBER_OF_CHANNELS; - if (tree->test_break && tree->test_break(tree->tbh)) { + if (isBreaked()) { breaked = true; } } diff --git a/source/blender/compositor/operations/COM_CompositorOperation.h b/source/blender/compositor/operations/COM_CompositorOperation.h index 13cb4f28324..36099b3eb91 100644 --- a/source/blender/compositor/operations/COM_CompositorOperation.h +++ b/source/blender/compositor/operations/COM_CompositorOperation.h @@ -37,11 +37,6 @@ private: const Scene *scene; /** - * @brief local reference to the node tree - */ - const bNodeTree *tree; - - /** * @brief reference to the output float buffer */ float *outputBuffer; @@ -59,7 +54,6 @@ public: CompositorOperation(); void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer** memoryBuffers); void setScene(const Scene *scene) {this->scene = scene;} - void setbNodeTree(const bNodeTree *tree) {this->tree = tree;} bool isOutputOperation(bool rendering) const {return true;} void initExecution(); void deinitExecution(); diff --git a/source/blender/compositor/operations/COM_DilateErodeOperation.cpp b/source/blender/compositor/operations/COM_DilateErodeOperation.cpp index 7bc49fa695c..bdd7362952a 100644 --- a/source/blender/compositor/operations/COM_DilateErodeOperation.cpp +++ b/source/blender/compositor/operations/COM_DilateErodeOperation.cpp @@ -24,7 +24,7 @@ #include "BLI_math.h" // DilateErode Distance Threshold -DilateErodeDistanceOperation::DilateErodeDistanceOperation(): NodeOperation() +DilateErodeThresholdOperation::DilateErodeThresholdOperation(): NodeOperation() { this->addInputSocket(COM_DT_VALUE); this->addOutputSocket(COM_DT_VALUE); @@ -34,7 +34,7 @@ DilateErodeDistanceOperation::DilateErodeDistanceOperation(): NodeOperation() this->_switch = 0.5f; this->distance = 0.0f; } -void DilateErodeDistanceOperation::initExecution() +void DilateErodeThresholdOperation::initExecution() { this->inputProgram = this->getInputSocketReader(0); if (this->distance < 0.0f) { @@ -53,13 +53,13 @@ void DilateErodeDistanceOperation::initExecution() } } -void *DilateErodeDistanceOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers) +void *DilateErodeThresholdOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers) { void *buffer = inputProgram->initializeTileData(NULL, memoryBuffers); return buffer; } -void DilateErodeDistanceOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data) +void DilateErodeThresholdOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data) { float inputValue[4]; const float sw = this->_switch; @@ -142,12 +142,12 @@ void DilateErodeDistanceOperation::executePixel(float *color, int x, int y, Memo } } -void DilateErodeDistanceOperation::deinitExecution() +void DilateErodeThresholdOperation::deinitExecution() { this->inputProgram = NULL; } -bool DilateErodeDistanceOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output) +bool DilateErodeThresholdOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output) { rcti newInput; @@ -167,6 +167,7 @@ DilateDistanceOperation::DilateDistanceOperation(): NodeOperation() this->setComplex(true); this->inputProgram = NULL; this->distance = 0.0f; + this->setOpenCL(true); } void DilateDistanceOperation::initExecution() { @@ -231,6 +232,28 @@ bool DilateDistanceOperation::determineDependingAreaOfInterest(rcti *input, Read return NodeOperation::determineDependingAreaOfInterest(&newInput, readOperation, output); } + +static cl_kernel dilateKernel = 0; +void DilateDistanceOperation::executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, + MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, + MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, + list<cl_kernel> *clKernelsToCleanUp) +{ + if (!dilateKernel) { + dilateKernel = COM_clCreateKernel(program, "dilateKernel", NULL); + } + cl_int distanceSquared = this->distance*this->distance; + cl_int scope = this->scope; + + COM_clAttachMemoryBufferToKernelParameter(context, dilateKernel, 0, 2, clMemToCleanUp, inputMemoryBuffers, this->inputProgram); + COM_clAttachOutputMemoryBufferToKernelParameter(dilateKernel, 1, clOutputBuffer); + COM_clAttachMemoryBufferOffsetToKernelParameter(dilateKernel, 3, outputMemoryBuffer); + clSetKernelArg(dilateKernel, 4, sizeof(cl_int), &scope); + clSetKernelArg(dilateKernel, 5, sizeof(cl_int), &distanceSquared); + COM_clAttachSizeToKernelParameter(dilateKernel, 6); + COM_clEnqueueRange(queue, dilateKernel, outputMemoryBuffer, 7); +} + // Erode Distance ErodeDistanceOperation::ErodeDistanceOperation() : DilateDistanceOperation() { @@ -268,6 +291,27 @@ void ErodeDistanceOperation::executePixel(float *color, int x, int y, MemoryBuff color[0] = value; } +static cl_kernel erodeKernel = 0; +void ErodeDistanceOperation::executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, + MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, + MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, + list<cl_kernel> *clKernelsToCleanUp) +{ + if (!erodeKernel) { + erodeKernel = COM_clCreateKernel(program, "erodeKernel", NULL); + } + cl_int distanceSquared = this->distance*this->distance; + cl_int scope = this->scope; + + COM_clAttachMemoryBufferToKernelParameter(context, erodeKernel, 0, 2, clMemToCleanUp, inputMemoryBuffers, this->inputProgram); + COM_clAttachOutputMemoryBufferToKernelParameter(erodeKernel, 1, clOutputBuffer); + COM_clAttachMemoryBufferOffsetToKernelParameter(erodeKernel, 3, outputMemoryBuffer); + clSetKernelArg(erodeKernel, 4, sizeof(cl_int), &scope); + clSetKernelArg(erodeKernel, 5, sizeof(cl_int), &distanceSquared); + COM_clAttachSizeToKernelParameter(erodeKernel, 6); + COM_clEnqueueRange(queue, erodeKernel, outputMemoryBuffer, 7); +} + // Dilate step DilateStepOperation::DilateStepOperation(): NodeOperation() { @@ -288,7 +332,7 @@ void *DilateStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB if (this->cached_buffer != NULL) { return this->cached_buffer; } - BLI_mutex_lock(getMutex()); + lockMutex(); if (this->cached_buffer == NULL) { MemoryBuffer *buffer = (MemoryBuffer*)inputProgram->initializeTileData(NULL, memoryBuffers); float *rectf = buffer->convertToValueBuffer(); @@ -327,7 +371,7 @@ void *DilateStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB } this->cached_buffer = rectf; } - BLI_mutex_unlock(getMutex()); + unlockMutex(); return this->cached_buffer; } @@ -374,7 +418,7 @@ void *ErodeStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu if (this->cached_buffer != NULL) { return this->cached_buffer; } - BLI_mutex_lock(getMutex()); + lockMutex(); if (this->cached_buffer == NULL) { MemoryBuffer *buffer = (MemoryBuffer*)inputProgram->initializeTileData(NULL, memoryBuffers); float *rectf = buffer->convertToValueBuffer(); @@ -413,6 +457,6 @@ void *ErodeStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu } this->cached_buffer = rectf; } - BLI_mutex_unlock(getMutex()); + unlockMutex(); return this->cached_buffer; } diff --git a/source/blender/compositor/operations/COM_DilateErodeOperation.h b/source/blender/compositor/operations/COM_DilateErodeOperation.h index 71bbab74a4b..63cee4c333a 100644 --- a/source/blender/compositor/operations/COM_DilateErodeOperation.h +++ b/source/blender/compositor/operations/COM_DilateErodeOperation.h @@ -25,7 +25,7 @@ #include "COM_NodeOperation.h" -class DilateErodeDistanceOperation : public NodeOperation { +class DilateErodeThresholdOperation : public NodeOperation { private: /** * Cached reference to the inputProgram @@ -42,7 +42,7 @@ private: */ int scope; public: - DilateErodeDistanceOperation(); + DilateErodeThresholdOperation(); /** * the inner loop of this program @@ -70,11 +70,11 @@ public: class DilateDistanceOperation : public NodeOperation { private: +protected: /** * Cached reference to the inputProgram */ SocketReader * inputProgram; -protected: float distance; int scope; public: @@ -98,6 +98,11 @@ public: void setDistance(float distance) {this->distance = distance;} bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output); + + void executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, + MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, + MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, + list<cl_kernel> *clKernelsToCleanUp); }; class ErodeDistanceOperation : public DilateDistanceOperation { public: @@ -107,6 +112,11 @@ public: * the inner loop of this program */ void executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data); + + void executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, + MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, + MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, + list<cl_kernel> *clKernelsToCleanUp); }; class DilateStepOperation : public NodeOperation { diff --git a/source/blender/compositor/operations/COM_DoubleEdgeMaskOperation.cpp b/source/blender/compositor/operations/COM_DoubleEdgeMaskOperation.cpp index 40f492b0f10..df04b889200 100644 --- a/source/blender/compositor/operations/COM_DoubleEdgeMaskOperation.cpp +++ b/source/blender/compositor/operations/COM_DoubleEdgeMaskOperation.cpp @@ -23,6 +23,7 @@ #include "COM_DoubleEdgeMaskOperation.h" #include "BLI_math.h" #include "DNA_node_types.h" +#include "MEM_guardedalloc.h" // this part has been copied from the double edge mask // Contributor(s): Peter Larabell. @@ -1215,12 +1216,12 @@ void DoubleEdgeMaskOperation::doDoubleEdgeMask(float *imask, float *omask, float gsz=rsize[2]; // by the do_*EdgeDetection() function. fsz=gsz+isz+osz; // calculate size of pixel index buffer needed - gbuf = new unsigned short[fsz*2]; // allocate edge/gradient pixel index buffer + gbuf = (unsigned short*)MEM_callocN(sizeof (unsigned short)*fsz*2, "DEM"); // allocate edge/gradient pixel index buffer do_createEdgeLocationBuffer(t,rw,lres,res,gbuf,&innerEdgeOffset,&outerEdgeOffset,isz,gsz); do_fillGradientBuffer(rw,res,gbuf,isz,osz,gsz,innerEdgeOffset,outerEdgeOffset); - delete [] gbuf; // free the gradient index buffer + MEM_freeN(gbuf); // free the gradient index buffer } } @@ -1263,7 +1264,7 @@ void *DoubleEdgeMaskOperation::initializeTileData(rcti *rect, MemoryBuffer **mem { if (this->cachedInstance) return this->cachedInstance; - BLI_mutex_lock(getMutex()); + lockMutex(); if (this->cachedInstance == NULL) { MemoryBuffer *innerMask = (MemoryBuffer*)inputInnerMask->initializeTileData(rect, memoryBuffers); MemoryBuffer *outerMask = (MemoryBuffer*)inputOuterMask->initializeTileData(rect, memoryBuffers); @@ -1275,7 +1276,7 @@ void *DoubleEdgeMaskOperation::initializeTileData(rcti *rect, MemoryBuffer **mem delete omask; this->cachedInstance = data; } - BLI_mutex_unlock(getMutex()); + unlockMutex(); return this->cachedInstance; } void DoubleEdgeMaskOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data) diff --git a/source/blender/compositor/operations/COM_FastGaussianBlurOperation.cpp b/source/blender/compositor/operations/COM_FastGaussianBlurOperation.cpp index ad58631f2c1..92e1aace9e2 100644 --- a/source/blender/compositor/operations/COM_FastGaussianBlurOperation.cpp +++ b/source/blender/compositor/operations/COM_FastGaussianBlurOperation.cpp @@ -79,7 +79,7 @@ void FastGaussianBlurOperation::deinitExecution() void *FastGaussianBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers) { - BLI_mutex_lock(this->getMutex()); + lockMutex(); if (!iirgaus) { MemoryBuffer *newBuf = (MemoryBuffer*)this->inputProgram->initializeTileData(rect, memoryBuffers); MemoryBuffer *copy = newBuf->duplicate(); @@ -105,7 +105,7 @@ void *FastGaussianBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **m } this->iirgaus = copy; } - BLI_mutex_unlock(this->getMutex()); + unlockMutex(); return iirgaus; } diff --git a/source/blender/compositor/operations/COM_GlareBaseOperation.cpp b/source/blender/compositor/operations/COM_GlareBaseOperation.cpp index fdfd19a10ae..83132963f0b 100644 --- a/source/blender/compositor/operations/COM_GlareBaseOperation.cpp +++ b/source/blender/compositor/operations/COM_GlareBaseOperation.cpp @@ -23,56 +23,41 @@ #include "COM_GlareBaseOperation.h" #include "BLI_math.h" -GlareBaseOperation::GlareBaseOperation(): NodeOperation() +GlareBaseOperation::GlareBaseOperation(): SingleThreadedNodeOperation() { this->addInputSocket(COM_DT_COLOR); this->addOutputSocket(COM_DT_COLOR); this->settings = NULL; - this->cachedInstance = NULL; - setComplex(true); } void GlareBaseOperation::initExecution() { - initMutex(); + SingleThreadedNodeOperation::initExecution(); this->inputProgram = getInputSocketReader(0); - this->cachedInstance = NULL; -} - -void GlareBaseOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data)\ -{ - float *buffer = (float*) data; - int index = (y*this->getWidth() + x) * COM_NUMBER_OF_CHANNELS; - color[0] = buffer[index]; - color[1] = buffer[index+1]; - color[2] = buffer[index+2]; - color[3] = buffer[index+3]; } void GlareBaseOperation::deinitExecution() { - deinitMutex(); this->inputProgram = NULL; - if (this->cachedInstance) { - delete cachedInstance; - this->cachedInstance = NULL; - } + SingleThreadedNodeOperation::deinitExecution(); } -void *GlareBaseOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers) + +MemoryBuffer *GlareBaseOperation::createMemoryBuffer(rcti *rect2, MemoryBuffer **memoryBuffers) { - BLI_mutex_lock(getMutex()); - if (this->cachedInstance == NULL) { - MemoryBuffer *tile = (MemoryBuffer*)inputProgram->initializeTileData(rect, memoryBuffers); - float *data = new float[this->getWidth()*this->getHeight()*COM_NUMBER_OF_CHANNELS]; - this->generateGlare(data, tile, this->settings); - this->cachedInstance = data; - } - BLI_mutex_unlock(getMutex()); - return this->cachedInstance; + MemoryBuffer *tile = (MemoryBuffer*)inputProgram->initializeTileData(rect2, memoryBuffers); + rcti rect; + rect.xmin = 0; + rect.ymin = 0; + rect.xmax = getWidth(); + rect.ymax = getHeight(); + MemoryBuffer *result = new MemoryBuffer(NULL, &rect); + float *data = result->getBuffer(); + this->generateGlare(data, tile, this->settings); + return result; } bool GlareBaseOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output) { - if (this->cachedInstance != NULL) { + if (isCached()) { return false; } else { diff --git a/source/blender/compositor/operations/COM_GlareBaseOperation.h b/source/blender/compositor/operations/COM_GlareBaseOperation.h index 2fa8afc9c4c..d79f449f426 100644 --- a/source/blender/compositor/operations/COM_GlareBaseOperation.h +++ b/source/blender/compositor/operations/COM_GlareBaseOperation.h @@ -22,7 +22,8 @@ #ifndef _COM_GlareBaseOperation_h #define _COM_GlareBaseOperation_h -#include "COM_NodeOperation.h" + +#include "COM_SingleThreadedNodeOperation.h" #include "DNA_node_types.h" @@ -55,7 +56,7 @@ typedef float fRGB[4]; } (void)0 -class GlareBaseOperation : public NodeOperation { +class GlareBaseOperation : public SingleThreadedNodeOperation { private: /** * @brief Cached reference to the inputProgram @@ -66,16 +67,7 @@ private: * @brief settings of the glare node. */ NodeGlare * settings; - - float *cachedInstance; - public: - - /** - * the inner loop of this program - */ - void executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data); - /** * Initialize the execution */ @@ -86,15 +78,15 @@ public: */ void deinitExecution(); - void *initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers); - void setGlareSettings(NodeGlare * settings) {this->settings = settings;} bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output); + protected: GlareBaseOperation(); virtual void generateGlare(float *data, MemoryBuffer *inputTile, NodeGlare *settings) = 0; + MemoryBuffer *createMemoryBuffer(rcti *rect, MemoryBuffer **memoryBuffers); }; #endif diff --git a/source/blender/compositor/operations/COM_GlareGhostOperation.cpp b/source/blender/compositor/operations/COM_GlareGhostOperation.cpp index c5b1d6caa89..383a13c54de 100644 --- a/source/blender/compositor/operations/COM_GlareGhostOperation.cpp +++ b/source/blender/compositor/operations/COM_GlareGhostOperation.cpp @@ -45,15 +45,21 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No MemoryBuffer *gbuf = inputTile->duplicate(); MemoryBuffer *tbuf1 = inputTile->duplicate(); + bool breaked = false; + FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 0, 3); - FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 1, 3); - FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 2, 3); + if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 1, 3); + if (isBreaked()) breaked = true; + if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 2, 3); MemoryBuffer *tbuf2 = tbuf1->duplicate(); - FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 0, 3); - FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 1, 3); - FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 2, 3); + if (isBreaked()) breaked = true; + if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 0, 3); + if (isBreaked()) breaked = true; + if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 1, 3); + if (isBreaked()) breaked = true; + if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 2, 3); if (settings->iter & 1) ofs = 0.5f; else ofs = 0.f; for (x=0; x<(settings->iter*4); x++) { @@ -68,7 +74,7 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No sc = 2.13; isc = -0.97; - for (y=0; y<gbuf->getHeight(); y++) { + for (y=0; y<gbuf->getHeight() &(!breaked); y++) { v = (float)(y+0.5f) / (float)gbuf->getHeight(); for (x=0; x<gbuf->getWidth(); x++) { u = (float)(x+0.5f) / (float)gbuf->getWidth(); @@ -83,11 +89,13 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No gbuf->writePixel(x, y, c); } + if (isBreaked()) breaked = true; + } memset(tbuf1->getBuffer(), 0, tbuf1->getWidth()*tbuf1->getHeight()*COM_NUMBER_OF_CHANNELS*sizeof(float)); - for (n=1; n<settings->iter; n++) { - for (y=0; y<gbuf->getHeight(); y++) { + for (n=1; n<settings->iter &(!breaked); n++) { + for (y=0; y<gbuf->getHeight()&(!breaked); y++) { v = (float)(y+0.5f) / (float)gbuf->getHeight(); for (x=0; x<gbuf->getWidth(); x++) { u = (float)(x+0.5f) / (float)gbuf->getWidth(); @@ -103,6 +111,7 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No } tbuf1->writePixel(x, y, tc); } + if (isBreaked()) breaked = true; } memcpy(gbuf->getBuffer(), tbuf1->getBuffer(), tbuf1->getWidth()*tbuf1->getHeight()*COM_NUMBER_OF_CHANNELS*sizeof(float)); } diff --git a/source/blender/compositor/operations/COM_GlareSimpleStarOperation.cpp b/source/blender/compositor/operations/COM_GlareSimpleStarOperation.cpp index fba3eca4af9..4a393a33073 100644 --- a/source/blender/compositor/operations/COM_GlareSimpleStarOperation.cpp +++ b/source/blender/compositor/operations/COM_GlareSimpleStarOperation.cpp @@ -32,10 +32,11 @@ void GlareSimpleStarOperation::generateGlare(float *data, MemoryBuffer *inputTil MemoryBuffer *tbuf1 = inputTile->duplicate(); MemoryBuffer *tbuf2 = inputTile->duplicate(); - for (i=0; i<settings->iter; i++) { + bool breaked = false; + for (i=0; i<settings->iter && (!breaked); i++) { // // (x || x-1, y-1) to (x || x+1, y+1) // // F - for (y=0; y<this->getHeight(); y++) { + for (y=0; y<this->getHeight() && (!breaked); y++) { ym = y - i; yp = y + i; for (x=0; x<this->getWidth(); x++) { @@ -58,11 +59,13 @@ void GlareSimpleStarOperation::generateGlare(float *data, MemoryBuffer *inputTil madd_v3_v3fl(c, tc, f2); c[3] = 1.0f; tbuf2->writePixel(x, y, c); - + } + if (isBreaked()) { + breaked = true; } } // // B - for (y=tbuf1->getHeight()-1; y>=0; y--) { + for (y=tbuf1->getHeight()-1 && (!breaked); y>=0; y--) { ym = y - i; yp = y + i; for (x=tbuf1->getWidth()-1; x>=0; x--) { @@ -86,6 +89,9 @@ void GlareSimpleStarOperation::generateGlare(float *data, MemoryBuffer *inputTil c[3] = 1.0f; tbuf2->writePixel(x, y, c); } + if (isBreaked()) { + breaked = true; + } } } diff --git a/source/blender/compositor/operations/COM_GlareStreaksOperation.cpp b/source/blender/compositor/operations/COM_GlareStreaksOperation.cpp index 42b6a2b5e50..e735893ed6d 100644 --- a/source/blender/compositor/operations/COM_GlareStreaksOperation.cpp +++ b/source/blender/compositor/operations/COM_GlareStreaksOperation.cpp @@ -33,22 +33,23 @@ void GlareStreaksOperation::generateGlare(float *data, MemoryBuffer *inputTile, int size = inputTile->getWidth()*inputTile->getHeight(); int size4 = size*4; + bool breaked = false; MemoryBuffer *tsrc = inputTile->duplicate(); MemoryBuffer *tdst = new MemoryBuffer(NULL, inputTile->getRect()); tdst->clear(); memset(data, 0, size4*sizeof(float)); - for (a=0.f; a<DEG2RADF(360.0f); a+=ang) { + for (a=0.f; a<DEG2RADF(360.0f) && (!breaked); a+=ang) { const float an = a + settings->angle_ofs; const float vx = cos((double)an), vy = sin((double)an); - for (n=0; n<settings->iter; ++n) { + for (n=0; n<settings->iter && (!breaked); ++n) { const float p4 = pow(4.0, (double)n); const float vxp = vx*p4, vyp = vy*p4; const float wt = pow((double)settings->fade, (double)p4); const float cmo = 1.f - (float)pow((double)settings->colmod, (double)n+1); // colormodulation amount relative to current pass float *tdstcol = tdst->getBuffer(); - for (y=0; y<tsrc->getHeight(); ++y) { + for (y=0; y<tsrc->getHeight() && (!breaked); ++y) { for (x=0; x<tsrc->getWidth(); ++x, tdstcol+=4) { // first pass no offset, always same for every pass, exact copy, // otherwise results in uneven brightness, only need once @@ -71,11 +72,13 @@ void GlareStreaksOperation::generateGlare(float *data, MemoryBuffer *inputTile, tdstcol[2] = 0.5f*(tdstcol[2] + c1[2] + wt*(c2[2] + wt*(c3[2] + wt*c4[2]))); tdstcol[3] = 1.0f; } + if (isBreaked()) { + breaked = true; + } } memcpy(tsrc->getBuffer(), tdst->getBuffer(), sizeof(float)*size4); } -// addImage(sbuf, tsrc, 1.f/(float)(6 - ndg->iter)); // add result to data @todo float *sourcebuffer = tsrc->getBuffer(); float factor = 1.f/(float)(6 - settings->iter); for (int i = 0 ; i < size4; i ++) { diff --git a/source/blender/compositor/operations/COM_MaskOperation.cpp b/source/blender/compositor/operations/COM_MaskOperation.cpp index 35174349a63..bfbf2b42e82 100644 --- a/source/blender/compositor/operations/COM_MaskOperation.cpp +++ b/source/blender/compositor/operations/COM_MaskOperation.cpp @@ -67,7 +67,7 @@ void *MaskOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers if (!this->mask) return NULL; - BLI_mutex_lock(getMutex()); + lockMutex(); if (this->rasterizedMask == NULL) { int width = this->getWidth(); int height = this->getHeight(); @@ -78,8 +78,7 @@ void *MaskOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers this->rasterizedMask = buffer; } - BLI_mutex_unlock(getMutex()); - + unlockMutex(); return this->rasterizedMask; } diff --git a/source/blender/compositor/operations/COM_MovieDistortionOperation.cpp b/source/blender/compositor/operations/COM_MovieDistortionOperation.cpp index ebea9e8b4a2..d9e8977871f 100644 --- a/source/blender/compositor/operations/COM_MovieDistortionOperation.cpp +++ b/source/blender/compositor/operations/COM_MovieDistortionOperation.cpp @@ -52,7 +52,7 @@ void MovieDistortionOperation::initExecution() BKE_movieclip_user_set_frame(&clipUser, this->framenumber); BKE_movieclip_get_size(this->movieClip, &clipUser, &calibration_width, &calibration_height); - for (int i = 0 ; i < s_cache.size() ; i ++) { + for (unsigned int i = 0 ; i < s_cache.size() ; i ++) { DistortionCache *c = (DistortionCache*)s_cache[i]; if (c->isCacheFor(this->movieClip, this->width, this->height, calibration_width, calibration_height, this->distortion)) diff --git a/source/blender/compositor/operations/COM_MultilayerImageOperation.cpp b/source/blender/compositor/operations/COM_MultilayerImageOperation.cpp index f95dd12a81a..5e81cd639dd 100644 --- a/source/blender/compositor/operations/COM_MultilayerImageOperation.cpp +++ b/source/blender/compositor/operations/COM_MultilayerImageOperation.cpp @@ -47,7 +47,7 @@ void MultilayerColorOperation::executePixel(float *color, float x, float y, Pixe { int yi = y; int xi = x; - if (this->imageBuffer == NULL || xi < 0 || yi < 0 || xi >= this->getWidth() || yi >= this->getHeight() ) { + if (this->imageBuffer == NULL || xi < 0 || yi < 0 || (unsigned int)xi >= this->getWidth() || (unsigned int)yi >= this->getHeight() ) { color[0] = 0.0f; color[1] = 0.0f; color[2] = 0.0f; @@ -80,7 +80,7 @@ void MultilayerValueOperation::executePixel(float *color, float x, float y, Pixe { int yi = y; int xi = x; - if (this->imageBuffer == NULL || xi < 0 || yi < 0 || xi >= this->getWidth() || yi >= this->getHeight() ) { + if (this->imageBuffer == NULL || xi < 0 || yi < 0 || (unsigned int)xi >= this->getWidth() || (unsigned int)yi >= this->getHeight() ) { color[0] = 0.0f; } else { @@ -93,7 +93,7 @@ void MultilayerVectorOperation::executePixel(float *color, float x, float y, Pix { int yi = y; int xi = x; - if (this->imageBuffer == NULL || xi < 0 || yi < 0 || xi >= this->getWidth() || yi >= this->getHeight() ) { + if (this->imageBuffer == NULL || xi < 0 || yi < 0 || (unsigned int)xi >= this->getWidth() || (unsigned int)yi >= this->getHeight() ) { color[0] = 0.0f; } else { diff --git a/source/blender/compositor/operations/COM_NormalizeOperation.cpp b/source/blender/compositor/operations/COM_NormalizeOperation.cpp index d144739f845..4dd94943f4e 100644 --- a/source/blender/compositor/operations/COM_NormalizeOperation.cpp +++ b/source/blender/compositor/operations/COM_NormalizeOperation.cpp @@ -76,8 +76,7 @@ bool NormalizeOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe void *NormalizeOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers) { - BLI_mutex_lock(getMutex()); - + lockMutex(); if (this->cachedInstance == NULL) { MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers); /* using generic two floats struct to store x: min y: mult */ @@ -105,7 +104,7 @@ void *NormalizeOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu this->cachedInstance = minmult; } - BLI_mutex_unlock(getMutex()); + unlockMutex(); return this->cachedInstance; } diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl b/source/blender/compositor/operations/COM_OpenCLKernels.cl index aeccfcab8b5..e1f175b318a 100644 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl +++ b/source/blender/compositor/operations/COM_OpenCLKernels.cl @@ -6,8 +6,8 @@ const sampler_t SAMPLER_NEAREST = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS __constant const int2 zero = {0,0}; // KERNEL --- BOKEH BLUR --- -__kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __global __read_only image2d_t inputImage, - __global __read_only image2d_t bokehImage, __global __write_only image2d_t output, +__kernel void bokehBlurKernel(__read_only image2d_t boundingBox, __read_only image2d_t inputImage, + __read_only image2d_t bokehImage, __write_only image2d_t output, int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset) { int2 coords = {get_global_id(0), get_global_id(1)}; @@ -50,3 +50,65 @@ __kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __glob write_imagef(output, coords, color); } + +// KERNEL --- DILATE --- +__kernel void dilateKernel(__read_only image2d_t inputImage, __write_only image2d_t output, + int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension, + int2 offset) +{ + int2 coords = {get_global_id(0), get_global_id(1)}; + coords += offset; + const int2 realCoordinate = coords + offsetOutput; + + const int2 minXY = max(realCoordinate - scope, zero); + const int2 maxXY = min(realCoordinate + scope, dimension); + + float value = 0.0f; + int nx, ny; + int2 inputXy; + + for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) { + for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) { + const float deltaX = (realCoordinate.x - nx); + const float deltaY = (realCoordinate.y - ny); + const float measuredDistance = deltaX*deltaX+deltaY*deltaY; + if (measuredDistance <= distanceSquared) { + value = max(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0); + } + } + } + + float4 color = {value,0.0f,0.0f,0.0f}; + write_imagef(output, coords, color); +} + +// KERNEL --- DILATE --- +__kernel void erodeKernel(__read_only image2d_t inputImage, __write_only image2d_t output, + int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension, + int2 offset) +{ + int2 coords = {get_global_id(0), get_global_id(1)}; + coords += offset; + const int2 realCoordinate = coords + offsetOutput; + + const int2 minXY = max(realCoordinate - scope, zero); + const int2 maxXY = min(realCoordinate + scope, dimension); + + float value = 1.0f; + int nx, ny; + int2 inputXy; + + for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) { + for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) { + const float deltaX = (realCoordinate.x - nx); + const float deltaY = (realCoordinate.y - ny); + const float measuredDistance = deltaX*deltaX+deltaY*deltaY; + if (measuredDistance <= distanceSquared) { + value = min(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0); + } + } + } + + float4 color = {value,0.0f,0.0f,0.0f}; + write_imagef(output, coords, color); +} diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl.h b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h index 3cf33c75272..ef8668f6f21 100644 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl.h +++ b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h @@ -8,8 +8,8 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope "__constant const int2 zero = {0,0};\n" \ "\n" \ "// KERNEL --- BOKEH BLUR ---\n" \ -"__kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __global __read_only image2d_t inputImage,\n" \ -" __global __read_only image2d_t bokehImage, __global __write_only image2d_t output,\n" \ +"__kernel void bokehBlurKernel(__read_only image2d_t boundingBox, __read_only image2d_t inputImage,\n" \ +" __read_only image2d_t bokehImage, __write_only image2d_t output,\n" \ " int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset)\n" \ "{\n" \ " int2 coords = {get_global_id(0), get_global_id(1)};\n" \ @@ -26,8 +26,8 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope " if (tempBoundingBox > 0.0f) {\n" \ " const int2 bokehImageDim = get_image_dim(bokehImage);\n" \ " const int2 bokehImageCenter = bokehImageDim/2;\n" \ -" const int2 minXY = max(realCoordinate - radius, zero);;\n" \ -" const int2 maxXY = min(realCoordinate + radius, dimension);;\n" \ +" const int2 minXY = max(realCoordinate - radius, zero);\n" \ +" const int2 maxXY = min(realCoordinate + radius, dimension);\n" \ " int nx, ny;\n" \ "\n" \ " float2 uv;\n" \ @@ -52,4 +52,66 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope "\n" \ " write_imagef(output, coords, color);\n" \ "}\n" \ +"\n" \ +"// KERNEL --- DILATE ---\n" \ +"__kernel void dilateKernel(__read_only image2d_t inputImage, __write_only image2d_t output,\n" \ +" int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,\n" \ +" int2 offset)\n" \ +"{\n" \ +" int2 coords = {get_global_id(0), get_global_id(1)};\n" \ +" coords += offset;\n" \ +" const int2 realCoordinate = coords + offsetOutput;\n" \ +"\n" \ +" const int2 minXY = max(realCoordinate - scope, zero);\n" \ +" const int2 maxXY = min(realCoordinate + scope, dimension);\n" \ +"\n" \ +" float value = 0.0f;\n" \ +" int nx, ny;\n" \ +" int2 inputXy;\n" \ +"\n" \ +" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\n" \ +" for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {\n" \ +" const float deltaX = (realCoordinate.x - nx);\n" \ +" const float deltaY = (realCoordinate.y - ny);\n" \ +" const float measuredDistance = deltaX*deltaX+deltaY*deltaY;\n" \ +" if (measuredDistance <= distanceSquared) {\n" \ +" value = max(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \ +" }\n" \ +" }\n" \ +" }\n" \ +"\n" \ +" float4 color = {value,0.0f,0.0f,0.0f};\n" \ +" write_imagef(output, coords, color);\n" \ +"}\n" \ +"\n" \ +"// KERNEL --- DILATE ---\n" \ +"__kernel void erodeKernel(__read_only image2d_t inputImage, __write_only image2d_t output,\n" \ +" int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,\n" \ +" int2 offset)\n" \ +"{\n" \ +" int2 coords = {get_global_id(0), get_global_id(1)};\n" \ +" coords += offset;\n" \ +" const int2 realCoordinate = coords + offsetOutput;\n" \ +"\n" \ +" const int2 minXY = max(realCoordinate - scope, zero);\n" \ +" const int2 maxXY = min(realCoordinate + scope, dimension);\n" \ +"\n" \ +" float value = 1.0f;\n" \ +" int nx, ny;\n" \ +" int2 inputXy;\n" \ +"\n" \ +" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\n" \ +" for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {\n" \ +" const float deltaX = (realCoordinate.x - nx);\n" \ +" const float deltaY = (realCoordinate.y - ny);\n" \ +" const float measuredDistance = deltaX*deltaX+deltaY*deltaY;\n" \ +" if (measuredDistance <= distanceSquared) {\n" \ +" value = min(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \ +" }\n" \ +" }\n" \ +" }\n" \ +"\n" \ +" float4 color = {value,0.0f,0.0f,0.0f};\n" \ +" write_imagef(output, coords, color);\n" \ +"}\n" \ "\0"; diff --git a/source/blender/compositor/operations/COM_OutputFileOperation.cpp b/source/blender/compositor/operations/COM_OutputFileOperation.cpp index 8d39e987bd4..1438116f313 100644 --- a/source/blender/compositor/operations/COM_OutputFileOperation.cpp +++ b/source/blender/compositor/operations/COM_OutputFileOperation.cpp @@ -177,7 +177,7 @@ void OutputOpenExrMultiLayerOperation::add_layer(const char *name, DataType data void OutputOpenExrMultiLayerOperation::initExecution() { - for (int i=0; i < layers.size(); ++i) { + for (unsigned int i=0; i < layers.size(); ++i) { layers[i].imageInput = getInputSocketReader(i); layers[i].outputBuffer = init_buffer(this->getWidth(), this->getHeight(), layers[i].datatype); } @@ -185,7 +185,7 @@ void OutputOpenExrMultiLayerOperation::initExecution() void OutputOpenExrMultiLayerOperation::executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer** memoryBuffers) { - for (int i=0; i < layers.size(); ++i) { + for (unsigned int i=0; i < layers.size(); ++i) { write_buffer_rect(rect, memoryBuffers, this->tree, layers[i].imageInput, layers[i].outputBuffer, this->getWidth(), layers[i].datatype); } } @@ -203,7 +203,7 @@ void OutputOpenExrMultiLayerOperation::deinitExecution() (this->scene->r.scemode & R_EXTENSION), true); BLI_make_existing_file(filename); - for (int i=0; i < layers.size(); ++i) { + for (unsigned int i=0; i < layers.size(); ++i) { char channelname[EXR_TOT_MAXNAME]; BLI_strncpy(channelname, layers[i].name, sizeof(channelname)-2); char *channelname_ext = channelname + strlen(channelname); @@ -251,7 +251,7 @@ void OutputOpenExrMultiLayerOperation::deinitExecution() } IMB_exr_close(exrhandle); - for (int i=0; i < layers.size(); ++i) { + for (unsigned int i=0; i < layers.size(); ++i) { if (layers[i].outputBuffer) { MEM_freeN(layers[i].outputBuffer); layers[i].outputBuffer = NULL; diff --git a/source/blender/compositor/operations/COM_PreviewOperation.h b/source/blender/compositor/operations/COM_PreviewOperation.h index 2b81b914746..32c0b6ecc77 100644 --- a/source/blender/compositor/operations/COM_PreviewOperation.h +++ b/source/blender/compositor/operations/COM_PreviewOperation.h @@ -34,7 +34,6 @@ protected: * @brief holds reference to the SDNA bNode, where this nodes will render the preview image for */ bNode *node; - const bNodeTree *tree; SocketReader *input; float divider; @@ -48,7 +47,6 @@ public: void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer **memoryBuffers); void determineResolution(unsigned int resolution[], unsigned int preferredResolution[]); void setbNode(bNode *node) { this->node = node;} - void setbNodeTree(const bNodeTree *tree) { this->tree = tree;} bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output); bool isPreviewOperation() {return true;} diff --git a/source/blender/compositor/operations/COM_TonemapOperation.cpp b/source/blender/compositor/operations/COM_TonemapOperation.cpp index 75adcf524c4..9a02d6a58ca 100644 --- a/source/blender/compositor/operations/COM_TonemapOperation.cpp +++ b/source/blender/compositor/operations/COM_TonemapOperation.cpp @@ -118,7 +118,7 @@ bool TonemapOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferO void *TonemapOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers) { - BLI_mutex_lock(getMutex()); + lockMutex(); if (this->cachedInstance == NULL) { MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers); AvgLogLum *data = new AvgLogLum(); @@ -150,7 +150,7 @@ void *TonemapOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuff data->igm = (this->data->gamma==0.f) ? 1 : (1.f / this->data->gamma); this->cachedInstance = data; } - BLI_mutex_unlock(getMutex()); + unlockMutex(); return this->cachedInstance; } diff --git a/source/blender/compositor/operations/COM_VectorBlurOperation.cpp b/source/blender/compositor/operations/COM_VectorBlurOperation.cpp index e6305dc26a2..3dac3547b8a 100644 --- a/source/blender/compositor/operations/COM_VectorBlurOperation.cpp +++ b/source/blender/compositor/operations/COM_VectorBlurOperation.cpp @@ -77,7 +77,7 @@ void *VectorBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB { if (this->cachedInstance) return this->cachedInstance; - BLI_mutex_lock(getMutex()); + lockMutex(); if (this->cachedInstance == NULL) { MemoryBuffer *tile = (MemoryBuffer*)inputImageProgram->initializeTileData(rect, memoryBuffers); MemoryBuffer *speed = (MemoryBuffer*)inputSpeedProgram->initializeTileData(rect, memoryBuffers); @@ -87,7 +87,7 @@ void *VectorBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB this->generateVectorBlur(data, tile, speed, z); this->cachedInstance = data; } - BLI_mutex_unlock(getMutex()); + unlockMutex(); return this->cachedInstance; } diff --git a/source/blender/compositor/operations/COM_ViewerBaseOperation.h b/source/blender/compositor/operations/COM_ViewerBaseOperation.h index 51fa8cecc0d..8864894f0ed 100644 --- a/source/blender/compositor/operations/COM_ViewerBaseOperation.h +++ b/source/blender/compositor/operations/COM_ViewerBaseOperation.h @@ -34,7 +34,6 @@ protected: ImageUser * imageUser; void *lock; bool active; - const bNodeTree *tree; float centerX; float centerY; OrderOfChunks chunkOrder; @@ -49,7 +48,6 @@ public: void setImageUser(ImageUser *imageUser) {this->imageUser = imageUser;} const bool isActiveViewerOutput() const {return active;} void setActive(bool active) {this->active = active;} - void setbNodeTree(const bNodeTree *tree) {this->tree = tree;} void setCenterX(float centerX) {this->centerX = centerX;} void setCenterY(float centerY) {this->centerY = centerY;} void setChunkOrder(OrderOfChunks tileOrder) {this->chunkOrder = tileOrder;} diff --git a/source/blender/compositor/operations/COM_ViewerOperation.cpp b/source/blender/compositor/operations/COM_ViewerOperation.cpp index 22e6511fbe7..3bec1ef48f5 100644 --- a/source/blender/compositor/operations/COM_ViewerOperation.cpp +++ b/source/blender/compositor/operations/COM_ViewerOperation.cpp @@ -104,7 +104,7 @@ void ViewerOperation::executeRegion(rcti *rect, unsigned int tileNumber, MemoryB offset +=4; } - if (tree->test_break && tree->test_break(tree->tbh)) { + if (isBreaked()) { breaked = true; } diff --git a/source/blender/compositor/operations/COM_WriteBufferOperation.cpp b/source/blender/compositor/operations/COM_WriteBufferOperation.cpp index 8888d30ba2f..087ab50cf47 100644 --- a/source/blender/compositor/operations/COM_WriteBufferOperation.cpp +++ b/source/blender/compositor/operations/COM_WriteBufferOperation.cpp @@ -30,7 +30,6 @@ WriteBufferOperation::WriteBufferOperation() :NodeOperation() this->memoryProxy = new MemoryProxy(); this->memoryProxy->setWriteBufferOperation(this); this->memoryProxy->setExecutor(NULL); - this->tree = NULL; } WriteBufferOperation::~WriteBufferOperation() { @@ -78,7 +77,7 @@ void WriteBufferOperation::executeRegion(rcti *rect, unsigned int tileNumber, Me offset4 +=COM_NUMBER_OF_CHANNELS; } - if (tree->test_break && tree->test_break(tree->tbh)) { + if (isBreaked()) { breaked = true; } @@ -103,7 +102,7 @@ void WriteBufferOperation::executeRegion(rcti *rect, unsigned int tileNumber, Me input->read(&(buffer[offset4]), x, y, COM_PS_NEAREST, memoryBuffers); offset4 +=COM_NUMBER_OF_CHANNELS; } - if (tree->test_break && tree->test_break(tree->tbh)) { + if (isBreaked()) { breaked = true; } } diff --git a/source/blender/compositor/operations/COM_WriteBufferOperation.h b/source/blender/compositor/operations/COM_WriteBufferOperation.h index 068adc03293..fff3212b410 100644 --- a/source/blender/compositor/operations/COM_WriteBufferOperation.h +++ b/source/blender/compositor/operations/COM_WriteBufferOperation.h @@ -33,7 +33,6 @@ class WriteBufferOperation: public NodeOperation { MemoryProxy *memoryProxy; NodeOperation *input; - const bNodeTree * tree; public: WriteBufferOperation(); ~WriteBufferOperation(); @@ -45,7 +44,6 @@ public: void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer** memoryBuffers); void initExecution(); void deinitExecution(); - void setbNodeTree(const bNodeTree *tree) {this->tree = tree;} void executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer** memoryBuffers, MemoryBuffer* outputBuffer); void readResolutionFromInputSocket(); |