diff options
Diffstat (limited to 'source/blender/compositor/intern')
20 files changed, 335 insertions, 326 deletions
diff --git a/source/blender/compositor/intern/COM_CompositorContext.cpp b/source/blender/compositor/intern/COM_CompositorContext.cpp index bb8e7d9606d..c3470f0a16e 100644 --- a/source/blender/compositor/intern/COM_CompositorContext.cpp +++ b/source/blender/compositor/intern/COM_CompositorContext.cpp @@ -26,7 +26,7 @@ CompositorContext::CompositorContext() { - this->scene = NULL; + this->rd = NULL; this->quality = COM_QUALITY_HIGH; this->hasActiveOpenCLDevices = false; this->activegNode = NULL; @@ -34,8 +34,8 @@ CompositorContext::CompositorContext() const int CompositorContext::getFramenumber() const { - if (this->scene) { - return this->scene->r.cfra; + if (this->rd) { + return this->rd->cfra; } else { return -1; /* this should never happen */ @@ -44,8 +44,8 @@ const int CompositorContext::getFramenumber() const const int CompositorContext::isColorManaged() const { - if (this->scene) { - return this->scene->r.color_mgt_flag & R_COLOR_MANAGEMENT; + if (this->rd) { + return this->rd->color_mgt_flag & R_COLOR_MANAGEMENT; } else { return 0; /* this should never happen */ diff --git a/source/blender/compositor/intern/COM_CompositorContext.h b/source/blender/compositor/intern/COM_CompositorContext.h index 93872f4839f..81fd81b4117 100644 --- a/source/blender/compositor/intern/COM_CompositorContext.h +++ b/source/blender/compositor/intern/COM_CompositorContext.h @@ -51,11 +51,11 @@ private: CompositorQuality quality; /** - * @brief Reference to the scene that is being composited. + * @brief Reference to the render data that is being composited. * This field is initialized in ExecutionSystem and must only be read from that point on. * @see ExecutionSystem */ - Scene *scene; + RenderData *rd; /** * @brief reference to the bNodeTree @@ -93,7 +93,7 @@ public: /** * @brief set the scene of the context */ - void setScene(Scene *scene) { this->scene = scene; } + void setRenderData(RenderData *rd) { this->rd = rd; } /** * @brief set the bnodetree of the context @@ -118,7 +118,7 @@ public: /** * @brief get the scene of the context */ - const Scene *getScene() const { return this->scene; } + const RenderData *getRenderData() const { return this->rd; } /** * @brief set the quality diff --git a/source/blender/compositor/intern/COM_Converter.cpp b/source/blender/compositor/intern/COM_Converter.cpp index fa30d965938..0a8862e4017 100644 --- a/source/blender/compositor/intern/COM_Converter.cpp +++ b/source/blender/compositor/intern/COM_Converter.cpp @@ -20,100 +20,102 @@ * Monique Dewanchand */ -#include "COM_Converter.h" +#include <string.h> + #include "BKE_node.h" -#include "COM_CompositorNode.h" -#include "COM_RenderLayersNode.h" + +#include "COM_AlphaOverNode.h" +#include "COM_BilateralBlurNode.h" +#include "COM_BlurNode.h" +#include "COM_BokehBlurNode.h" +#include "COM_BokehImageNode.h" +#include "COM_BoxMaskNode.h" +#include "COM_BrightnessNode.h" +#include "COM_ChannelMatteNode.h" +#include "COM_ChromaMatteNode.h" +#include "COM_ColorBalanceNode.h" +#include "COM_ColorCorrectionNode.h" +#include "COM_ColorCurveNode.h" +#include "COM_ColorMatteNode.h" +#include "COM_ColorNode.h" +#include "COM_ColorRampNode.h" +#include "COM_ColorSpillNode.h" #include "COM_ColorToBWNode.h" -#include "string.h" -#include "COM_SocketConnection.h" +#include "COM_CombineHSVANode.h" +#include "COM_CombineRGBANode.h" +#include "COM_CombineYCCANode.h" +#include "COM_CombineYUVANode.h" +#include "COM_CompositorNode.h" +#include "COM_ConvertAlphaNode.h" +#include "COM_ConvertColorToVectorOperation.h" #include "COM_ConvertColourToValueProg.h" #include "COM_ConvertValueToColourProg.h" -#include "COM_ConvertColorToVectorOperation.h" #include "COM_ConvertValueToVectorOperation.h" #include "COM_ConvertVectorToColorOperation.h" #include "COM_ConvertVectorToValueOperation.h" +#include "COM_Converter.h" +#include "COM_CropNode.h" +#include "COM_DefocusNode.h" +#include "COM_DifferenceMatteNode.h" +#include "COM_DilateErodeNode.h" +#include "COM_DirectionalBlurNode.h" +#include "COM_DisplaceNode.h" +#include "COM_DistanceMatteNode.h" +#include "COM_DoubleEdgeMaskNode.h" +#include "COM_EllipseMaskNode.h" #include "COM_ExecutionSystem.h" +#include "COM_ExecutionSystemHelper.h" +#include "COM_FilterNode.h" +#include "COM_FlipNode.h" +#include "COM_GammaNode.h" +#include "COM_GlareNode.h" +#include "COM_GroupNode.h" +#include "COM_HueSaturationValueCorrectNode.h" +#include "COM_HueSaturationValueNode.h" +#include "COM_IDMaskNode.h" +#include "COM_ImageNode.h" +#include "COM_InvertNode.h" +#include "COM_KeyingNode.h" +#include "COM_KeyingScreenNode.h" +#include "COM_LensDistortionNode.h" +#include "COM_LuminanceMatteNode.h" +#include "COM_MapUVNode.h" +#include "COM_MapValueNode.h" +#include "COM_MaskNode.h" +#include "COM_MathNode.h" #include "COM_MixNode.h" +#include "COM_MovieClipNode.h" +#include "COM_MovieDistortionNode.h" #include "COM_MuteNode.h" -#include "COM_TranslateNode.h" +#include "COM_NormalNode.h" +#include "COM_NormalizeNode.h" +#include "COM_OutputFileNode.h" +#include "COM_RenderLayersNode.h" #include "COM_RotateNode.h" #include "COM_ScaleNode.h" -#include "COM_FlipNode.h" -#include "COM_IDMaskNode.h" -#include "COM_FilterNode.h" -#include "COM_BrightnessNode.h" -#include "COM_SeparateRGBANode.h" -#include "COM_CombineRGBANode.h" +#include "COM_ScaleOperation.h" #include "COM_SeparateHSVANode.h" -#include "COM_CombineHSVANode.h" -#include "COM_SeparateYUVANode.h" -#include "COM_CombineYUVANode.h" +#include "COM_SeparateRGBANode.h" #include "COM_SeparateYCCANode.h" -#include "COM_CombineYCCANode.h" -#include "COM_AlphaOverNode.h" -#include "COM_ColorBalanceNode.h" -#include "COM_ViewerNode.h" -#include "COM_SplitViewerNode.h" -#include "COM_InvertNode.h" -#include "COM_GroupNode.h" -#include "COM_NormalNode.h" -#include "COM_NormalizeNode.h" -#include "COM_ImageNode.h" -#include "COM_BokehImageNode.h" -#include "COM_ColorCurveNode.h" -#include "COM_VectorCurveNode.h" +#include "COM_SeparateYUVANode.h" #include "COM_SetAlphaNode.h" -#include "COM_ConvertAlphaNode.h" -#include "COM_MapUVNode.h" -#include "COM_DisplaceNode.h" -#include "COM_MathNode.h" -#include "COM_HueSaturationValueNode.h" -#include "COM_HueSaturationValueCorrectNode.h" -#include "COM_ColorCorrectionNode.h" -#include "COM_BoxMaskNode.h" -#include "COM_EllipseMaskNode.h" -#include "COM_GammaNode.h" -#include "COM_ColorRampNode.h" -#include "COM_DifferenceMatteNode.h" -#include "COM_LuminanceMatteNode.h" -#include "COM_DistanceMatteNode.h" -#include "COM_ChromaMatteNode.h" -#include "COM_ColorMatteNode.h" -#include "COM_ChannelMatteNode.h" -#include "COM_BlurNode.h" -#include "COM_BokehBlurNode.h" -#include "COM_DilateErodeNode.h" -#include "COM_TranslateOperation.h" -#include "COM_LensDistortionNode.h" +#include "COM_SetValueOperation.h" +#include "COM_SocketConnection.h" +#include "COM_SplitViewerNode.h" +#include "COM_Stabilize2dNode.h" +#include "COM_SwitchNode.h" #include "COM_TextureNode.h" -#include "COM_ColorNode.h" -#include "COM_ValueNode.h" #include "COM_TimeNode.h" -#include "COM_DirectionalBlurNode.h" -#include "COM_ZCombineNode.h" -#include "COM_SetValueOperation.h" -#include "COM_ScaleOperation.h" -#include "COM_ExecutionSystemHelper.h" #include "COM_TonemapNode.h" -#include "COM_SwitchNode.h" -#include "COM_GlareNode.h" -#include "COM_MovieClipNode.h" -#include "COM_ColorSpillNode.h" -#include "COM_OutputFileNode.h" -#include "COM_MapValueNode.h" #include "COM_TransformNode.h" -#include "COM_Stabilize2dNode.h" -#include "COM_BilateralBlurNode.h" +#include "COM_TranslateNode.h" +#include "COM_TranslateOperation.h" +#include "COM_ValueNode.h" #include "COM_VectorBlurNode.h" -#include "COM_MovieDistortionNode.h" +#include "COM_VectorCurveNode.h" #include "COM_ViewLevelsNode.h" -#include "COM_DefocusNode.h" -#include "COM_DoubleEdgeMaskNode.h" -#include "COM_CropNode.h" -#include "COM_MaskNode.h" -#include "COM_KeyingScreenNode.h" -#include "COM_KeyingNode.h" +#include "COM_ViewerNode.h" +#include "COM_ZCombineNode.h" Node *Converter::convert(bNode *bNode) { diff --git a/source/blender/compositor/intern/COM_Device.h b/source/blender/compositor/intern/COM_Device.h index 08fdb5bb578..2a86382a191 100644 --- a/source/blender/compositor/intern/COM_Device.h +++ b/source/blender/compositor/intern/COM_Device.h @@ -23,11 +23,7 @@ #ifndef _COM_Device_h #define _COM_Device_h -#include "COM_ExecutionSystem.h" #include "COM_WorkPackage.h" -#include "COM_NodeOperation.h" -#include "BLI_rect.h" -#include "COM_MemoryBuffer.h" /** * @brief Abstract class for device implementations to be used by the Compositor. diff --git a/source/blender/compositor/intern/COM_ExecutionGroup.cpp b/source/blender/compositor/intern/COM_ExecutionGroup.cpp index 2a790da0354..4dfb9c7d26c 100644 --- a/source/blender/compositor/intern/COM_ExecutionGroup.cpp +++ b/source/blender/compositor/intern/COM_ExecutionGroup.cpp @@ -20,24 +20,25 @@ * Monique Dewanchand */ +#include <algorithm> +#include <math.h> +#include <sstream> +#include <stdlib.h> + +#include "BLI_math.h" +#include "PIL_time.h" + #include "COM_ExecutionGroup.h" #include "COM_InputSocket.h" #include "COM_SocketConnection.h" #include "COM_defines.h" -#include "math.h" #include "COM_ExecutionSystem.h" -#include <sstream> #include "COM_ReadBufferOperation.h" #include "COM_WriteBufferOperation.h" #include "COM_ReadBufferOperation.h" #include "COM_WorkScheduler.h" #include "COM_ViewerOperation.h" -#include <stdlib.h> -#include "BLI_math.h" -#include "PIL_time.h" #include "COM_ChunkOrder.h" -#include <algorithm> -#include "BLI_math.h" #include "COM_ExecutionSystemHelper.h" ExecutionGroup::ExecutionGroup() diff --git a/source/blender/compositor/intern/COM_ExecutionSystem.cpp b/source/blender/compositor/intern/COM_ExecutionSystem.cpp index 7250a851f7b..7e09486fd0b 100644 --- a/source/blender/compositor/intern/COM_ExecutionSystem.cpp +++ b/source/blender/compositor/intern/COM_ExecutionSystem.cpp @@ -22,16 +22,18 @@ #include "COM_ExecutionSystem.h" +#include <sstream> +#include <stdio.h> + #include "PIL_time.h" #include "BKE_node.h" + #include "COM_Converter.h" -#include <sstream> #include "COM_NodeOperation.h" #include "COM_ExecutionGroup.h" #include "COM_NodeBase.h" #include "COM_WorkScheduler.h" #include "COM_ReadBufferOperation.h" -#include "stdio.h" #include "COM_GroupNode.h" #include "COM_WriteBufferOperation.h" #include "COM_ReadBufferOperation.h" @@ -39,7 +41,7 @@ #include "BKE_global.h" -ExecutionSystem::ExecutionSystem(bNodeTree *editingtree, bool rendering) +ExecutionSystem::ExecutionSystem(RenderData *rd, bNodeTree *editingtree, bool rendering) { context.setbNodeTree(editingtree); bNode *gnode; @@ -60,22 +62,18 @@ ExecutionSystem::ExecutionSystem(bNodeTree *editingtree, bool rendering) context.setRendering(rendering); context.setHasActiveOpenCLDevices(WorkScheduler::hasGPUDevices() && (editingtree->flag & NTREE_COM_OPENCL)); - Node *mainOutputNode = NULL; + ExecutionSystemHelper::addbNodeTree(*this, 0, editingtree, NULL); - mainOutputNode = ExecutionSystemHelper::addbNodeTree(*this, 0, editingtree, NULL); - - if (mainOutputNode) { - context.setScene((Scene *)mainOutputNode->getbNode()->id); - this->convertToOperations(); - this->groupOperations(); /* group operations in ExecutionGroups */ - unsigned int index; - unsigned int resolution[2]; - for (index = 0; index < this->groups.size(); index++) { - resolution[0] = 0; - resolution[1] = 0; - ExecutionGroup *executionGroup = groups[index]; - executionGroup->determineResolution(resolution); - } + context.setRenderData(rd); + this->convertToOperations(); + this->groupOperations(); /* group operations in ExecutionGroups */ + unsigned int index; + unsigned int resolution[2]; + for (index = 0; index < this->groups.size(); index++) { + resolution[0] = 0; + resolution[1] = 0; + ExecutionGroup *executionGroup = groups[index]; + executionGroup->determineResolution(resolution); } #ifdef COM_DEBUG diff --git a/source/blender/compositor/intern/COM_ExecutionSystem.h b/source/blender/compositor/intern/COM_ExecutionSystem.h index 70fd94ca57f..48ff2ef6af9 100644 --- a/source/blender/compositor/intern/COM_ExecutionSystem.h +++ b/source/blender/compositor/intern/COM_ExecutionSystem.h @@ -156,7 +156,7 @@ public: * @param editingtree [bNodeTree*] * @param rendering [true false] */ - ExecutionSystem(bNodeTree *editingtree, bool rendering); + ExecutionSystem(RenderData *rd, bNodeTree *editingtree, bool rendering); /** * Destructor diff --git a/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp b/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp index bcb606316ab..e5376567077 100644 --- a/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp +++ b/source/blender/compositor/intern/COM_ExecutionSystemHelper.cpp @@ -22,36 +22,34 @@ #include "COM_ExecutionSystemHelper.h" +#include <sstream> +#include <stdio.h> + #include "PIL_time.h" #include "BKE_node.h" + #include "COM_Converter.h" -#include <sstream> #include "COM_NodeOperation.h" #include "COM_ExecutionGroup.h" #include "COM_NodeBase.h" #include "COM_WorkScheduler.h" #include "COM_ReadBufferOperation.h" -#include "stdio.h" #include "COM_GroupNode.h" #include "COM_WriteBufferOperation.h" #include "COM_ReadBufferOperation.h" #include "COM_ViewerBaseOperation.h" -Node *ExecutionSystemHelper::addbNodeTree(ExecutionSystem &system, int nodes_start, bNodeTree *tree, bNode *groupnode) +void ExecutionSystemHelper::addbNodeTree(ExecutionSystem &system, int nodes_start, bNodeTree *tree, bNode *groupnode) { vector<Node *>& nodes = system.getNodes(); vector<SocketConnection *>& links = system.getConnections(); - Node *mainnode = NULL; const bNode *activeGroupNode = system.getContext().getActivegNode(); bool isActiveGroup = activeGroupNode == groupnode; /* add all nodes of the tree to the node list */ bNode *node = (bNode *)tree->nodes.first; while (node != NULL) { - Node *execnode = addNode(nodes, node, isActiveGroup); - if (node->type == CMP_NODE_COMPOSITE) { - mainnode = execnode; - } + addNode(nodes, node, isActiveGroup); node = (bNode *)node->next; } @@ -72,8 +70,6 @@ Node *ExecutionSystemHelper::addbNodeTree(ExecutionSystem &system, int nodes_sta groupNode->ungroup(system); } } - - return mainnode; } void ExecutionSystemHelper::addNode(vector<Node *>& nodes, Node *node) diff --git a/source/blender/compositor/intern/COM_ExecutionSystemHelper.h b/source/blender/compositor/intern/COM_ExecutionSystemHelper.h index 99a05472075..8dbd308153b 100644 --- a/source/blender/compositor/intern/COM_ExecutionSystemHelper.h +++ b/source/blender/compositor/intern/COM_ExecutionSystemHelper.h @@ -48,7 +48,7 @@ public: * @param tree bNodeTree to add * @return Node representing the "Compositor node" of the maintree. or NULL when a subtree is added */ - static Node *addbNodeTree(ExecutionSystem &system, int nodes_start, bNodeTree *tree, bNode *groupnode); + static void addbNodeTree(ExecutionSystem &system, int nodes_start, bNodeTree *tree, bNode *groupnode); /** * @brief add an editor node to the system. diff --git a/source/blender/compositor/intern/COM_Node.cpp b/source/blender/compositor/intern/COM_Node.cpp index 62e030b777c..06b6164be3c 100644 --- a/source/blender/compositor/intern/COM_Node.cpp +++ b/source/blender/compositor/intern/COM_Node.cpp @@ -20,11 +20,12 @@ * Monique Dewanchand */ -#include "COM_Node.h" -#include "string.h" +#include <string.h> -#include "COM_NodeOperation.h" #include "BKE_node.h" + +#include "COM_Node.h" +#include "COM_NodeOperation.h" #include "COM_SetValueOperation.h" #include "COM_SetVectorOperation.h" #include "COM_SetColorOperation.h" @@ -35,7 +36,7 @@ #include "COM_SocketProxyNode.h" -//#include "stdio.h" +//#include <stdio.h> #include "COM_defines.h" Node::Node(bNode *editorNode, bool create_sockets) @@ -86,11 +87,15 @@ void Node::addSetValueOperation(ExecutionSystem *graph, InputSocket *inputsocket void Node::addPreviewOperation(ExecutionSystem *system, OutputSocket *outputSocket) { if (this->isInActiveGroup()) { - PreviewOperation *operation = new PreviewOperation(); - system->addOperation(operation); - operation->setbNode(this->getbNode()); - operation->setbNodeTree(system->getContext().getbNodeTree()); - this->addLink(system, outputSocket, operation->getInputSocket(0)); + if (!(this->getbNode()->flag & NODE_HIDDEN)) { // do not calculate previews of hidden nodes. + if (this->getbNode()->flag & NODE_PREVIEW) { + PreviewOperation *operation = new PreviewOperation(); + system->addOperation(operation); + operation->setbNode(this->getbNode()); + operation->setbNodeTree(system->getContext().getbNodeTree()); + this->addLink(system, outputSocket, operation->getInputSocket(0)); + } + } } } diff --git a/source/blender/compositor/intern/COM_Node.h b/source/blender/compositor/intern/COM_Node.h index 12baa26cd6e..090b1455440 100644 --- a/source/blender/compositor/intern/COM_Node.h +++ b/source/blender/compositor/intern/COM_Node.h @@ -29,6 +29,7 @@ #include "COM_CompositorContext.h" #include "DNA_node_types.h" #include "BKE_text.h" +#include "COM_ExecutionSystem.h" #include <vector> #include <string> diff --git a/source/blender/compositor/intern/COM_NodeBase.cpp b/source/blender/compositor/intern/COM_NodeBase.cpp index 42946d7315e..1a895cf93b1 100644 --- a/source/blender/compositor/intern/COM_NodeBase.cpp +++ b/source/blender/compositor/intern/COM_NodeBase.cpp @@ -20,10 +20,12 @@ * Monique Dewanchand */ +#include <string.h> + +#include "BKE_node.h" + #include "COM_NodeBase.h" -#include "string.h" #include "COM_NodeOperation.h" -#include "BKE_node.h" #include "COM_SetValueOperation.h" #include "COM_SetColorOperation.h" #include "COM_SocketConnection.h" @@ -81,12 +83,14 @@ const bool NodeBase::isInputNode() const return this->inputsockets.size() == 0; } -OutputSocket *NodeBase::getOutputSocket(int index) +OutputSocket *NodeBase::getOutputSocket(unsigned int index) { + BLI_assert(index < this->outputsockets.size()); return this->outputsockets[index]; } -InputSocket *NodeBase::getInputSocket(int index) +InputSocket *NodeBase::getInputSocket(unsigned int index) { + BLI_assert(index < this->inputsockets.size()); return this->inputsockets[index]; } diff --git a/source/blender/compositor/intern/COM_NodeBase.h b/source/blender/compositor/intern/COM_NodeBase.h index 5e3a4fa5531..54f80926b84 100644 --- a/source/blender/compositor/intern/COM_NodeBase.h +++ b/source/blender/compositor/intern/COM_NodeBase.h @@ -103,7 +103,7 @@ public: * @param index * the index of the needed outputsocket */ - OutputSocket *getOutputSocket(const int index); + OutputSocket *getOutputSocket(const unsigned int index); /** * get the reference to the first outputsocket @@ -117,7 +117,7 @@ public: * @param index * the index of the needed inputsocket */ - InputSocket *getInputSocket(const int index); + InputSocket *getInputSocket(const unsigned int index); virtual bool isStatic() const { return false; } diff --git a/source/blender/compositor/intern/COM_NodeOperation.cpp b/source/blender/compositor/intern/COM_NodeOperation.cpp index 114d9f44cef..33989fa5787 100644 --- a/source/blender/compositor/intern/COM_NodeOperation.cpp +++ b/source/blender/compositor/intern/COM_NodeOperation.cpp @@ -20,12 +20,13 @@ * Monique Dewanchand */ -#include "COM_NodeOperation.h" #include <typeinfo> +#include <stdio.h> + +#include "COM_NodeOperation.h" #include "COM_InputSocket.h" #include "COM_SocketConnection.h" #include "COM_defines.h" -#include "stdio.h" NodeOperation::NodeOperation() { @@ -139,116 +140,3 @@ bool NodeOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferOper 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; - - 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 && (!breaked); 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); - if (isBreaked()) { - breaked = false; - } - } - } -} - -cl_kernel NodeOperation::COM_clCreateKernel(cl_program program, const char *kernelname, list<cl_kernel> *clKernelsToCleanUp) -{ - cl_int error; - cl_kernel kernel = clCreateKernel(program, kernelname, &error); - if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } - else { - if (clKernelsToCleanUp) clKernelsToCleanUp->push_back(kernel); - } - return kernel; - -} diff --git a/source/blender/compositor/intern/COM_NodeOperation.h b/source/blender/compositor/intern/COM_NodeOperation.h index 30731572712..f96b994685a 100644 --- a/source/blender/compositor/intern/COM_NodeOperation.h +++ b/source/blender/compositor/intern/COM_NodeOperation.h @@ -22,9 +22,7 @@ #ifndef _COM_NodeOperation_h #define _COM_NodeOperation_h - -class NodeOperation; - +class OpenCLDevice; #include "COM_Node.h" #include <string> #include <sstream> @@ -150,7 +148,7 @@ public: * @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, + virtual void executeOpenCLRegion(OpenCLDevice* device, rcti *rect, unsigned int chunkNumber, MemoryBuffer **memoryBuffers, MemoryBuffer *outputBuffer) {} /** @@ -165,7 +163,7 @@ public: * @param clMemToCleanUp all created cl_mem references must be added to this list. Framework will clean this after execution * @param clKernelsToCleanUp all created cl_kernel references must be added to this list. Framework will clean this after execution */ - 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 executeOpenCL(OpenCLDevice* device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, list<cl_kernel> *clKernelsToCleanUp) {} virtual void deinitExecution(); bool isResolutionSet() { @@ -272,15 +270,6 @@ protected: * @brief set if this NodeOperation can be scheduled on a OpenCLDevice */ 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); - 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); - }; #endif diff --git a/source/blender/compositor/intern/COM_OpenCLDevice.cpp b/source/blender/compositor/intern/COM_OpenCLDevice.cpp index 9d005804098..c9d27b8543c 100644 --- a/source/blender/compositor/intern/COM_OpenCLDevice.cpp +++ b/source/blender/compositor/intern/COM_OpenCLDevice.cpp @@ -23,13 +23,15 @@ #include "COM_OpenCLDevice.h" #include "COM_WorkScheduler.h" +typedef enum COM_VendorID {NVIDIA=0x10DE, AMD=0x1002} COM_VendorID; -OpenCLDevice::OpenCLDevice(cl_context context, cl_device_id device, cl_program program) +OpenCLDevice::OpenCLDevice(cl_context context, cl_device_id device, cl_program program, cl_int vendorId) { this->device = device; this->context = context; this->program = program; this->queue = NULL; + this->vendorID = vendorId; } bool OpenCLDevice::initialize() @@ -56,10 +58,126 @@ void OpenCLDevice::execute(WorkPackage *work) MemoryBuffer **inputBuffers = executionGroup->getInputBuffersOpenCL(chunkNumber); MemoryBuffer *outputBuffer = executionGroup->allocateOutputBuffer(chunkNumber, &rect); - executionGroup->getOutputNodeOperation()->executeOpenCLRegion(this->context, this->program, this->queue, &rect, + executionGroup->getOutputNodeOperation()->executeOpenCLRegion(this, &rect, chunkNumber, inputBuffers, outputBuffer); delete outputBuffer; executionGroup->finalizeChunkExecution(chunkNumber, inputBuffers); } + +cl_mem OpenCLDevice::COM_clAttachMemoryBufferToKernelParameter(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(this->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 OpenCLDevice::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 OpenCLDevice::COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex, NodeOperation* operation) +{ + if (offsetIndex != -1) { + cl_int error; + cl_int2 offset = {operation->getWidth(), operation->getHeight()}; + + error = clSetKernelArg(kernel, offsetIndex, sizeof(cl_int2), &offset); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + } +} + +void OpenCLDevice::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 OpenCLDevice::COM_clEnqueueRange(cl_kernel kernel, MemoryBuffer *outputMemoryBuffer) +{ + cl_int error; + const size_t size[] = {outputMemoryBuffer->getWidth(), outputMemoryBuffer->getHeight()}; + + error = clEnqueueNDRangeKernel(this->queue, kernel, 2, NULL, size, 0, 0, 0, NULL); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } +} + +void OpenCLDevice::COM_clEnqueueRange(cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex, NodeOperation* operation) +{ + cl_int error; + const int width = outputMemoryBuffer->getWidth(); + const int height = outputMemoryBuffer->getHeight(); + int offsetx; + int offsety; + int localSize = 1024; + size_t size[2]; + cl_int2 offset; + + if (this->vendorID == NVIDIA){localSize = 32;} + 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 && (!breaked); 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(this->queue, kernel, 2, NULL, size, 0, 0, 0, NULL); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + clFlush(this->queue); + if (operation->isBreaked()) { + breaked = false; + } + } + } +} + +cl_kernel OpenCLDevice::COM_clCreateKernel(const char *kernelname, list<cl_kernel> *clKernelsToCleanUp) +{ + cl_int error; + cl_kernel kernel = clCreateKernel(this->program, kernelname, &error); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + else { + if (clKernelsToCleanUp) clKernelsToCleanUp->push_back(kernel); + } + return kernel; + +} diff --git a/source/blender/compositor/intern/COM_OpenCLDevice.h b/source/blender/compositor/intern/COM_OpenCLDevice.h index 83ce8cec811..d132f330651 100644 --- a/source/blender/compositor/intern/COM_OpenCLDevice.h +++ b/source/blender/compositor/intern/COM_OpenCLDevice.h @@ -29,7 +29,6 @@ class OpenCLDevice; #include "OCL_opencl.h" #include "COM_WorkScheduler.h" - /** * @brief device representing an GPU OpenCL device. * an instance of this class represents a single cl_device @@ -55,13 +54,21 @@ private: * @brief opencl command queue */ cl_command_queue queue; + + /** + * @brief opencl vendor ID + */ + cl_int vendorID; + public: /** * @brief constructor with opencl device * @param context * @param device + * @param program + * @param vendorID */ - OpenCLDevice(cl_context context, cl_device_id device, cl_program program); + OpenCLDevice(cl_context context, cl_device_id device, cl_program program, cl_int vendorId); /** @@ -83,6 +90,18 @@ public: * @param work the WorkPackage to execute */ void execute(WorkPackage *work); + + cl_context getContext(){return this->context;} + + cl_command_queue getQueue(){return this->queue;} + + cl_mem COM_clAttachMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, int offsetIndex, list<cl_mem> *cleanup, MemoryBuffer **inputMemoryBuffers, SocketReader *reader); + void COM_clAttachMemoryBufferOffsetToKernelParameter(cl_kernel kernel, int offsetIndex, MemoryBuffer *memoryBuffers); + void COM_clAttachOutputMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, cl_mem clOutputMemoryBuffer); + void COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex, NodeOperation* operation); + void COM_clEnqueueRange(cl_kernel kernel, MemoryBuffer *outputMemoryBuffer); + void COM_clEnqueueRange(cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex, NodeOperation* operation); + cl_kernel COM_clCreateKernel(const char *kernelname, list<cl_kernel> *clKernelsToCleanUp); }; #endif diff --git a/source/blender/compositor/intern/COM_WorkPackage.h b/source/blender/compositor/intern/COM_WorkPackage.h index 18d83cc151c..fed87186d20 100644 --- a/source/blender/compositor/intern/COM_WorkPackage.h +++ b/source/blender/compositor/intern/COM_WorkPackage.h @@ -24,7 +24,7 @@ class WorkPackage; #ifndef _COM_WorkPackage_h_ #define _COM_WorkPackage_h_ - +class ExecutionGroup; #include "COM_ExecutionGroup.h" /** diff --git a/source/blender/compositor/intern/COM_WorkScheduler.cpp b/source/blender/compositor/intern/COM_WorkScheduler.cpp index fb7a8f8a764..619c5bea71e 100644 --- a/source/blender/compositor/intern/COM_WorkScheduler.cpp +++ b/source/blender/compositor/intern/COM_WorkScheduler.cpp @@ -21,15 +21,18 @@ */ #include <list> +#include <stdio.h> + +#include "BKE_global.h" + #include "COM_WorkScheduler.h" -#include "PIL_time.h" -#include "BLI_threads.h" #include "COM_CPUDevice.h" #include "COM_OpenCLDevice.h" -#include "OCL_opencl.h" -#include "stdio.h" #include "COM_OpenCLKernels.cl.h" -#include "BKE_global.h" +#include "OCL_opencl.h" + +#include "PIL_time.h" +#include "BLI_threads.h" #if COM_CURRENT_THREADING_MODEL == COM_TM_NOTHREAD #warning COM_CURRENT_THREADING_MODEL COM_TM_NOTHREAD is activated. Use only for debugging. @@ -214,60 +217,46 @@ void WorkScheduler::initialize() cl_platform_id *platforms = new cl_platform_id[numberOfPlatforms]; error = clGetPlatformIDs(numberOfPlatforms, platforms, 0); unsigned int indexPlatform; - cl_uint totalNumberOfDevices = 0; for (indexPlatform = 0; indexPlatform < numberOfPlatforms; indexPlatform++) { cl_platform_id platform = platforms[indexPlatform]; - cl_uint numberOfDevices; + cl_uint numberOfDevices = 0; clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, 0, &numberOfDevices); - totalNumberOfDevices += numberOfDevices; - } + if (numberOfDevices>0) { + cl_device_id *cldevices = new cl_device_id[numberOfDevices]; + clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numberOfDevices, cldevices, 0); - cl_device_id *cldevices = new cl_device_id[totalNumberOfDevices]; - unsigned int numberOfDevicesReceived = 0; - for (indexPlatform = 0; indexPlatform < numberOfPlatforms; indexPlatform++) { - cl_platform_id platform = platforms[indexPlatform]; - cl_uint numberOfDevices; - clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, 0, &numberOfDevices); - clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numberOfDevices, cldevices + numberOfDevicesReceived * sizeof (cl_device_id), 0); - numberOfDevicesReceived += numberOfDevices; - } - if (totalNumberOfDevices > 0) { - context = clCreateContext(NULL, totalNumberOfDevices, cldevices, clContextError, NULL, &error); - if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } - program = clCreateProgramWithSource(context, 1, &clkernelstoh_COM_OpenCLKernels_cl, 0, &error); - error = clBuildProgram(program, totalNumberOfDevices, cldevices, 0, 0, 0); - if (error != CL_SUCCESS) { - cl_int error2; - size_t ret_val_size = 0; - printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); - error2 = clGetProgramBuildInfo(program, cldevices[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size); - if (error2 != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } - char *build_log = new char[ret_val_size + 1]; - error2 = clGetProgramBuildInfo(program, cldevices[0], CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL); - if (error2 != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } - build_log[ret_val_size] = '\0'; - printf("%s", build_log); - delete build_log; - - } - else { - unsigned int indexDevices; - for (indexDevices = 0; indexDevices < totalNumberOfDevices; indexDevices++) { - cl_device_id device = cldevices[indexDevices]; - OpenCLDevice *clDevice = new OpenCLDevice(context, device, program); - clDevice->initialize(), - gpudevices.push_back(clDevice); - if (G.f & G_DEBUG) { - char resultString[32]; - error = clGetDeviceInfo(device, CL_DEVICE_NAME, 32, resultString, 0); - printf("OPENCL_DEVICE: %s, ", resultString); - error = clGetDeviceInfo(device, CL_DEVICE_VENDOR, 32, resultString, 0); - printf("%s\n", resultString); + context = clCreateContext(NULL, numberOfDevices, cldevices, clContextError, NULL, &error); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + program = clCreateProgramWithSource(context, 1, &clkernelstoh_COM_OpenCLKernels_cl, 0, &error); + error = clBuildProgram(program, numberOfDevices, cldevices, 0, 0, 0); + if (error != CL_SUCCESS) { + cl_int error2; + size_t ret_val_size = 0; + printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); + error2 = clGetProgramBuildInfo(program, cldevices[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size); + if (error2 != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + char *build_log = new char[ret_val_size + 1]; + error2 = clGetProgramBuildInfo(program, cldevices[0], CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL); + if (error2 != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + build_log[ret_val_size] = '\0'; + printf("%s", build_log); + delete build_log; + } + else { + unsigned int indexDevices; + for (indexDevices = 0; indexDevices < numberOfDevices; indexDevices++) { + cl_device_id device = cldevices[indexDevices]; + cl_int vendorID = 0; + cl_int error = clGetDeviceInfo(device, CL_DEVICE_VENDOR_ID, sizeof(cl_int), &vendorID, NULL); + if (error!= CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + OpenCLDevice *clDevice = new OpenCLDevice(context, device, program, vendorID); + clDevice->initialize(); + gpudevices.push_back(clDevice); } } + delete cldevices; } } - delete[] cldevices; delete[] platforms; } #endif diff --git a/source/blender/compositor/intern/COM_compositor.cpp b/source/blender/compositor/intern/COM_compositor.cpp index 2bbfd18e7c5..bec9ff95eed 100644 --- a/source/blender/compositor/intern/COM_compositor.cpp +++ b/source/blender/compositor/intern/COM_compositor.cpp @@ -25,6 +25,8 @@ extern "C" { #include "BLI_threads.h" } +#include "BKE_main.h" +#include "BKE_global.h" #include "COM_compositor.h" #include "COM_ExecutionSystem.h" @@ -32,7 +34,7 @@ extern "C" { #include "OCL_opencl.h" static ThreadMutex *compositorMutex; -void COM_execute(bNodeTree *editingtree, int rendering) +void COM_execute(RenderData *rd, bNodeTree *editingtree, int rendering) { if (compositorMutex == NULL) { /// TODO: move to blender startup phase compositorMutex = new ThreadMutex(); @@ -41,7 +43,7 @@ void COM_execute(bNodeTree *editingtree, int rendering) WorkScheduler::initialize(); ///TODO: call workscheduler.deinitialize somewhere } BLI_mutex_lock(compositorMutex); - if (editingtree->test_break && editingtree->test_break(editingtree->tbh)) { + if (editingtree->test_break(editingtree->tbh)) { // during editing multiple calls to this method can be triggered. // make sure one the last one will be doing the work. BLI_mutex_unlock(compositorMutex); @@ -49,11 +51,12 @@ void COM_execute(bNodeTree *editingtree, int rendering) } + /* set progress bar to 0% and status to init compositing*/ editingtree->progress(editingtree->prh, 0.0); /* initialize execution system */ - ExecutionSystem *system = new ExecutionSystem(editingtree, rendering); + ExecutionSystem *system = new ExecutionSystem(rd, editingtree, rendering); system->execute(); delete system; |