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