archive/blender-archive
Archived
1
1
Fork 0
Code Pull Requests Packages Activity

* optimized threading

Browse Source 
* break out with glare node
 * Added OpenCL kernels compatible with AMD still need some testing.
This commit is contained in:
Jeroen Bakker
2012-06-13 12:34:56 +00:00
parent 4ba456d175
commit be1b5f82ce
40 changed files with 483 additions and 163 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
source/blender/compositor/CMakeLists.txt
View File

@@ -97,6 +97,9 @@ set(SRC
intern/COM_CompositorContext.h
intern/COM_ChannelInfo.cpp
intern/COM_ChannelInfo.h
intern/COM_SingleThreadedNodeOperation.cpp
intern/COM_SingleThreadedNodeOperation.h
operations/COM_QualityStepHelper.h
operations/COM_QualityStepHelper.cpp

31
source/blender/compositor/intern/COM_ExecutionGroup.cpp
View File

@@ -54,6 +54,7 @@ ExecutionGroup::ExecutionGroup()
this->numberOfChunks = 0;
this->initialized = false;
this->openCL = false;
this->singleThreaded = false;
this->chunksFinished = 0;
}
@@ -100,6 +101,7 @@ void ExecutionGroup::addOperation(ExecutionSystem *system, NodeOperation *operat
if (!operation->isBufferOperation()) {
this->complex = operation->isComplex();
this->openCL = operation->isOpenCL();
this->singleThreaded = operation->isSingleThreaded();
this->initialized = true;
}
this->operations.push_back(operation);
@@ -191,10 +193,17 @@ void ExecutionGroup::determineResolution(unsigned int resolution[])
void ExecutionGroup::determineNumberOfChunks()
{
const float chunkSizef = this->chunkSize;
this->numberOfXChunks = ceil(this->width / chunkSizef);
this->numberOfYChunks = ceil(this->height / chunkSizef);
this->numberOfChunks = this->numberOfXChunks * this->numberOfYChunks;
if (singleThreaded) {
this->numberOfXChunks = 1;
this->numberOfYChunks = 1;
this->numberOfChunks = 1;
}
else {
const float chunkSizef = this->chunkSize;
this->numberOfXChunks = ceil(this->width / chunkSizef);
this->numberOfYChunks = ceil(this->height / chunkSizef);
this->numberOfChunks = this->numberOfXChunks * this->numberOfYChunks;
}
}
/**
@@ -435,9 +444,14 @@ void ExecutionGroup::finalizeChunkExecution(int chunkNumber, MemoryBuffer** memo
inline void ExecutionGroup::determineChunkRect(rcti *rect, const unsigned int xChunk, const unsigned int yChunk ) const
{
const unsigned int minx = xChunk * chunkSize;
const unsigned int miny = yChunk * chunkSize;
BLI_init_rcti(rect, minx, min(minx + this->chunkSize, this->width), miny, min(miny + this->chunkSize, this->height));
if (singleThreaded) {
BLI_init_rcti(rect, 0, this->width, 0, this->height);
}
else {
const unsigned int minx = xChunk * chunkSize;
const unsigned int miny = yChunk * chunkSize;
BLI_init_rcti(rect, minx, min(minx + this->chunkSize, this->width), miny, min(miny + this->chunkSize, this->height));
}
}
void ExecutionGroup::determineChunkRect(rcti *rect, const unsigned int chunkNumber) const
@@ -462,6 +476,9 @@ MemoryBuffer *ExecutionGroup::allocateOutputBuffer(int chunkNumber, rcti *rect)
bool ExecutionGroup::scheduleAreaWhenPossible(ExecutionSystem * graph, rcti *area)
{
if (singleThreaded) {
return scheduleChunkWhenPossible(graph, 0, 0);
}
// find all chunks inside the rect
// determine minxchunk, minychunk, maxxchunk, maxychunk where x and y are chunknumbers

22
source/blender/compositor/intern/COM_ExecutionGroup.h
View File

@@ -63,10 +63,6 @@ class Device;
class ExecutionGroup {
private:
// fields
/**
* @brief unique identifier of this node.
*/
string id;
/**
* @brief list of operations in this ExecutionGroup
@@ -120,6 +116,11 @@ private:
*/
bool openCL;
/**
* @brief Is this Execution group SingleThreaded
*/
bool singleThreaded;
/**
* @brief what is the maximum number field of all ReadBufferOperation in this ExecutionGroup.
* @note this is used to construct the MemoryBuffers that will be passed during execution.
@@ -233,18 +234,7 @@ private:
public:
// constructors
ExecutionGroup();
/**
* @brief set the id of this ExecutionGroup
* @param id
*/
void setId(string id) {this->id = id;}
/**
* @brief return the id of this ExecutionGroup
*/
const string getId() const {return this->id;}
// methods
/**
* @brief check to see if a NodeOperation is already inside this execution group

4
source/blender/compositor/intern/COM_ExecutionSystem.cpp
View File

@@ -124,6 +124,7 @@ void ExecutionSystem::execute()
for (index = 0 ; index < this->operations.size() ; index ++) {
NodeOperation * operation = this->operations[index];
operation->setbNodeTree(this->context.getbNodeTree());
operation->initExecution();
}
for (index = 0 ; index < this->groups.size() ; index ++) {
@@ -153,7 +154,7 @@ void ExecutionSystem::execute()
void ExecutionSystem::executeGroups(CompositorPriority priority)
{
int index;
unsigned int index;
vector<ExecutionGroup*> executionGroups;
this->findOutputExecutionGroup(&executionGroups, priority);
@@ -166,6 +167,7 @@ void ExecutionSystem::executeGroups(CompositorPriority priority)
void ExecutionSystem::addOperation(NodeOperation *operation)
{
ExecutionSystemHelper::addOperation(this->operations, operation);
// operation->setBTree
}
void ExecutionSystem::addReadWriteBufferOperations(NodeOperation *operation)

2
source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp
View File

@@ -65,7 +65,7 @@ Node *ExecutionSystemHelper::addbNodeTree(ExecutionSystem &system, int nodes_sta
}
/* Expand group nodes */
for (int i=nodes_start; i < nodes.size(); ++i) {
for (unsigned int i=nodes_start; i < nodes.size(); ++i) {
Node *execnode = nodes[i];
if (execnode->isGroupNode()) {
GroupNode * groupNode = (GroupNode*)execnode;

21
source/blender/compositor/intern/COM_NodeOperation.cpp
View File

@@ -34,6 +34,7 @@ NodeOperation::NodeOperation()
this->width = 0;
this->height = 0;
this->openCL = false;
this->btree = NULL;
}
void NodeOperation::determineResolution(unsigned int resolution[], unsigned int preferredResolution[])
@@ -74,10 +75,22 @@ void NodeOperation::initMutex()
{
BLI_mutex_init(&mutex);
}
void NodeOperation::lockMutex()
{
BLI_mutex_lock(&mutex);
}
void NodeOperation::unlockMutex()
{
BLI_mutex_unlock(&mutex);
}
void NodeOperation::deinitMutex()
{
BLI_mutex_end(&mutex);
}
void NodeOperation::deinitExecution()
{
}
@@ -196,14 +209,15 @@ void NodeOperation::COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel,
size_t size[2];
cl_int2 offset;
for (offsety = 0 ; offsety < height; offsety+=localSize) {
bool breaked = false;
for (offsety = 0 ; offsety < height && (!breaked); offsety+=localSize) {
offset[1] = offsety;
if (offsety+localSize < height) {
size[1] = localSize;
} else {
size[1] = height - offsety;
}
for (offsetx = 0 ; offsetx < width ; offsetx+=localSize) {
for (offsetx = 0 ; offsetx < width && (!breaked) ; offsetx+=localSize) {
if (offsetx+localSize < width) {
size[0] = localSize;
} else {
@@ -216,6 +230,9 @@ void NodeOperation::COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel,
error = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, size, 0, 0, 0, NULL);
if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); }
clFlush(queue);
if (isBreaked()) {
breaked = false;
}
}
}
}

22
source/blender/compositor/intern/COM_NodeOperation.h
View File

@@ -77,6 +77,11 @@ private:
* @see NodeOperation.getMutex retrieve a pointer to this mutex.
*/
ThreadMutex mutex;
/**
* @brief reference to the editing bNodeTree only used for break callback
*/
const bNodeTree *btree;
public:
/**
@@ -119,9 +124,10 @@ public:
* for all other operations this will result in false.
*/
virtual int isBufferOperation() {return false;}
virtual int isSingleThreaded() {return false;}
void setbNodeTree(const bNodeTree * tree) {this->btree = tree;}
virtual void initExecution();
void initMutex();
/**
* @brief when a chunk is executed by a CPUDevice, this method is called
@@ -161,7 +167,6 @@ public:
*/
virtual void executeOpenCL(cl_context context,cl_program program, cl_command_queue queue, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer** inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, list<cl_kernel> *clKernelsToCleanUp) {}
virtual void deinitExecution();
void deinitMutex();
bool isResolutionSet() {
return this->width != 0 && height != 0;
@@ -236,6 +241,11 @@ public:
virtual bool isViewerOperation() {return false;}
virtual bool isPreviewOperation() {return false;}
inline bool isBreaked() {
return btree->test_break(btree->tbh);
}
protected:
NodeOperation();
@@ -244,7 +254,11 @@ protected:
SocketReader *getInputSocketReader(unsigned int inputSocketindex);
NodeOperation *getInputOperation(unsigned int inputSocketindex);
inline ThreadMutex *getMutex() {return &this->mutex;}
void deinitMutex();
void initMutex();
void lockMutex();
void unlockMutex();
/**
* @brief set whether this operation is complex
@@ -264,7 +278,7 @@ protected:
static void COM_clAttachOutputMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, cl_mem clOutputMemoryBuffer);
void COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex);
static void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer* outputMemoryBuffer);
static void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex);
void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex);
cl_kernel COM_clCreateKernel(cl_program program, const char* kernelname, list<cl_kernel> *clKernelsToCleanUp);
};

60
source/blender/compositor/intern/COM_SingleThreadedNodeOperation.cpp Normal file
View File

@@ -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;
}

60
source/blender/compositor/intern/COM_SingleThreadedNodeOperation.h Normal file
View File

@@ -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

1
source/blender/compositor/intern/COM_compositor.cpp
View File

@@ -51,7 +51,6 @@ void COM_execute(bNodeTree *editingtree, int rendering)
/* set progress bar to 0% and status to init compositing*/
editingtree->progress(editingtree->prh, 0.0);
editingtree->stats_draw(editingtree->sdh, (char*)"Compositing");
/* initialize execution system */
ExecutionSystem *system = new ExecutionSystem(editingtree, rendering);

2
source/blender/compositor/nodes/COM_DilateErodeNode.cpp
View File

@@ -36,7 +36,7 @@ void DilateErodeNode::convertToOperations(ExecutionSystem *graph, CompositorCont
bNode *editorNode = this->getbNode();
if (editorNode->custom1 == CMP_NODE_DILATEERODE_DISTANCE_THRESH) {
DilateErodeDistanceOperation *operation = new DilateErodeDistanceOperation();
DilateErodeThresholdOperation *operation = new DilateErodeThresholdOperation();
operation->setDistance(editorNode->custom2);
operation->setInset(editorNode->custom3);

3
source/blender/compositor/nodes/COM_MuteNode.cpp
View File

@@ -44,7 +44,7 @@ void MuteNode::reconnect(ExecutionSystem * graph, OutputSocket * output)
}
}
NodeOperation *operation;
NodeOperation *operation = NULL;
switch (output->getDataType()) {
case COM_DT_VALUE:
{
@@ -74,7 +74,6 @@ void MuteNode::reconnect(ExecutionSystem * graph, OutputSocket * output)
}
/* quiet warnings */
case COM_DT_UNKNOWN:
operation = NULL;
break;
}

6
source/blender/compositor/operations/COM_AntiAliasOperation.cpp
View File

@@ -44,7 +44,7 @@ void AntiAliasOperation::initExecution()
void AntiAliasOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void * data)
{
if (y < 0 || y >= this->height || x < 0 || x >= this->width) {
if (y < 0 || (unsigned int)y >= this->height || x < 0 || (unsigned int)x >= this->width) {
color[0] = 0.0f;
}
else {
@@ -85,7 +85,7 @@ bool AntiAliasOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe
void *AntiAliasOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
{
if (this->buffer) {return buffer;}
BLI_mutex_lock(getMutex());
lockMutex();
if (this->buffer == NULL) {
MemoryBuffer *tile = (MemoryBuffer*)valueReader->initializeTileData(rect, memoryBuffers);
int size = tile->getHeight()*tile->getWidth();
@@ -100,6 +100,6 @@ void *AntiAliasOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu
antialias_tagbuf(tile->getWidth(), tile->getHeight(), valuebuffer);
this->buffer = valuebuffer;
}
BLI_mutex_unlock(getMutex());
unlockMutex();
return this->buffer;
}

4
source/blender/compositor/operations/COM_CalculateMeanOperation.cpp
View File

@@ -72,13 +72,13 @@ bool CalculateMeanOperation::determineDependingAreaOfInterest(rcti *input, ReadB
void *CalculateMeanOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
{
BLI_mutex_lock(getMutex());
lockMutex();
if (!this->iscalculated) {
MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers);
calculateMean(tile);
this->iscalculated = true;
}
BLI_mutex_unlock(getMutex());
unlockMutex();
return NULL;
}

4
source/blender/compositor/operations/COM_CalculateStandardDeviationOperation.cpp
View File

@@ -37,7 +37,7 @@ void CalculateStandardDeviationOperation::executePixel(float *color, int x, int
void *CalculateStandardDeviationOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
{
BLI_mutex_lock(getMutex());
lockMutex();
if (!this->iscalculated) {
MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers);
CalculateMeanOperation::calculateMean(tile);
@@ -92,6 +92,6 @@ void *CalculateStandardDeviationOperation::initializeTileData(rcti *rect, Memory
this->standardDeviation = sqrt(sum / (float)(pixels-1));
this->iscalculated = true;
}
BLI_mutex_unlock(getMutex());
unlockMutex();
return NULL;
}

4
source/blender/compositor/operations/COM_CompositorOperation.cpp
View File

@@ -59,7 +59,7 @@ void CompositorOperation::initExecution()
void CompositorOperation::deinitExecution()
{
if (tree->test_break && !tree->test_break(tree->tbh)) {
if (isBreaked()) {
const Scene * scene = this->scene;
Render *re = RE_GetRender(scene->id.name);
RenderResult *rr = RE_AcquireResultWrite(re);
@@ -118,7 +118,7 @@ void CompositorOperation::executeRegion(rcti *rect, unsigned int tileNumber, Mem
buffer[offset+2] = color[2];
buffer[offset+3] = color[3];
offset +=COM_NUMBER_OF_CHANNELS;
if (tree->test_break && tree->test_break(tree->tbh)) {
if (isBreaked()) {
breaked = true;
}
}

6
source/blender/compositor/operations/COM_CompositorOperation.h
View File

@@ -36,11 +36,6 @@ private:
*/
const Scene *scene;
/**
* @brief local reference to the node tree
*/
const bNodeTree *tree;
/**
* @brief reference to the output float buffer
*/
@@ -59,7 +54,6 @@ public:
CompositorOperation();
void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer** memoryBuffers);
void setScene(const Scene *scene) {this->scene = scene;}
void setbNodeTree(const bNodeTree *tree) {this->tree = tree;}
bool isOutputOperation(bool rendering) const {return true;}
void initExecution();
void deinitExecution();

64
source/blender/compositor/operations/COM_DilateErodeOperation.cpp
View File

@@ -24,7 +24,7 @@
#include "BLI_math.h"
// DilateErode Distance Threshold
DilateErodeDistanceOperation::DilateErodeDistanceOperation(): NodeOperation()
DilateErodeThresholdOperation::DilateErodeThresholdOperation(): NodeOperation()
{
this->addInputSocket(COM_DT_VALUE);
this->addOutputSocket(COM_DT_VALUE);
@@ -34,7 +34,7 @@ DilateErodeDistanceOperation::DilateErodeDistanceOperation(): NodeOperation()
this->_switch = 0.5f;
this->distance = 0.0f;
}
void DilateErodeDistanceOperation::initExecution()
void DilateErodeThresholdOperation::initExecution()
{
this->inputProgram = this->getInputSocketReader(0);
if (this->distance < 0.0f) {
@@ -53,13 +53,13 @@ void DilateErodeDistanceOperation::initExecution()
}
}
void *DilateErodeDistanceOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
void *DilateErodeThresholdOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
{
void *buffer = inputProgram->initializeTileData(NULL, memoryBuffers);
return buffer;
}
void DilateErodeDistanceOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data)
void DilateErodeThresholdOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data)
{
float inputValue[4];
const float sw = this->_switch;
@@ -142,12 +142,12 @@ void DilateErodeDistanceOperation::executePixel(float *color, int x, int y, Memo
}
}
void DilateErodeDistanceOperation::deinitExecution()
void DilateErodeThresholdOperation::deinitExecution()
{
this->inputProgram = NULL;
}
bool DilateErodeDistanceOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output)
bool DilateErodeThresholdOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output)
{
rcti newInput;
@@ -167,6 +167,7 @@ DilateDistanceOperation::DilateDistanceOperation(): NodeOperation()
this->setComplex(true);
this->inputProgram = NULL;
this->distance = 0.0f;
this->setOpenCL(true);
}
void DilateDistanceOperation::initExecution()
{
@@ -231,6 +232,28 @@ bool DilateDistanceOperation::determineDependingAreaOfInterest(rcti *input, Read
return NodeOperation::determineDependingAreaOfInterest(&newInput, readOperation, output);
}
static cl_kernel dilateKernel = 0;
void DilateDistanceOperation::executeOpenCL(cl_context context, cl_program program, cl_command_queue queue,
MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer,
MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp,
list<cl_kernel> *clKernelsToCleanUp)
{
if (!dilateKernel) {
dilateKernel = COM_clCreateKernel(program, "dilateKernel", NULL);
}
cl_int distanceSquared = this->distance*this->distance;
cl_int scope = this->scope;
COM_clAttachMemoryBufferToKernelParameter(context, dilateKernel, 0, 2, clMemToCleanUp, inputMemoryBuffers, this->inputProgram);
COM_clAttachOutputMemoryBufferToKernelParameter(dilateKernel, 1, clOutputBuffer);
COM_clAttachMemoryBufferOffsetToKernelParameter(dilateKernel, 3, outputMemoryBuffer);
clSetKernelArg(dilateKernel, 4, sizeof(cl_int), &scope);
clSetKernelArg(dilateKernel, 5, sizeof(cl_int), &distanceSquared);
COM_clAttachSizeToKernelParameter(dilateKernel, 6);
COM_clEnqueueRange(queue, dilateKernel, outputMemoryBuffer, 7);
}
// Erode Distance
ErodeDistanceOperation::ErodeDistanceOperation() : DilateDistanceOperation()
{
@@ -268,6 +291,27 @@ void ErodeDistanceOperation::executePixel(float *color, int x, int y, MemoryBuff
color[0] = value;
}
static cl_kernel erodeKernel = 0;
void ErodeDistanceOperation::executeOpenCL(cl_context context, cl_program program, cl_command_queue queue,
MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer,
MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp,
list<cl_kernel> *clKernelsToCleanUp)
{
if (!erodeKernel) {
erodeKernel = COM_clCreateKernel(program, "erodeKernel", NULL);
}
cl_int distanceSquared = this->distance*this->distance;
cl_int scope = this->scope;
COM_clAttachMemoryBufferToKernelParameter(context, erodeKernel, 0, 2, clMemToCleanUp, inputMemoryBuffers, this->inputProgram);
COM_clAttachOutputMemoryBufferToKernelParameter(erodeKernel, 1, clOutputBuffer);
COM_clAttachMemoryBufferOffsetToKernelParameter(erodeKernel, 3, outputMemoryBuffer);
clSetKernelArg(erodeKernel, 4, sizeof(cl_int), &scope);
clSetKernelArg(erodeKernel, 5, sizeof(cl_int), &distanceSquared);
COM_clAttachSizeToKernelParameter(erodeKernel, 6);
COM_clEnqueueRange(queue, erodeKernel, outputMemoryBuffer, 7);
}
// Dilate step
DilateStepOperation::DilateStepOperation(): NodeOperation()
{
@@ -288,7 +332,7 @@ void *DilateStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB
if (this->cached_buffer != NULL) {
return this->cached_buffer;
}
BLI_mutex_lock(getMutex());
lockMutex();
if (this->cached_buffer == NULL) {
MemoryBuffer *buffer = (MemoryBuffer*)inputProgram->initializeTileData(NULL, memoryBuffers);
float *rectf = buffer->convertToValueBuffer();
@@ -327,7 +371,7 @@ void *DilateStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB
}
this->cached_buffer = rectf;
}
BLI_mutex_unlock(getMutex());
unlockMutex();
return this->cached_buffer;
}
@@ -374,7 +418,7 @@ void *ErodeStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu
if (this->cached_buffer != NULL) {
return this->cached_buffer;
}
BLI_mutex_lock(getMutex());
lockMutex();
if (this->cached_buffer == NULL) {
MemoryBuffer *buffer = (MemoryBuffer*)inputProgram->initializeTileData(NULL, memoryBuffers);
float *rectf = buffer->convertToValueBuffer();
@@ -413,6 +457,6 @@ void *ErodeStepOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu
}
this->cached_buffer = rectf;
}
BLI_mutex_unlock(getMutex());
unlockMutex();
return this->cached_buffer;
}

16
source/blender/compositor/operations/COM_DilateErodeOperation.h
View File

@@ -25,7 +25,7 @@
#include "COM_NodeOperation.h"
class DilateErodeDistanceOperation : public NodeOperation {
class DilateErodeThresholdOperation : public NodeOperation {
private:
/**
* Cached reference to the inputProgram
@@ -42,7 +42,7 @@ private:
*/
int scope;
public:
DilateErodeDistanceOperation();
DilateErodeThresholdOperation();
/**
* the inner loop of this program
@@ -70,11 +70,11 @@ public:
class DilateDistanceOperation : public NodeOperation {
private:
protected:
/**
* Cached reference to the inputProgram
*/
SocketReader * inputProgram;
protected:
float distance;
int scope;
public:
@@ -98,6 +98,11 @@ public:
void setDistance(float distance) {this->distance = distance;}
bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output);
void executeOpenCL(cl_context context, cl_program program, cl_command_queue queue,
MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer,
MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp,
list<cl_kernel> *clKernelsToCleanUp);
};
class ErodeDistanceOperation : public DilateDistanceOperation {
public:
@@ -107,6 +112,11 @@ public:
* the inner loop of this program
*/
void executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data);
void executeOpenCL(cl_context context, cl_program program, cl_command_queue queue,
MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer,
MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp,
list<cl_kernel> *clKernelsToCleanUp);
};
class DilateStepOperation : public NodeOperation {

9
source/blender/compositor/operations/COM_DoubleEdgeMaskOperation.cpp
View File

@@ -23,6 +23,7 @@
#include "COM_DoubleEdgeMaskOperation.h"
#include "BLI_math.h"
#include "DNA_node_types.h"
#include "MEM_guardedalloc.h"
// this part has been copied from the double edge mask
// Contributor(s): Peter Larabell.
@@ -1215,12 +1216,12 @@ void DoubleEdgeMaskOperation::doDoubleEdgeMask(float *imask, float *omask, float
gsz=rsize[2]; // by the do_*EdgeDetection() function.
fsz=gsz+isz+osz; // calculate size of pixel index buffer needed
gbuf = new unsigned short[fsz*2]; // allocate edge/gradient pixel index buffer
gbuf = (unsigned short*)MEM_callocN(sizeof (unsigned short)*fsz*2, "DEM"); // allocate edge/gradient pixel index buffer
do_createEdgeLocationBuffer(t,rw,lres,res,gbuf,&innerEdgeOffset,&outerEdgeOffset,isz,gsz);
do_fillGradientBuffer(rw,res,gbuf,isz,osz,gsz,innerEdgeOffset,outerEdgeOffset);
delete [] gbuf; // free the gradient index buffer
MEM_freeN(gbuf); // free the gradient index buffer
}
}
@@ -1263,7 +1264,7 @@ void *DoubleEdgeMaskOperation::initializeTileData(rcti *rect, MemoryBuffer **mem
{
if (this->cachedInstance) return this->cachedInstance;
BLI_mutex_lock(getMutex());
lockMutex();
if (this->cachedInstance == NULL) {
MemoryBuffer *innerMask = (MemoryBuffer*)inputInnerMask->initializeTileData(rect, memoryBuffers);
MemoryBuffer *outerMask = (MemoryBuffer*)inputOuterMask->initializeTileData(rect, memoryBuffers);
@@ -1275,7 +1276,7 @@ void *DoubleEdgeMaskOperation::initializeTileData(rcti *rect, MemoryBuffer **mem
delete omask;
this->cachedInstance = data;
}
BLI_mutex_unlock(getMutex());
unlockMutex();
return this->cachedInstance;
}
void DoubleEdgeMaskOperation::executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data)

4
source/blender/compositor/operations/COM_FastGaussianBlurOperation.cpp
View File

@@ -79,7 +79,7 @@ void FastGaussianBlurOperation::deinitExecution()
void *FastGaussianBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
{
BLI_mutex_lock(this->getMutex());
lockMutex();
if (!iirgaus) {
MemoryBuffer *newBuf = (MemoryBuffer*)this->inputProgram->initializeTileData(rect, memoryBuffers);
MemoryBuffer *copy = newBuf->duplicate();
@@ -105,7 +105,7 @@ void *FastGaussianBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **m
}
this->iirgaus = copy;
}
BLI_mutex_unlock(this->getMutex());
unlockMutex();
return iirgaus;
}

47
source/blender/compositor/operations/COM_GlareBaseOperation.cpp
View File

@@ -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 {

18
source/blender/compositor/operations/COM_GlareBaseOperation.h
View File

@@ -22,7 +22,8 @@
#ifndef _COM_GlareBaseOperation_h
#define _COM_GlareBaseOperation_h
#include "COM_NodeOperation.h"
#include "COM_SingleThreadedNodeOperation.h"
#include "DNA_node_types.h"
@@ -55,7 +56,7 @@ typedef float fRGB[4];
} (void)0
class GlareBaseOperation : public NodeOperation {
class GlareBaseOperation : public SingleThreadedNodeOperation {
private:
/**
* @brief Cached reference to the inputProgram
@@ -66,16 +67,7 @@ private:
* @brief settings of the glare node.
*/
NodeGlare * settings;
float *cachedInstance;
public:
/**
* the inner loop of this program
*/
void executePixel(float *color, int x, int y, MemoryBuffer *inputBuffers[], void *data);
/**
* Initialize the execution
*/
@@ -86,15 +78,15 @@ public:
*/
void deinitExecution();
void *initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers);
void setGlareSettings(NodeGlare * settings) {this->settings = settings;}
bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output);
protected:
GlareBaseOperation();
virtual void generateGlare(float *data, MemoryBuffer *inputTile, NodeGlare *settings) = 0;
MemoryBuffer *createMemoryBuffer(rcti *rect, MemoryBuffer **memoryBuffers);
};
#endif

25
source/blender/compositor/operations/COM_GlareGhostOperation.cpp
View File

@@ -45,15 +45,21 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No
MemoryBuffer *gbuf = inputTile->duplicate();
MemoryBuffer *tbuf1 = inputTile->duplicate();
bool breaked = false;
FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 0, 3);
FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 1, 3);
FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 2, 3);
if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 1, 3);
if (isBreaked()) breaked = true;
if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf1, s1, 2, 3);
MemoryBuffer *tbuf2 = tbuf1->duplicate();
FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 0, 3);
FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 1, 3);
FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 2, 3);
if (isBreaked()) breaked = true;
if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 0, 3);
if (isBreaked()) breaked = true;
if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 1, 3);
if (isBreaked()) breaked = true;
if (!breaked) FastGaussianBlurOperation::IIR_gauss(tbuf2, s2, 2, 3);
if (settings->iter & 1) ofs = 0.5f; else ofs = 0.f;
for (x=0; x<(settings->iter*4); x++) {
@@ -68,7 +74,7 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No
sc = 2.13;
isc = -0.97;
for (y=0; y<gbuf->getHeight(); y++) {
for (y=0; y<gbuf->getHeight() &(!breaked); y++) {
v = (float)(y+0.5f) / (float)gbuf->getHeight();
for (x=0; x<gbuf->getWidth(); x++) {
u = (float)(x+0.5f) / (float)gbuf->getWidth();
@@ -83,11 +89,13 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No
gbuf->writePixel(x, y, c);
}
if (isBreaked()) breaked = true;
}
memset(tbuf1->getBuffer(), 0, tbuf1->getWidth()*tbuf1->getHeight()*COM_NUMBER_OF_CHANNELS*sizeof(float));
for (n=1; n<settings->iter; n++) {
for (y=0; y<gbuf->getHeight(); y++) {
for (n=1; n<settings->iter &(!breaked); n++) {
for (y=0; y<gbuf->getHeight()&(!breaked); y++) {
v = (float)(y+0.5f) / (float)gbuf->getHeight();
for (x=0; x<gbuf->getWidth(); x++) {
u = (float)(x+0.5f) / (float)gbuf->getWidth();
@@ -103,6 +111,7 @@ void GlareGhostOperation::generateGlare(float *data, MemoryBuffer *inputTile, No
}
tbuf1->writePixel(x, y, tc);
}
if (isBreaked()) breaked = true;
}
memcpy(gbuf->getBuffer(), tbuf1->getBuffer(), tbuf1->getWidth()*tbuf1->getHeight()*COM_NUMBER_OF_CHANNELS*sizeof(float));
}

14
source/blender/compositor/operations/COM_GlareSimpleStarOperation.cpp
View File

@@ -32,10 +32,11 @@ void GlareSimpleStarOperation::generateGlare(float *data, MemoryBuffer *inputTil
MemoryBuffer *tbuf1 = inputTile->duplicate();
MemoryBuffer *tbuf2 = inputTile->duplicate();
for (i=0; i<settings->iter; i++) {
bool breaked = false;
for (i=0; i<settings->iter && (!breaked); i++) {
// // (x || x-1, y-1) to (x || x+1, y+1)
// // F
for (y=0; y<this->getHeight(); y++) {
for (y=0; y<this->getHeight() && (!breaked); y++) {
ym = y - i;
yp = y + i;
for (x=0; x<this->getWidth(); x++) {
@@ -58,11 +59,13 @@ void GlareSimpleStarOperation::generateGlare(float *data, MemoryBuffer *inputTil
madd_v3_v3fl(c, tc, f2);
c[3] = 1.0f;
tbuf2->writePixel(x, y, c);
}
if (isBreaked()) {
breaked = true;
}
}
// // B
for (y=tbuf1->getHeight()-1; y>=0; y--) {
for (y=tbuf1->getHeight()-1 && (!breaked); y>=0; y--) {
ym = y - i;
yp = y + i;
for (x=tbuf1->getWidth()-1; x>=0; x--) {
@@ -86,6 +89,9 @@ void GlareSimpleStarOperation::generateGlare(float *data, MemoryBuffer *inputTil
c[3] = 1.0f;
tbuf2->writePixel(x, y, c);
}
if (isBreaked()) {
breaked = true;
}
}
}

11
source/blender/compositor/operations/COM_GlareStreaksOperation.cpp
View File

@@ -33,22 +33,23 @@ void GlareStreaksOperation::generateGlare(float *data, MemoryBuffer *inputTile,
int size = inputTile->getWidth()*inputTile->getHeight();
int size4 = size*4;
bool breaked = false;
MemoryBuffer *tsrc = inputTile->duplicate();
MemoryBuffer *tdst = new MemoryBuffer(NULL, inputTile->getRect());
tdst->clear();
memset(data, 0, size4*sizeof(float));
for (a=0.f; a<DEG2RADF(360.0f); a+=ang) {
for (a=0.f; a<DEG2RADF(360.0f) && (!breaked); a+=ang) {
const float an = a + settings->angle_ofs;
const float vx = cos((double)an), vy = sin((double)an);
for (n=0; n<settings->iter; ++n) {
for (n=0; n<settings->iter && (!breaked); ++n) {
const float p4 = pow(4.0, (double)n);
const float vxp = vx*p4, vyp = vy*p4;
const float wt = pow((double)settings->fade, (double)p4);
const float cmo = 1.f - (float)pow((double)settings->colmod, (double)n+1); // colormodulation amount relative to current pass
float *tdstcol = tdst->getBuffer();
for (y=0; y<tsrc->getHeight(); ++y) {
for (y=0; y<tsrc->getHeight() && (!breaked); ++y) {
for (x=0; x<tsrc->getWidth(); ++x, tdstcol+=4) {
// first pass no offset, always same for every pass, exact copy,
// otherwise results in uneven brightness, only need once
@@ -71,11 +72,13 @@ void GlareStreaksOperation::generateGlare(float *data, MemoryBuffer *inputTile,
tdstcol[2] = 0.5f*(tdstcol[2] + c1[2] + wt*(c2[2] + wt*(c3[2] + wt*c4[2])));
tdstcol[3] = 1.0f;
}
if (isBreaked()) {
breaked = true;
}
}
memcpy(tsrc->getBuffer(), tdst->getBuffer(), sizeof(float)*size4);
}
// addImage(sbuf, tsrc, 1.f/(float)(6 - ndg->iter)); // add result to data @todo
float *sourcebuffer = tsrc->getBuffer();
float factor = 1.f/(float)(6 - settings->iter);
for (int i = 0 ; i < size4; i ++) {

5
source/blender/compositor/operations/COM_MaskOperation.cpp
View File

@@ -67,7 +67,7 @@ void *MaskOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers
if (!this->mask)
return NULL;
BLI_mutex_lock(getMutex());
lockMutex();
if (this->rasterizedMask == NULL) {
int width = this->getWidth();
int height = this->getHeight();
@@ -78,8 +78,7 @@ void *MaskOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers
this->rasterizedMask = buffer;
}
BLI_mutex_unlock(getMutex());
unlockMutex();
return this->rasterizedMask;
}

2
source/blender/compositor/operations/COM_MovieDistortionOperation.cpp
View File

@@ -52,7 +52,7 @@ void MovieDistortionOperation::initExecution()
BKE_movieclip_user_set_frame(&clipUser, this->framenumber);
BKE_movieclip_get_size(this->movieClip, &clipUser, &calibration_width, &calibration_height);
for (int i = 0 ; i < s_cache.size() ; i ++) {
for (unsigned int i = 0 ; i < s_cache.size() ; i ++) {
DistortionCache *c = (DistortionCache*)s_cache[i];
if (c->isCacheFor(this->movieClip, this->width, this->height,
calibration_width, calibration_height, this->distortion))

6
source/blender/compositor/operations/COM_MultilayerImageOperation.cpp
View File

@@ -47,7 +47,7 @@ void MultilayerColorOperation::executePixel(float *color, float x, float y, Pixe
{
int yi = y;
int xi = x;
if (this->imageBuffer == NULL || xi < 0 || yi < 0 || xi >= this->getWidth() || yi >= this->getHeight() ) {
if (this->imageBuffer == NULL || xi < 0 || yi < 0 || (unsigned int)xi >= this->getWidth() || (unsigned int)yi >= this->getHeight() ) {
color[0] = 0.0f;
color[1] = 0.0f;
color[2] = 0.0f;
@@ -80,7 +80,7 @@ void MultilayerValueOperation::executePixel(float *color, float x, float y, Pixe
{
int yi = y;
int xi = x;
if (this->imageBuffer == NULL || xi < 0 || yi < 0 || xi >= this->getWidth() || yi >= this->getHeight() ) {
if (this->imageBuffer == NULL || xi < 0 || yi < 0 || (unsigned int)xi >= this->getWidth() || (unsigned int)yi >= this->getHeight() ) {
color[0] = 0.0f;
}
else {
@@ -93,7 +93,7 @@ void MultilayerVectorOperation::executePixel(float *color, float x, float y, Pix
{
int yi = y;
int xi = x;
if (this->imageBuffer == NULL || xi < 0 || yi < 0 || xi >= this->getWidth() || yi >= this->getHeight() ) {
if (this->imageBuffer == NULL || xi < 0 || yi < 0 || (unsigned int)xi >= this->getWidth() || (unsigned int)yi >= this->getHeight() ) {
color[0] = 0.0f;
}
else {

5
source/blender/compositor/operations/COM_NormalizeOperation.cpp
View File

@@ -76,8 +76,7 @@ bool NormalizeOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe
void *NormalizeOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
{
BLI_mutex_lock(getMutex());
lockMutex();
if (this->cachedInstance == NULL) {
MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers);
/* using generic two floats struct to store x: min y: mult */
@@ -105,7 +104,7 @@ void *NormalizeOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBu
this->cachedInstance = minmult;
}
BLI_mutex_unlock(getMutex());
unlockMutex();
return this->cachedInstance;
}

66
source/blender/compositor/operations/COM_OpenCLKernels.cl
View File

@@ -6,8 +6,8 @@ const sampler_t SAMPLER_NEAREST = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS
__constant const int2 zero = {0,0};
// KERNEL --- BOKEH BLUR ---
__kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __global __read_only image2d_t inputImage,
__global __read_only image2d_t bokehImage, __global __write_only image2d_t output,
__kernel void bokehBlurKernel(__read_only image2d_t boundingBox, __read_only image2d_t inputImage,
__read_only image2d_t bokehImage, __write_only image2d_t output,
int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset)
{
int2 coords = {get_global_id(0), get_global_id(1)};
@@ -50,3 +50,65 @@ __kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __glob
write_imagef(output, coords, color);
}
// KERNEL --- DILATE ---
__kernel void dilateKernel(__read_only image2d_t inputImage, __write_only image2d_t output,
int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,
int2 offset)
{
int2 coords = {get_global_id(0), get_global_id(1)};
coords += offset;
const int2 realCoordinate = coords + offsetOutput;
const int2 minXY = max(realCoordinate - scope, zero);
const int2 maxXY = min(realCoordinate + scope, dimension);
float value = 0.0f;
int nx, ny;
int2 inputXy;
for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {
for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {
const float deltaX = (realCoordinate.x - nx);
const float deltaY = (realCoordinate.y - ny);
const float measuredDistance = deltaX*deltaX+deltaY*deltaY;
if (measuredDistance <= distanceSquared) {
value = max(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);
}
}
}
float4 color = {value,0.0f,0.0f,0.0f};
write_imagef(output, coords, color);
}
// KERNEL --- DILATE ---
__kernel void erodeKernel(__read_only image2d_t inputImage, __write_only image2d_t output,
int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,
int2 offset)
{
int2 coords = {get_global_id(0), get_global_id(1)};
coords += offset;
const int2 realCoordinate = coords + offsetOutput;
const int2 minXY = max(realCoordinate - scope, zero);
const int2 maxXY = min(realCoordinate + scope, dimension);
float value = 1.0f;
int nx, ny;
int2 inputXy;
for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {
for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {
const float deltaX = (realCoordinate.x - nx);
const float deltaY = (realCoordinate.y - ny);
const float measuredDistance = deltaX*deltaX+deltaY*deltaY;
if (measuredDistance <= distanceSquared) {
value = min(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);
}
}
}
float4 color = {value,0.0f,0.0f,0.0f};
write_imagef(output, coords, color);
}

70
source/blender/compositor/operations/COM_OpenCLKernels.cl.h
View File

@@ -8,8 +8,8 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope
"__constant const int2 zero = {0,0};\n" \
"\n" \
"// KERNEL --- BOKEH BLUR ---\n" \
"__kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __global __read_only image2d_t inputImage,\n" \
" __global __read_only image2d_t bokehImage, __global __write_only image2d_t output,\n" \
"__kernel void bokehBlurKernel(__read_only image2d_t boundingBox, __read_only image2d_t inputImage,\n" \
" __read_only image2d_t bokehImage, __write_only image2d_t output,\n" \
" int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset)\n" \
"{\n" \
" int2 coords = {get_global_id(0), get_global_id(1)};\n" \
@@ -26,8 +26,8 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope
" if (tempBoundingBox > 0.0f) {\n" \
" const int2 bokehImageDim = get_image_dim(bokehImage);\n" \
" const int2 bokehImageCenter = bokehImageDim/2;\n" \
" const int2 minXY = max(realCoordinate - radius, zero);;\n" \
" const int2 maxXY = min(realCoordinate + radius, dimension);;\n" \
" const int2 minXY = max(realCoordinate - radius, zero);\n" \
" const int2 maxXY = min(realCoordinate + radius, dimension);\n" \
" int nx, ny;\n" \
"\n" \
" float2 uv;\n" \
@@ -52,4 +52,66 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope
"\n" \
" write_imagef(output, coords, color);\n" \
"}\n" \
"\n" \
"// KERNEL --- DILATE ---\n" \
"__kernel void dilateKernel(__read_only image2d_t inputImage, __write_only image2d_t output,\n" \
" int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,\n" \
" int2 offset)\n" \
"{\n" \
" int2 coords = {get_global_id(0), get_global_id(1)};\n" \
" coords += offset;\n" \
" const int2 realCoordinate = coords + offsetOutput;\n" \
"\n" \
" const int2 minXY = max(realCoordinate - scope, zero);\n" \
" const int2 maxXY = min(realCoordinate + scope, dimension);\n" \
"\n" \
" float value = 0.0f;\n" \
" int nx, ny;\n" \
" int2 inputXy;\n" \
"\n" \
" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\n" \
" for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {\n" \
" const float deltaX = (realCoordinate.x - nx);\n" \
" const float deltaY = (realCoordinate.y - ny);\n" \
" const float measuredDistance = deltaX*deltaX+deltaY*deltaY;\n" \
" if (measuredDistance <= distanceSquared) {\n" \
" value = max(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \
" }\n" \
" }\n" \
" }\n" \
"\n" \
" float4 color = {value,0.0f,0.0f,0.0f};\n" \
" write_imagef(output, coords, color);\n" \
"}\n" \
"\n" \
"// KERNEL --- DILATE ---\n" \
"__kernel void erodeKernel(__read_only image2d_t inputImage, __write_only image2d_t output,\n" \
" int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,\n" \
" int2 offset)\n" \
"{\n" \
" int2 coords = {get_global_id(0), get_global_id(1)};\n" \
" coords += offset;\n" \
" const int2 realCoordinate = coords + offsetOutput;\n" \
"\n" \
" const int2 minXY = max(realCoordinate - scope, zero);\n" \
" const int2 maxXY = min(realCoordinate + scope, dimension);\n" \
"\n" \
" float value = 1.0f;\n" \
" int nx, ny;\n" \
" int2 inputXy;\n" \
"\n" \
" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\n" \
" for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {\n" \
" const float deltaX = (realCoordinate.x - nx);\n" \
" const float deltaY = (realCoordinate.y - ny);\n" \
" const float measuredDistance = deltaX*deltaX+deltaY*deltaY;\n" \
" if (measuredDistance <= distanceSquared) {\n" \
" value = min(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \
" }\n" \
" }\n" \
" }\n" \
"\n" \
" float4 color = {value,0.0f,0.0f,0.0f};\n" \
" write_imagef(output, coords, color);\n" \
"}\n" \
"\0";

8
source/blender/compositor/operations/COM_OutputFileOperation.cpp
View File

@@ -177,7 +177,7 @@ void OutputOpenExrMultiLayerOperation::add_layer(const char *name, DataType data
void OutputOpenExrMultiLayerOperation::initExecution()
{
for (int i=0; i < layers.size(); ++i) {
for (unsigned int i=0; i < layers.size(); ++i) {
layers[i].imageInput = getInputSocketReader(i);
layers[i].outputBuffer = init_buffer(this->getWidth(), this->getHeight(), layers[i].datatype);
}
@@ -185,7 +185,7 @@ void OutputOpenExrMultiLayerOperation::initExecution()
void OutputOpenExrMultiLayerOperation::executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer** memoryBuffers)
{
for (int i=0; i < layers.size(); ++i) {
for (unsigned int i=0; i < layers.size(); ++i) {
write_buffer_rect(rect, memoryBuffers, this->tree, layers[i].imageInput, layers[i].outputBuffer, this->getWidth(), layers[i].datatype);
}
}
@@ -203,7 +203,7 @@ void OutputOpenExrMultiLayerOperation::deinitExecution()
(this->scene->r.scemode & R_EXTENSION), true);
BLI_make_existing_file(filename);
for (int i=0; i < layers.size(); ++i) {
for (unsigned int i=0; i < layers.size(); ++i) {
char channelname[EXR_TOT_MAXNAME];
BLI_strncpy(channelname, layers[i].name, sizeof(channelname)-2);
char *channelname_ext = channelname + strlen(channelname);
@@ -251,7 +251,7 @@ void OutputOpenExrMultiLayerOperation::deinitExecution()
}
IMB_exr_close(exrhandle);
for (int i=0; i < layers.size(); ++i) {
for (unsigned int i=0; i < layers.size(); ++i) {
if (layers[i].outputBuffer) {
MEM_freeN(layers[i].outputBuffer);
layers[i].outputBuffer = NULL;

2
source/blender/compositor/operations/COM_PreviewOperation.h
View File

@@ -34,7 +34,6 @@ protected:
* @brief holds reference to the SDNA bNode, where this nodes will render the preview image for
*/
bNode *node;
const bNodeTree *tree;
SocketReader *input;
float divider;
@@ -48,7 +47,6 @@ public:
void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer **memoryBuffers);
void determineResolution(unsigned int resolution[], unsigned int preferredResolution[]);
void setbNode(bNode *node) { this->node = node;}
void setbNodeTree(const bNodeTree *tree) { this->tree = tree;}
bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output);
bool isPreviewOperation() {return true;}

4
source/blender/compositor/operations/COM_TonemapOperation.cpp
View File

@@ -118,7 +118,7 @@ bool TonemapOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferO
void *TonemapOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuffers)
{
BLI_mutex_lock(getMutex());
lockMutex();
if (this->cachedInstance == NULL) {
MemoryBuffer *tile = (MemoryBuffer*)imageReader->initializeTileData(rect, memoryBuffers);
AvgLogLum *data = new AvgLogLum();
@@ -150,7 +150,7 @@ void *TonemapOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryBuff
data->igm = (this->data->gamma==0.f) ? 1 : (1.f / this->data->gamma);
this->cachedInstance = data;
}
BLI_mutex_unlock(getMutex());
unlockMutex();
return this->cachedInstance;
}

4
source/blender/compositor/operations/COM_VectorBlurOperation.cpp
View File

@@ -77,7 +77,7 @@ void *VectorBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB
{
if (this->cachedInstance) return this->cachedInstance;
BLI_mutex_lock(getMutex());
lockMutex();
if (this->cachedInstance == NULL) {
MemoryBuffer *tile = (MemoryBuffer*)inputImageProgram->initializeTileData(rect, memoryBuffers);
MemoryBuffer *speed = (MemoryBuffer*)inputSpeedProgram->initializeTileData(rect, memoryBuffers);
@@ -87,7 +87,7 @@ void *VectorBlurOperation::initializeTileData(rcti *rect, MemoryBuffer **memoryB
this->generateVectorBlur(data, tile, speed, z);
this->cachedInstance = data;
}
BLI_mutex_unlock(getMutex());
unlockMutex();
return this->cachedInstance;
}

2
source/blender/compositor/operations/COM_ViewerBaseOperation.h
View File

@@ -34,7 +34,6 @@ protected:
ImageUser * imageUser;
void *lock;
bool active;
const bNodeTree *tree;
float centerX;
float centerY;
OrderOfChunks chunkOrder;
@@ -49,7 +48,6 @@ public:
void setImageUser(ImageUser *imageUser) {this->imageUser = imageUser;}
const bool isActiveViewerOutput() const {return active;}
void setActive(bool active) {this->active = active;}
void setbNodeTree(const bNodeTree *tree) {this->tree = tree;}
void setCenterX(float centerX) {this->centerX = centerX;}
void setCenterY(float centerY) {this->centerY = centerY;}
void setChunkOrder(OrderOfChunks tileOrder) {this->chunkOrder = tileOrder;}

2
source/blender/compositor/operations/COM_ViewerOperation.cpp
View File

@@ -104,7 +104,7 @@ void ViewerOperation::executeRegion(rcti *rect, unsigned int tileNumber, MemoryB
offset +=4;
}
if (tree->test_break && tree->test_break(tree->tbh)) {
if (isBreaked()) {
breaked = true;
}

5
source/blender/compositor/operations/COM_WriteBufferOperation.cpp
View File

@@ -30,7 +30,6 @@ WriteBufferOperation::WriteBufferOperation() :NodeOperation()
this->memoryProxy = new MemoryProxy();
this->memoryProxy->setWriteBufferOperation(this);
this->memoryProxy->setExecutor(NULL);
this->tree = NULL;
}
WriteBufferOperation::~WriteBufferOperation()
{
@@ -78,7 +77,7 @@ void WriteBufferOperation::executeRegion(rcti *rect, unsigned int tileNumber, Me
offset4 +=COM_NUMBER_OF_CHANNELS;
}
if (tree->test_break && tree->test_break(tree->tbh)) {
if (isBreaked()) {
breaked = true;
}
@@ -103,7 +102,7 @@ void WriteBufferOperation::executeRegion(rcti *rect, unsigned int tileNumber, Me
input->read(&(buffer[offset4]), x, y, COM_PS_NEAREST, memoryBuffers);
offset4 +=COM_NUMBER_OF_CHANNELS;
}
if (tree->test_break && tree->test_break(tree->tbh)) {
if (isBreaked()) {
breaked = true;
}
}

2
source/blender/compositor/operations/COM_WriteBufferOperation.h
View File

@@ -33,7 +33,6 @@
class WriteBufferOperation: public NodeOperation {
MemoryProxy *memoryProxy;
NodeOperation *input;
const bNodeTree * tree;
public:
WriteBufferOperation();
~WriteBufferOperation();
@@ -45,7 +44,6 @@ public:
void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer** memoryBuffers);
void initExecution();
void deinitExecution();
void setbNodeTree(const bNodeTree *tree) {this->tree = tree;}
void executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer** memoryBuffers, MemoryBuffer* outputBuffer);
void readResolutionFromInputSocket();