blender
/
blender
136
407
Fork 590
Code Issues 6.3k Pull Requests 779 Projects 19 Wiki Activity

Compositor: Remove now-unused "tiled" implementation #118819

Merged
Aras Pranckevicius merged 20 commits from aras_p/blender:com_remove_tiled into main 2024-02-28 16:59:26 +01:00
Conversation 1 Commits 20 Files Changed 296 +256 -20633
296 changed files with 256 additions and 20633 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
extern/CMakeLists.txt vendored
View File

@ -55,8 +55,7 @@ if(WITH_LZMA)
add_subdirectory(lzma)
endif()
if(WITH_CYCLES OR WITH_COMPOSITOR_CPU OR WITH_OPENSUBDIV)
add_subdirectory(clew)
if(WITH_CYCLES OR WITH_OPENSUBDIV)
aras_p marked this conversation as resolved Outdated
Sergey Sharybin commented 2024-02-28 13:58:24 +01:00
Outdated
Review
Copy Link

WITH_COMPOSITOR_CPU ca be removed from the condition, as neither CUDA nor HIP are used by compositor.

`WITH_COMPOSITOR_CPU ` ca be removed from the condition, as neither CUDA nor HIP are used by compositor.
if((WITH_CYCLES_DEVICE_CUDA OR WITH_CYCLES_DEVICE_OPTIX) AND WITH_CUDA_DYNLOAD)
add_subdirectory(cuew)
endif()

24
extern/clew/CMakeLists.txt vendored
View File

@ -1,24 +0,0 @@
# SPDX-FileCopyrightText: 2006 Blender Foundation
#
# SPDX-License-Identifier: GPL-2.0-or-later
set(INC
.
include
)
set(INC_SYS
)
set(SRC
include/clew.h
src/clew.c
)
set(LIB
)
add_definitions(-DCL_USE_DEPRECATED_OPENCL_1_1_APIS)
blender_add_lib(extern_clew "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")

5
extern/clew/README.blender vendored
View File

@ -1,5 +0,0 @@
Project: OpenCL Wrangler
URL: https://github.com/OpenCLWrangler/clew
License: Apache 2.0
Upstream version: 27a6867
Local modifications: None

2822
extern/clew/include/clew.h vendored
View File

File diff suppressed because it is too large Load Diff

407
extern/clew/src/clew.c vendored
View File

@ -1,407 +0,0 @@
//////////////////////////////////////////////////////////////////////////
// Copyright (c) 2009 Organic Vectory B.V.
// Written by George van Venrooij
//
// Distributed under the Boost Software License, Version 1.0.
// (See accompanying file license.txt)
//////////////////////////////////////////////////////////////////////////
#include "clew.h"
#ifdef _WIN32
#define WIN32_LEAN_AND_MEAN
#define VC_EXTRALEAN
#include <windows.h>
typedef HMODULE CLEW_DYNLIB_HANDLE;
#define CLEW_DYNLIB_OPEN LoadLibraryA
#define CLEW_DYNLIB_CLOSE FreeLibrary
#define CLEW_DYNLIB_IMPORT GetProcAddress
#else
#include <dlfcn.h>
typedef void* CLEW_DYNLIB_HANDLE;
#define CLEW_DYNLIB_OPEN(path) dlopen(path, RTLD_NOW | RTLD_GLOBAL)
#define CLEW_DYNLIB_CLOSE dlclose
#define CLEW_DYNLIB_IMPORT dlsym
#endif
#include <stdlib.h>
//! \brief module handle
static CLEW_DYNLIB_HANDLE module = NULL;
// Variables holding function entry points
PFNCLGETPLATFORMIDS __clewGetPlatformIDs = NULL;
PFNCLGETPLATFORMINFO __clewGetPlatformInfo = NULL;
PFNCLGETDEVICEIDS __clewGetDeviceIDs = NULL;
PFNCLGETDEVICEINFO __clewGetDeviceInfo = NULL;
PFNCLCREATESUBDEVICES __clewCreateSubDevices = NULL;
PFNCLRETAINDEVICE __clewRetainDevice = NULL;
PFNCLRELEASEDEVICE __clewReleaseDevice = NULL;
PFNCLCREATECONTEXT __clewCreateContext = NULL;
PFNCLCREATECONTEXTFROMTYPE __clewCreateContextFromType = NULL;
PFNCLRETAINCONTEXT __clewRetainContext = NULL;
PFNCLRELEASECONTEXT __clewReleaseContext = NULL;
PFNCLGETCONTEXTINFO __clewGetContextInfo = NULL;
PFNCLCREATECOMMANDQUEUE __clewCreateCommandQueue = NULL;
PFNCLRETAINCOMMANDQUEUE __clewRetainCommandQueue = NULL;
PFNCLRELEASECOMMANDQUEUE __clewReleaseCommandQueue = NULL;
PFNCLGETCOMMANDQUEUEINFO __clewGetCommandQueueInfo = NULL;
#ifdef CL_USE_DEPRECATED_OPENCL_1_0_APIS
PFNCLSETCOMMANDQUEUEPROPERTY __clewSetCommandQueueProperty = NULL;
#endif
PFNCLCREATEBUFFER __clewCreateBuffer = NULL;
PFNCLCREATESUBBUFFER __clewCreateSubBuffer = NULL;
PFNCLCREATEIMAGE __clewCreateImage = NULL;
PFNCLRETAINMEMOBJECT __clewRetainMemObject = NULL;
PFNCLRELEASEMEMOBJECT __clewReleaseMemObject = NULL;
PFNCLGETSUPPORTEDIMAGEFORMATS __clewGetSupportedImageFormats = NULL;
PFNCLGETMEMOBJECTINFO __clewGetMemObjectInfo = NULL;
PFNCLGETIMAGEINFO __clewGetImageInfo = NULL;
PFNCLSETMEMOBJECTDESTRUCTORCALLBACK __clewSetMemObjectDestructorCallback = NULL;
PFNCLCREATESAMPLER __clewCreateSampler = NULL;
PFNCLRETAINSAMPLER __clewRetainSampler = NULL;
PFNCLRELEASESAMPLER __clewReleaseSampler = NULL;
PFNCLGETSAMPLERINFO __clewGetSamplerInfo = NULL;
PFNCLCREATEPROGRAMWITHSOURCE __clewCreateProgramWithSource = NULL;
PFNCLCREATEPROGRAMWITHBINARY __clewCreateProgramWithBinary = NULL;
PFNCLCREATEPROGRAMWITHBUILTINKERNELS __clewCreateProgramWithBuiltInKernels = NULL;
PFNCLRETAINPROGRAM __clewRetainProgram = NULL;
PFNCLRELEASEPROGRAM __clewReleaseProgram = NULL;
PFNCLBUILDPROGRAM __clewBuildProgram = NULL;
PFNCLGETPROGRAMINFO __clewGetProgramInfo = NULL;
PFNCLGETPROGRAMBUILDINFO __clewGetProgramBuildInfo = NULL;
PFNCLCREATEKERNEL __clewCreateKernel = NULL;
PFNCLCREATEKERNELSINPROGRAM __clewCreateKernelsInProgram = NULL;
PFNCLRETAINKERNEL __clewRetainKernel = NULL;
PFNCLRELEASEKERNEL __clewReleaseKernel = NULL;
PFNCLSETKERNELARG __clewSetKernelArg = NULL;
PFNCLGETKERNELINFO __clewGetKernelInfo = NULL;
PFNCLGETKERNELWORKGROUPINFO __clewGetKernelWorkGroupInfo = NULL;
PFNCLWAITFOREVENTS __clewWaitForEvents = NULL;
PFNCLGETEVENTINFO __clewGetEventInfo = NULL;
PFNCLCREATEUSEREVENT __clewCreateUserEvent = NULL;
PFNCLRETAINEVENT __clewRetainEvent = NULL;
PFNCLRELEASEEVENT __clewReleaseEvent = NULL;
PFNCLSETUSEREVENTSTATUS __clewSetUserEventStatus = NULL;
PFNCLSETEVENTCALLBACK __clewSetEventCallback = NULL;
PFNCLGETEVENTPROFILINGINFO __clewGetEventProfilingInfo = NULL;
PFNCLFLUSH __clewFlush = NULL;
PFNCLFINISH __clewFinish = NULL;
PFNCLENQUEUEREADBUFFER __clewEnqueueReadBuffer = NULL;
PFNCLENQUEUEREADBUFFERRECT __clewEnqueueReadBufferRect = NULL;
PFNCLENQUEUEWRITEBUFFER __clewEnqueueWriteBuffer = NULL;
PFNCLENQUEUEWRITEBUFFERRECT __clewEnqueueWriteBufferRect = NULL;
PFNCLENQUEUECOPYBUFFER __clewEnqueueCopyBuffer = NULL;
PFNCLENQUEUEREADIMAGE __clewEnqueueReadImage = NULL;
PFNCLENQUEUEWRITEIMAGE __clewEnqueueWriteImage = NULL;
PFNCLENQUEUECOPYIMAGE __clewEnqueueCopyImage = NULL;
PFNCLENQUEUECOPYBUFFERRECT __clewEnqueueCopyBufferRect = NULL;
PFNCLENQUEUECOPYIMAGETOBUFFER __clewEnqueueCopyImageToBuffer = NULL;
PFNCLENQUEUECOPYBUFFERTOIMAGE __clewEnqueueCopyBufferToImage = NULL;
PFNCLENQUEUEMAPBUFFER __clewEnqueueMapBuffer = NULL;
PFNCLENQUEUEMAPIMAGE __clewEnqueueMapImage = NULL;
PFNCLENQUEUEUNMAPMEMOBJECT __clewEnqueueUnmapMemObject = NULL;
PFNCLENQUEUENDRANGEKERNEL __clewEnqueueNDRangeKernel = NULL;
PFNCLENQUEUETASK __clewEnqueueTask = NULL;
PFNCLENQUEUENATIVEKERNEL __clewEnqueueNativeKernel = NULL;
PFNCLGETEXTENSIONFUNCTIONADDRESSFORPLATFORM __clewGetExtensionFunctionAddressForPlatform = NULL;
#ifdef CL_USE_DEPRECATED_OPENCL_1_1_APIS
PFNCLCREATEIMAGE2D __clewCreateImage2D = NULL;
PFNCLCREATEIMAGE3D __clewCreateImage3D = NULL;
PFNCLENQUEUEMARKER __clewEnqueueMarker = NULL;
PFNCLENQUEUEWAITFOREVENTS __clewEnqueueWaitForEvents = NULL;
PFNCLENQUEUEBARRIER __clewEnqueueBarrier = NULL;
PFNCLUNLOADCOMPILER __clewUnloadCompiler = NULL;
PFNCLGETEXTENSIONFUNCTIONADDRESS __clewGetExtensionFunctionAddress = NULL;
#endif
/* cl_gl */
PFNCLCREATEFROMGLBUFFER __clewCreateFromGLBuffer = NULL;
PFNCLCREATEFROMGLTEXTURE __clewCreateFromGLTexture = NULL;
PFNCLCREATEFROMGLRENDERBUFFER __clewCreateFromGLRenderbuffer = NULL;
PFNCLGETGLOBJECTINFO __clewGetGLObjectInfo = NULL;
PFNCLGETGLTEXTUREINFO __clewGetGLTextureInfo = NULL;
PFNCLENQUEUEACQUIREGLOBJECTS __clewEnqueueAcquireGLObjects = NULL;
PFNCLENQUEUERELEASEGLOBJECTS __clewEnqueueReleaseGLObjects = NULL;
#ifdef CL_USE_DEPRECATED_OPENCL_1_1_APIS
PFNCLCREATEFROMGLTEXTURE2D __clewCreateFromGLTexture2D = NULL;
PFNCLCREATEFROMGLTEXTURE3D __clewCreateFromGLTexture3D = NULL;
#endif
PFNCLGETGLCONTEXTINFOKHR __clewGetGLContextInfoKHR = NULL;
static CLEW_DYNLIB_HANDLE dynamic_library_open_find(const char **paths) {
int i = 0;
while (paths[i] != NULL) {
CLEW_DYNLIB_HANDLE lib = CLEW_DYNLIB_OPEN(paths[i]);
if (lib != NULL) {
return lib;
}
++i;
}
return NULL;
}
static void clewExit(void)
{
if (module != NULL)
{
// Ignore errors
CLEW_DYNLIB_CLOSE(module);
module = NULL;
}
}
int clewInit()
{
#ifdef _WIN32
const char *paths[] = {"OpenCL.dll", NULL};
#elif defined(__APPLE__)
const char *paths[] = {"/Library/Frameworks/OpenCL.framework/OpenCL", NULL};
#else
const char *paths[] = {"libOpenCL.so",
"libOpenCL.so.0",
"libOpenCL.so.1",
"libOpenCL.so.2",
NULL};
#endif
int error = 0;
// Check if already initialized
if (module != NULL)
{
return CLEW_SUCCESS;
}
// Load library
module = dynamic_library_open_find(paths);
// Check for errors
if (module == NULL)
{
return CLEW_ERROR_OPEN_FAILED;
}
// Set unloading
error = atexit(clewExit);
if (error)
{
// Failure queuing atexit, shutdown with error
CLEW_DYNLIB_CLOSE(module);
module = NULL;
return CLEW_ERROR_ATEXIT_FAILED;
}
// Determine function entry-points
__clewGetPlatformIDs = (PFNCLGETPLATFORMIDS )CLEW_DYNLIB_IMPORT(module, "clGetPlatformIDs");
__clewGetPlatformInfo = (PFNCLGETPLATFORMINFO )CLEW_DYNLIB_IMPORT(module, "clGetPlatformInfo");
__clewGetDeviceIDs = (PFNCLGETDEVICEIDS )CLEW_DYNLIB_IMPORT(module, "clGetDeviceIDs");
__clewGetDeviceInfo = (PFNCLGETDEVICEINFO )CLEW_DYNLIB_IMPORT(module, "clGetDeviceInfo");
__clewCreateSubDevices = (PFNCLCREATESUBDEVICES )CLEW_DYNLIB_IMPORT(module, "clCreateSubDevices");
__clewRetainDevice = (PFNCLRETAINDEVICE )CLEW_DYNLIB_IMPORT(module, "clRetainDevice");
__clewReleaseDevice = (PFNCLRELEASEDEVICE )CLEW_DYNLIB_IMPORT(module, "clReleaseDevice");
__clewCreateContext = (PFNCLCREATECONTEXT )CLEW_DYNLIB_IMPORT(module, "clCreateContext");
__clewCreateContextFromType = (PFNCLCREATECONTEXTFROMTYPE )CLEW_DYNLIB_IMPORT(module, "clCreateContextFromType");
__clewRetainContext = (PFNCLRETAINCONTEXT )CLEW_DYNLIB_IMPORT(module, "clRetainContext");
__clewReleaseContext = (PFNCLRELEASECONTEXT )CLEW_DYNLIB_IMPORT(module, "clReleaseContext");
__clewGetContextInfo = (PFNCLGETCONTEXTINFO )CLEW_DYNLIB_IMPORT(module, "clGetContextInfo");
__clewCreateCommandQueue = (PFNCLCREATECOMMANDQUEUE )CLEW_DYNLIB_IMPORT(module, "clCreateCommandQueue");
__clewRetainCommandQueue = (PFNCLRETAINCOMMANDQUEUE )CLEW_DYNLIB_IMPORT(module, "clRetainCommandQueue");
__clewReleaseCommandQueue = (PFNCLRELEASECOMMANDQUEUE )CLEW_DYNLIB_IMPORT(module, "clReleaseCommandQueue");
__clewGetCommandQueueInfo = (PFNCLGETCOMMANDQUEUEINFO )CLEW_DYNLIB_IMPORT(module, "clGetCommandQueueInfo");
#ifdef CL_USE_DEPRECATED_OPENCL_1_0_APIS
__clewSetCommandQueueProperty = (PFNCLSETCOMMANDQUEUEPROPERTY )CLEW_DYNLIB_IMPORT(module, "clSetCommandQueueProperty");
#endif
__clewCreateBuffer = (PFNCLCREATEBUFFER )CLEW_DYNLIB_IMPORT(module, "clCreateBuffer");
__clewCreateSubBuffer = (PFNCLCREATESUBBUFFER )CLEW_DYNLIB_IMPORT(module, "clCreateSubBuffer");
__clewCreateImage = (PFNCLCREATEIMAGE )CLEW_DYNLIB_IMPORT(module, "clCreateImage");
__clewRetainMemObject = (PFNCLRETAINMEMOBJECT )CLEW_DYNLIB_IMPORT(module, "clRetainMemObject");
__clewReleaseMemObject = (PFNCLRELEASEMEMOBJECT )CLEW_DYNLIB_IMPORT(module, "clReleaseMemObject");
__clewGetSupportedImageFormats = (PFNCLGETSUPPORTEDIMAGEFORMATS )CLEW_DYNLIB_IMPORT(module, "clGetSupportedImageFormats");
__clewGetMemObjectInfo = (PFNCLGETMEMOBJECTINFO )CLEW_DYNLIB_IMPORT(module, "clGetMemObjectInfo");
__clewGetImageInfo = (PFNCLGETIMAGEINFO )CLEW_DYNLIB_IMPORT(module, "clGetImageInfo");
__clewSetMemObjectDestructorCallback = (PFNCLSETMEMOBJECTDESTRUCTORCALLBACK)CLEW_DYNLIB_IMPORT(module, "clSetMemObjectDestructorCallback");
__clewCreateSampler = (PFNCLCREATESAMPLER )CLEW_DYNLIB_IMPORT(module, "clCreateSampler");
__clewRetainSampler = (PFNCLRETAINSAMPLER )CLEW_DYNLIB_IMPORT(module, "clRetainSampler");
__clewReleaseSampler = (PFNCLRELEASESAMPLER )CLEW_DYNLIB_IMPORT(module, "clReleaseSampler");
__clewGetSamplerInfo = (PFNCLGETSAMPLERINFO )CLEW_DYNLIB_IMPORT(module, "clGetSamplerInfo");
__clewCreateProgramWithSource = (PFNCLCREATEPROGRAMWITHSOURCE )CLEW_DYNLIB_IMPORT(module, "clCreateProgramWithSource");
__clewCreateProgramWithBinary = (PFNCLCREATEPROGRAMWITHBINARY )CLEW_DYNLIB_IMPORT(module, "clCreateProgramWithBinary");
__clewCreateProgramWithBuiltInKernels =(PFNCLCREATEPROGRAMWITHBUILTINKERNELS)CLEW_DYNLIB_IMPORT(module, "clCreateProgramWithBuiltInKernels");
__clewRetainProgram = (PFNCLRETAINPROGRAM )CLEW_DYNLIB_IMPORT(module, "clRetainProgram");
__clewReleaseProgram = (PFNCLRELEASEPROGRAM )CLEW_DYNLIB_IMPORT(module, "clReleaseProgram");
__clewBuildProgram = (PFNCLBUILDPROGRAM )CLEW_DYNLIB_IMPORT(module, "clBuildProgram");
__clewGetProgramInfo = (PFNCLGETPROGRAMINFO )CLEW_DYNLIB_IMPORT(module, "clGetProgramInfo");
__clewGetProgramBuildInfo = (PFNCLGETPROGRAMBUILDINFO )CLEW_DYNLIB_IMPORT(module, "clGetProgramBuildInfo");
__clewCreateKernel = (PFNCLCREATEKERNEL )CLEW_DYNLIB_IMPORT(module, "clCreateKernel");
__clewCreateKernelsInProgram = (PFNCLCREATEKERNELSINPROGRAM )CLEW_DYNLIB_IMPORT(module, "clCreateKernelsInProgram");
__clewRetainKernel = (PFNCLRETAINKERNEL )CLEW_DYNLIB_IMPORT(module, "clRetainKernel");
__clewReleaseKernel = (PFNCLRELEASEKERNEL )CLEW_DYNLIB_IMPORT(module, "clReleaseKernel");
__clewSetKernelArg = (PFNCLSETKERNELARG )CLEW_DYNLIB_IMPORT(module, "clSetKernelArg");
__clewGetKernelInfo = (PFNCLGETKERNELINFO )CLEW_DYNLIB_IMPORT(module, "clGetKernelInfo");
__clewGetKernelWorkGroupInfo = (PFNCLGETKERNELWORKGROUPINFO )CLEW_DYNLIB_IMPORT(module, "clGetKernelWorkGroupInfo");
__clewWaitForEvents = (PFNCLWAITFOREVENTS )CLEW_DYNLIB_IMPORT(module, "clWaitForEvents");
__clewGetEventInfo = (PFNCLGETEVENTINFO )CLEW_DYNLIB_IMPORT(module, "clGetEventInfo");
__clewCreateUserEvent = (PFNCLCREATEUSEREVENT )CLEW_DYNLIB_IMPORT(module, "clCreateUserEvent");
__clewRetainEvent = (PFNCLRETAINEVENT )CLEW_DYNLIB_IMPORT(module, "clRetainEvent");
__clewReleaseEvent = (PFNCLRELEASEEVENT )CLEW_DYNLIB_IMPORT(module, "clReleaseEvent");
__clewSetUserEventStatus = (PFNCLSETUSEREVENTSTATUS )CLEW_DYNLIB_IMPORT(module, "clSetUserEventStatus");
__clewSetEventCallback = (PFNCLSETEVENTCALLBACK )CLEW_DYNLIB_IMPORT(module, "clSetEventCallback");
__clewGetEventProfilingInfo = (PFNCLGETEVENTPROFILINGINFO )CLEW_DYNLIB_IMPORT(module, "clGetEventProfilingInfo");
__clewFlush = (PFNCLFLUSH )CLEW_DYNLIB_IMPORT(module, "clFlush");
__clewFinish = (PFNCLFINISH )CLEW_DYNLIB_IMPORT(module, "clFinish");
__clewEnqueueReadBuffer = (PFNCLENQUEUEREADBUFFER )CLEW_DYNLIB_IMPORT(module, "clEnqueueReadBuffer");
__clewEnqueueReadBufferRect = (PFNCLENQUEUEREADBUFFERRECT )CLEW_DYNLIB_IMPORT(module, "clEnqueueReadBufferRect");
__clewEnqueueWriteBuffer = (PFNCLENQUEUEWRITEBUFFER )CLEW_DYNLIB_IMPORT(module, "clEnqueueWriteBuffer");
__clewEnqueueWriteBufferRect = (PFNCLENQUEUEWRITEBUFFERRECT )CLEW_DYNLIB_IMPORT(module, "clEnqueueWriteBufferRect");
__clewEnqueueCopyBuffer = (PFNCLENQUEUECOPYBUFFER )CLEW_DYNLIB_IMPORT(module, "clEnqueueCopyBuffer");
__clewEnqueueCopyBufferRect = (PFNCLENQUEUECOPYBUFFERRECT )CLEW_DYNLIB_IMPORT(module, "clEnqueueCopyBufferRect");
__clewEnqueueReadImage = (PFNCLENQUEUEREADIMAGE )CLEW_DYNLIB_IMPORT(module, "clEnqueueReadImage");
__clewEnqueueWriteImage = (PFNCLENQUEUEWRITEIMAGE )CLEW_DYNLIB_IMPORT(module, "clEnqueueWriteImage");
__clewEnqueueCopyImage = (PFNCLENQUEUECOPYIMAGE )CLEW_DYNLIB_IMPORT(module, "clEnqueueCopyImage");
__clewEnqueueCopyImageToBuffer = (PFNCLENQUEUECOPYIMAGETOBUFFER )CLEW_DYNLIB_IMPORT(module, "clEnqueueCopyImageToBuffer");
__clewEnqueueCopyBufferToImage = (PFNCLENQUEUECOPYBUFFERTOIMAGE )CLEW_DYNLIB_IMPORT(module, "clEnqueueCopyBufferToImage");
__clewEnqueueMapBuffer = (PFNCLENQUEUEMAPBUFFER )CLEW_DYNLIB_IMPORT(module, "clEnqueueMapBuffer");
__clewEnqueueMapImage = (PFNCLENQUEUEMAPIMAGE )CLEW_DYNLIB_IMPORT(module, "clEnqueueMapImage");
__clewEnqueueUnmapMemObject = (PFNCLENQUEUEUNMAPMEMOBJECT )CLEW_DYNLIB_IMPORT(module, "clEnqueueUnmapMemObject");
__clewEnqueueNDRangeKernel = (PFNCLENQUEUENDRANGEKERNEL )CLEW_DYNLIB_IMPORT(module, "clEnqueueNDRangeKernel");
__clewEnqueueTask = (PFNCLENQUEUETASK )CLEW_DYNLIB_IMPORT(module, "clEnqueueTask");
__clewEnqueueNativeKernel = (PFNCLENQUEUENATIVEKERNEL )CLEW_DYNLIB_IMPORT(module, "clEnqueueNativeKernel");
__clewGetExtensionFunctionAddressForPlatform = (PFNCLGETEXTENSIONFUNCTIONADDRESSFORPLATFORM)CLEW_DYNLIB_IMPORT(module, "clGetExtensionFunctionAddressForPlatform");
#ifdef CL_USE_DEPRECATED_OPENCL_1_1_APIS
__clewCreateImage2D = (PFNCLCREATEIMAGE2D )CLEW_DYNLIB_IMPORT(module, "clCreateImage2D");
__clewCreateImage3D = (PFNCLCREATEIMAGE3D )CLEW_DYNLIB_IMPORT(module, "clCreateImage3D");
__clewEnqueueMarker = (PFNCLENQUEUEMARKER )CLEW_DYNLIB_IMPORT(module, "clEnqueueMarker");
__clewEnqueueWaitForEvents = (PFNCLENQUEUEWAITFOREVENTS )CLEW_DYNLIB_IMPORT(module, "clEnqueueWaitForEvents");
__clewEnqueueBarrier = (PFNCLENQUEUEBARRIER )CLEW_DYNLIB_IMPORT(module, "clEnqueueBarrier");
__clewUnloadCompiler = (PFNCLUNLOADCOMPILER )CLEW_DYNLIB_IMPORT(module, "clUnloadCompiler");
__clewGetExtensionFunctionAddress = (PFNCLGETEXTENSIONFUNCTIONADDRESS )CLEW_DYNLIB_IMPORT(module, "clGetExtensionFunctionAddress");
#endif
/* cl_gl */
__clewCreateFromGLBuffer = (PFNCLCREATEFROMGLBUFFER )CLEW_DYNLIB_IMPORT(module, "clCreateFromGLBuffer");
__clewCreateFromGLTexture = (PFNCLCREATEFROMGLTEXTURE )CLEW_DYNLIB_IMPORT(module, "clCreateFromGLTexture");
__clewCreateFromGLRenderbuffer = (PFNCLCREATEFROMGLRENDERBUFFER )CLEW_DYNLIB_IMPORT(module, "clCreateFromGLRenderbuffer");
__clewGetGLObjectInfo = (PFNCLGETGLOBJECTINFO )CLEW_DYNLIB_IMPORT(module, "clGetGLObjectInfo");
__clewGetGLTextureInfo = (PFNCLGETGLTEXTUREINFO )CLEW_DYNLIB_IMPORT(module, "clGetGLTextureInfo");
__clewEnqueueAcquireGLObjects = (PFNCLENQUEUEACQUIREGLOBJECTS )CLEW_DYNLIB_IMPORT(module, "clEnqueueAcquireGLObjects");
__clewEnqueueReleaseGLObjects = (PFNCLENQUEUERELEASEGLOBJECTS )CLEW_DYNLIB_IMPORT(module, "clEnqueueReleaseGLObjects");
#ifdef CL_USE_DEPRECATED_OPENCL_1_1_APIS
__clewCreateFromGLTexture2D = (PFNCLCREATEFROMGLTEXTURE2D )CLEW_DYNLIB_IMPORT(module, "clCreateFromGLTexture2D");
__clewCreateFromGLTexture3D = (PFNCLCREATEFROMGLTEXTURE3D )CLEW_DYNLIB_IMPORT(module, "clCreateFromGLTexture3D");
#endif
__clewGetGLContextInfoKHR = (PFNCLGETGLCONTEXTINFOKHR )CLEW_DYNLIB_IMPORT(module, "clGetGLContextInfoKHR");
if(__clewGetPlatformIDs == NULL) return 0;
if(__clewGetPlatformInfo == NULL) return 0;
if(__clewGetDeviceIDs == NULL) return 0;
if(__clewGetDeviceInfo == NULL) return 0;
return CLEW_SUCCESS;
}
const char* clewErrorString(cl_int error)
{
static const char* strings[] =
{
// Error Codes
"CL_SUCCESS" // 0
, "CL_DEVICE_NOT_FOUND" // -1
, "CL_DEVICE_NOT_AVAILABLE" // -2
, "CL_COMPILER_NOT_AVAILABLE" // -3
, "CL_MEM_OBJECT_ALLOCATION_FAILURE" // -4
, "CL_OUT_OF_RESOURCES" // -5
, "CL_OUT_OF_HOST_MEMORY" // -6
, "CL_PROFILING_INFO_NOT_AVAILABLE" // -7
, "CL_MEM_COPY_OVERLAP" // -8
, "CL_IMAGE_FORMAT_MISMATCH" // -9
, "CL_IMAGE_FORMAT_NOT_SUPPORTED" // -10
, "CL_BUILD_PROGRAM_FAILURE" // -11
, "CL_MAP_FAILURE" // -12
, "CL_MISALIGNED_SUB_BUFFER_OFFSET" // -13
, "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"// -14
, "CL_COMPILE_PROGRAM_FAILURE" // -15
, "CL_LINKER_NOT_AVAILABLE" // -16
, "CL_LINK_PROGRAM_FAILURE" // -17
, "CL_DEVICE_PARTITION_FAILED" // -18
, "CL_KERNEL_ARG_INFO_NOT_AVAILABLE" // -19
, "" // -20
, "" // -21
, "" // -22
, "" // -23
, "" // -24
, "" // -25
, "" // -26
, "" // -27
, "" // -28
, "" // -29
, "CL_INVALID_VALUE" // -30
, "CL_INVALID_DEVICE_TYPE" // -31
, "CL_INVALID_PLATFORM" // -32
, "CL_INVALID_DEVICE" // -33
, "CL_INVALID_CONTEXT" // -34
, "CL_INVALID_QUEUE_PROPERTIES" // -35
, "CL_INVALID_COMMAND_QUEUE" // -36
, "CL_INVALID_HOST_PTR" // -37
, "CL_INVALID_MEM_OBJECT" // -38
, "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR" // -39
, "CL_INVALID_IMAGE_SIZE" // -40
, "CL_INVALID_SAMPLER" // -41
, "CL_INVALID_BINARY" // -42
, "CL_INVALID_BUILD_OPTIONS" // -43
, "CL_INVALID_PROGRAM" // -44
, "CL_INVALID_PROGRAM_EXECUTABLE" // -45
, "CL_INVALID_KERNEL_NAME" // -46
, "CL_INVALID_KERNEL_DEFINITION" // -47
, "CL_INVALID_KERNEL" // -48
, "CL_INVALID_ARG_INDEX" // -49
, "CL_INVALID_ARG_VALUE" // -50
, "CL_INVALID_ARG_SIZE" // -51
, "CL_INVALID_KERNEL_ARGS" // -52
, "CL_INVALID_WORK_DIMENSION" // -53
, "CL_INVALID_WORK_GROUP_SIZE" // -54
, "CL_INVALID_WORK_ITEM_SIZE" // -55
, "CL_INVALID_GLOBAL_OFFSET" // -56
, "CL_INVALID_EVENT_WAIT_LIST" // -57
, "CL_INVALID_EVENT" // -58
, "CL_INVALID_OPERATION" // -59
, "CL_INVALID_GL_OBJECT" // -60
, "CL_INVALID_BUFFER_SIZE" // -61
, "CL_INVALID_MIP_LEVEL" // -62
, "CL_INVALID_GLOBAL_WORK_SIZE" // -63
, "CL_INVALID_PROPERTY" // -64
, "CL_INVALID_IMAGE_DESCRIPTOR" // -65
, "CL_INVALID_COMPILER_OPTIONS" // -66
, "CL_INVALID_LINKER_OPTIONS" // -67
, "CL_INVALID_DEVICE_PARTITION_COUNT" // -68
};
static const int num_errors = sizeof(strings) / sizeof(strings[0]);
if (error == -1001) {
return "CL_PLATFORM_NOT_FOUND_KHR";
}
if (error > 0 || -error >= num_errors) {
return "Unknown OpenCL error";
}
return strings[-error];
}

1
release/license/THIRD-PARTY-LICENSES.txt
View File

@ -32,7 +32,6 @@ https://github.com/AcademySoftwareFoundation/MaterialX
** meson; version 0.63 -- https://github.com/mesonbuild/meson
** oneAPI Threading Building Block; version 2020_U3 --
https://software.intel.com/en-us/oneapi/onetbb
** OpenCL Wrangler; version 27a6867 -- https://github.com/OpenCLWrangler/clew
** OpenImageDenoise; version 1.4.3 -- https://www.openimagedenoise.org/
** OpenImageIO; version 2.4.15.0 -- http://www.openimageio.org
** OpenSSL; version 3.1.2 -- https://www.openssl.org/

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

@ -22,7 +22,6 @@ if(WITH_COMPOSITOR_CPU)
../nodes/intern
../render
../render/intern
../../../extern/clew/include
# RNA_prototypes.h
${CMAKE_BINARY_DIR}/source/blender/makesrna
@ -38,16 +37,10 @@ if(WITH_COMPOSITOR_CPU)
COM_profile.hh
intern/COM_BufferArea.h
intern/COM_BufferOperation.cc
intern/COM_BufferOperation.h
intern/COM_BufferRange.h
intern/COM_BuffersIterator.h
intern/COM_CPUDevice.cc
intern/COM_CPUDevice.h
intern/COM_ChunkOrder.cc
intern/COM_ChunkOrder.h
intern/COM_ChunkOrderHotspot.cc
intern/COM_ChunkOrderHotspot.h
intern/COM_CompositorContext.cc
intern/COM_CompositorContext.h
intern/COM_ConstantFolder.cc
@ -60,8 +53,6 @@ if(WITH_COMPOSITOR_CPU)
intern/COM_Device.h
intern/COM_Enums.cc
intern/COM_Enums.h
intern/COM_ExecutionGroup.cc
intern/COM_ExecutionGroup.h
intern/COM_ExecutionModel.cc
intern/COM_ExecutionModel.h
intern/COM_ExecutionSystem.cc
@ -70,8 +61,6 @@ if(WITH_COMPOSITOR_CPU)
intern/COM_FullFrameExecutionModel.h
intern/COM_MemoryBuffer.cc
intern/COM_MemoryBuffer.h
intern/COM_MemoryProxy.cc
intern/COM_MemoryProxy.h
intern/COM_MetaData.cc
intern/COM_MetaData.h
intern/COM_MultiThreadedOperation.cc
@ -88,15 +77,8 @@ if(WITH_COMPOSITOR_CPU)
intern/COM_NodeOperation.h
intern/COM_NodeOperationBuilder.cc
intern/COM_NodeOperationBuilder.h
intern/COM_OpenCLDevice.cc
intern/COM_OpenCLDevice.h
intern/COM_SharedOperationBuffers.cc
intern/COM_SharedOperationBuffers.h
intern/COM_SingleThreadedOperation.cc
intern/COM_SingleThreadedOperation.h
intern/COM_TiledExecutionModel.cc
intern/COM_TiledExecutionModel.h
intern/COM_WorkPackage.cc
intern/COM_WorkPackage.h
intern/COM_WorkScheduler.cc
intern/COM_WorkScheduler.h
@ -335,18 +317,10 @@ if(WITH_COMPOSITOR_CPU)
operations/COM_GammaCorrectOperation.h
operations/COM_GaussianAlphaBlurBaseOperation.cc
operations/COM_GaussianAlphaBlurBaseOperation.h
operations/COM_GaussianAlphaXBlurOperation.cc
operations/COM_GaussianAlphaXBlurOperation.h
operations/COM_GaussianAlphaYBlurOperation.cc
operations/COM_GaussianAlphaYBlurOperation.h
operations/COM_GaussianBlurBaseOperation.cc
operations/COM_GaussianBlurBaseOperation.h
operations/COM_GaussianBokehBlurOperation.cc
operations/COM_GaussianBokehBlurOperation.h
operations/COM_GaussianXBlurOperation.cc
operations/COM_GaussianXBlurOperation.h
operations/COM_GaussianYBlurOperation.cc
operations/COM_GaussianYBlurOperation.h
operations/COM_KuwaharaAnisotropicOperation.cc
operations/COM_KuwaharaAnisotropicOperation.h
operations/COM_KuwaharaAnisotropicStructureTensorOperation.cc
@ -473,16 +447,12 @@ if(WITH_COMPOSITOR_CPU)
operations/COM_GammaOperation.h
operations/COM_MixOperation.cc
operations/COM_MixOperation.h
operations/COM_ReadBufferOperation.cc
operations/COM_ReadBufferOperation.h
operations/COM_SetColorOperation.cc
operations/COM_SetColorOperation.h
operations/COM_SetValueOperation.cc
operations/COM_SetValueOperation.h
operations/COM_SetVectorOperation.cc
operations/COM_SetVectorOperation.h
operations/COM_WriteBufferOperation.cc
operations/COM_WriteBufferOperation.h
operations/COM_MathBaseOperation.cc
operations/COM_MathBaseOperation.h
@ -536,8 +506,6 @@ if(WITH_COMPOSITOR_CPU)
operations/COM_TransformOperation.h
operations/COM_TranslateOperation.cc
operations/COM_TranslateOperation.h
operations/COM_WrapOperation.cc
operations/COM_WrapOperation.h
# Filter operations
operations/COM_ConvolutionEdgeFilterOperation.cc
@ -607,7 +575,6 @@ if(WITH_COMPOSITOR_CPU)
PRIVATE bf::intern::clog
PRIVATE bf::intern::guardedalloc
bf_realtime_compositor
extern_clew
PRIVATE bf::intern::atomic
)
@ -615,15 +582,6 @@ if(WITH_COMPOSITOR_CPU)
${CMAKE_CURRENT_BINARY_DIR}/operations
)
data_to_c(
${CMAKE_CURRENT_SOURCE_DIR}/operations/COM_OpenCLKernels.cl
${CMAKE_CURRENT_BINARY_DIR}/operations/COM_OpenCLKernels.cl.h
SRC
STRIP_LEADING_C_COMMENTS
)
add_definitions(-DCL_USE_DEPRECATED_OPENCL_1_1_APIS)
set(GENSRC_DIR ${CMAKE_CURRENT_BINARY_DIR}/operations)
set(GENSRC ${GENSRC_DIR}/COM_SMAAAreaTexture.h)
add_custom_command(

230
source/blender/compositor/COM_compositor.hh
View File

@ -33,49 +33,6 @@ struct Render;
* \defgroup Operation All operations of the compositor
* \ingroup compositor
*
* \page Introduction of the Blender Compositor
*
* \section bcomp Blender compositor
* This project redesigns the internals of Blender's compositor.
* The project has been executed in 2011 by At Mind.
* At Mind is a technology company located in Amsterdam, The Netherlands.
* The project has been crowd-funded. This code has been released under GPL2 to be used in Blender.
*
* \section goals The goals of the project
* the new compositor has 2 goals.
* - Make a faster compositor (speed of calculation)
* - Make the compositor work faster for you (workflow)
*
* \section speed Faster compositor
* The speedup has been done by making better use of the hardware Blenders is working on.
* The previous compositor only used a single threaded model to calculate a node.
* The only exception to this is the Defocus node.
* Only when it is possible to calculate two full nodes in parallel a second thread was used.
* Current workstations have 8-16 threads available, and most of the time these are idle.
*
* In the new compositor we want to use as much of threads as possible.
* Even new OpenCL capable GPU-hardware can be used for calculation.
*
* \section workflow Work faster
* The previous compositor only showed the final image.
* The compositor could wait a long time before seeing the result of his work.
* The new compositor will work in a way that it will focus on
* getting information back to the user. It will prioritize its work to get earlier user feedback.
*
* \page memory Memory model
* The main issue is the type of memory model to use.
* Blender is used by consumers and professionals.
* Ranging from low-end machines to very high-end machines.
* The system should work on high-end machines and on low-end machines.
* \page executing Executing
* \section prepare Prepare execution
*
* during the preparation of the execution All ReadBufferOperation will receive an offset.
* This offset is used during execution as an optimization trick
* Next all operations will be initialized for execution \see NodeOperation.init_execution
* Next all ExecutionGroup's will be initialized for execution \see ExecutionGroup.init_execution
* this all is controlled from \see ExecutionSystem.execute
*
* \section priority Render priority
* Render priority is an priority of an output node.
* A user has a different need of Render priorities of output nodes
@ -85,152 +42,6 @@ struct Render;
* All NodeOperation has a setting for their render-priority,
* but only for output NodeOperation these have effect.
* In ExecutionSystem.execute all priorities are checked.
* For every priority the ExecutionGroup's are check if the
* priority do match.
* When match the ExecutionGroup will be executed (this happens in serial)
*
* \see ExecutionSystem.execute control of the Render priority
* \see NodeOperation.get_render_priority receive the render priority
* \see ExecutionGroup.execute the main loop to execute a whole ExecutionGroup
*
* \section order Chunk order
*
* When a ExecutionGroup is executed, first the order of chunks are determined.
* The settings are stored in the ViewerNode inside the ExecutionGroup.
* ExecutionGroups that have no viewer-node,
* will use a default one.
* There are several possible chunk orders
* - [@ref ChunkOrdering.CenterOut]:
* Start calculating from a configurable point and order by nearest chunk.
* - [@ref ChunkOrdering.Random]:
* Randomize all chunks.
* - [@ref ChunkOrdering.TopDown]:
* Start calculation from the bottom to the top of the image.
* - [@ref ChunkOrdering.RuleOfThirds]:
* Experimental order based on 9 hot-spots in the image.
*
* When the chunk-order is determined, the first few chunks will be checked if they can be scheduled.
* Chunks can have three states:
* - [@ref eWorkPackageState.NotScheduled]:
* Chunk is not yet scheduled, or dependencies are not met.
* - [@ref eWorkPackageState.Scheduled]:
* All dependencies are met, chunk is scheduled, but not finished.
* - [@ref eWorkPackageState.Executed]:
* Chunk is finished.
*
* \see ExecutionGroup.execute
* \see ViewerOperation.get_chunk_order
* \see ChunkOrdering
*
* \section interest Area of interest
* An ExecutionGroup can have dependencies to other ExecutionGroup's.
* Data passing from one ExecutionGroup to another one are stored in 'chunks'.
* If not all input chunks are available the chunk execution will not be scheduled.
* <pre>
* +-------------------------------------+ +--------------------------------------+
* | ExecutionGroup A | | ExecutionGroup B |
* | +----------------+ +-------------+ | | +------------+ +-----------------+ |
* | | NodeOperation a| | WriteBuffer | | | | ReadBuffer | | ViewerOperation | |
* | | *==* Operation | | | | Operation *===* | |
* | | | | | | | | | | | |
* | +----------------+ +-------------+ | | +------------+ +-----------------+ |
* | | | | | |
* +--------------------------------|----+ +---|----------------------------------+
* | |
* | |
* +---------------------------+
* | MemoryProxy |
* | +----------+ +---------+ |
* | | Chunk a | | Chunk b | |
* | | | | | |
* | +----------+ +---------+ |
* | |
* +---------------------------+
* </pre>
*
* In the above example ExecutionGroup B has an outputoperation (ViewerOperation)
* and is being executed.
* The first chunk is evaluated [@ref ExecutionGroup.schedule_chunk_when_possible],
* but not all input chunks are available.
* The relevant ExecutionGroup (that can calculate the missing chunks; ExecutionGroup A)
* is asked to calculate the area ExecutionGroup B is missing.
* [@ref ExecutionGroup.schedule_area_when_possible]
* ExecutionGroup B checks what chunks the area spans, and tries to schedule these chunks.
* If all input data is available these chunks are scheduled [@ref ExecutionGroup.schedule_chunk]
*
* <pre>
*
* +-------------------------+ +----------------+ +----------------+
* | ExecutionSystem.execute | | ExecutionGroup | | ExecutionGroup |
* +-------------------------+ | (B) | | (A) |
* O +----------------+ +----------------+
* O | |
* O ExecutionGroup.execute | |
* O------------------------------->O |
* . O |
* . O-------\ |
* . . | ExecutionGroup.schedule_chunk_when_possible
* . . O----/ (*) |
* . . O |
* . . O |
* . . O ExecutionGroup.schedule_area_when_possible|
* . . O---------------------------------------->O
* . . . O----------\ ExecutionGroup.schedule_chunk_when_possible
* . . . . | (*)
* . . . . O-------/
* . . . . O
* . . . . O
* . . . . O-------\ ExecutionGroup.schedule_chunk
* . . . . . |
* . . . . . O----/
* . . . . O<=O
* . . . O<=O
* . . . O
* . . O<========================================O
* . . O |
* . O<=O |
* . O |
* . O |
* </pre>
*
* This happens until all chunks of (ExecutionGroup B) are finished executing or the user break's the process.
*
* NodeOperation like the ScaleOperation can influence the area of interest by reimplementing the
* [@ref NodeOperation.determine_area_of_interest] method
*
* <pre>
*
* +--------------------------+ +---------------------------------+
* | ExecutionGroup A | | ExecutionGroup B |
* | | | |
* +--------------------------+ +---------------------------------+
* Needed chunks from ExecutionGroup A | Chunk of ExecutionGroup B (to be evaluated)
* +-------+ +-------+ | +--------+
* |Chunk 1| |Chunk 2| +----------------+ |Chunk 1 |
* | | | | | ScaleOperation | | |
* +-------+ +-------+ +----------------+ +--------+
*
* +-------+ +-------+
* |Chunk 3| |Chunk 4|
* | | | |
* +-------+ +-------+
*
* </pre>
*
* \see ExecutionGroup.execute Execute a complete ExecutionGroup.
* Halts until finished or breaked by user
* \see ExecutionGroup.schedule_chunk_when_possible Tries to schedule a single chunk,
* checks if all input data is available. Can trigger dependent chunks to be calculated
* \see ExecutionGroup.schedule_area_when_possible
* Tries to schedule an area. This can be multiple chunks
* (is called from [@ref ExecutionGroup.schedule_chunk_when_possible])
* \see ExecutionGroup.schedule_chunk Schedule a chunk on the WorkScheduler
* \see NodeOperation.determine_depending_area_of_interest Influence the area of interest of a chunk.
* \see WriteBufferOperation Operation to write to a MemoryProxy/MemoryBuffer
* \see ReadBufferOperation Operation to read from a MemoryProxy/MemoryBuffer
* \see MemoryProxy proxy for information about memory image
* (a image consist out of multiple chunks)
* \see MemoryBuffer Allocated memory for a single chunk
*
* \section workscheduler WorkScheduler
* the WorkScheduler is implemented as a static class. the responsibility of the WorkScheduler
@ -251,47 +62,6 @@ struct Render;
* This is done by changing the `COM_threading_model`
* to `ThreadingModel::SingleThreaded`. When compiling the work-scheduler
* will be changes to support no threading and run everything on the CPU.
*
* \section devices Devices
* A Device within the compositor context is a Hardware component that can used to calculate chunks.
* This chunk is encapsulated in a WorkPackage.
* the WorkScheduler controls the devices and selects the device where a
* WorkPackage will be calculated.
*
* \subsection WS_Devices Work-scheduler
* The WorkScheduler controls all Devices.
* When initializing the compositor the WorkScheduler selects all
* devices that will be used during compositor.
* There are two types of Devices, CPUDevice and OpenCLDevice.
* When an ExecutionGroup schedules a Chunk the schedule method of the WorkScheduler
* The Workscheduler determines if the chunk can be run on an OpenCLDevice
* (and that there are available OpenCLDevice).
* If this is the case the chunk will be added to the work-list for OpenCLDevice's
* otherwise the chunk will be added to the work-list of CPUDevices.
*
* A thread will read the work-list and sends a work-package to its device.
*
* \see WorkScheduler.schedule method that is called to schedule a chunk
* \see Device.execute method called to execute a chunk
*
* \subsection CPUDevice CPUDevice
* When a CPUDevice gets a WorkPackage the Device will get the input-buffer that is needed to
* calculate the chunk. Allocation is already done by the ExecutionGroup.
* The output-buffer of the chunk is being created.
* The OutputOperation of the ExecutionGroup is called to execute the area of the output-buffer.
*
* \see ExecutionGroup
* \see NodeOperation.execute_region executes a single chunk of a NodeOperation
* \see CPUDevice.execute
*
* \subsection GPUDevice OpenCLDevice
*
* To be completed!
* \see NodeOperation.execute_opencl_region
* \see OpenCLDevice.execute
*
* \section execute_pixel executing a pixel
* Finally the last step, the node functionality :)
*/
/**

32
source/blender/compositor/COM_defines.h
View File

@ -12,16 +12,6 @@ namespace blender::compositor {
using Size2f = float2;
enum class eExecutionModel {
/**
* Operations are executed from outputs to inputs grouped in execution groups and rendered
* in tiles.
*/
Tiled,
/** Operations are fully rendered in order from inputs to outputs. */
FullFrame
};
enum class eDimension { X, Y };
/**
@ -84,28 +74,6 @@ constexpr DataType COM_num_channels_data_type(const int num_channels)
}
}
/* Configurable items.
*
* Chunk size determination.
*
* Chunk order. */
/**
* \brief The order of chunks to be scheduled
* \ingroup Execution
*/
enum class ChunkOrdering {
/** \brief order from a distance to centerX/centerY */
CenterOut = 0,
/** \brief order randomly */
Random = 1,
/** \brief no ordering */
TopDown = 2,
/** \brief experimental ordering with 9 hot-spots. */
RuleOfThirds = 3,
Default = ChunkOrdering::CenterOut,
};
constexpr float COM_PREVIEW_SIZE = 140.f;
constexpr float COM_RULE_OF_THIRDS_DIVIDER = 100.0f;
constexpr float COM_BLUR_BOKEH_PIXELS = 512;

2
source/blender/compositor/COM_precomp.h
View File

@ -23,12 +23,10 @@
#include "COM_ConvertOperation.h"
#include "COM_Debug.h"
#include "COM_Enums.h"
#include "COM_ExecutionGroup.h"
#include "COM_ExecutionSystem.h"
#include "COM_MultiThreadedOperation.h"
#include "COM_Node.h"
#include "COM_NodeOperation.h"
#include "COM_OpenCLDevice.h"
#include "COM_SetAlphaMultiplyOperation.h"
#include "COM_SetColorOperation.h"
#include "COM_SetSamplerOperation.h"

82
source/blender/compositor/intern/COM_BufferOperation.cc
View File

@ -1,82 +0,0 @@
/* SPDX-FileCopyrightText: 2021 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#include "COM_BufferOperation.h"
namespace blender::compositor {
BufferOperation::BufferOperation(MemoryBuffer *buffer, DataType data_type)
{
buffer_ = buffer;
inflated_buffer_ = nullptr;
set_canvas(buffer->get_rect());
add_output_socket(data_type);
flags_.is_constant_operation = buffer_->is_a_single_elem();
flags_.is_fullframe_operation = false;
}
const float *BufferOperation::get_constant_elem()
{
BLI_assert(buffer_->is_a_single_elem());
return buffer_->get_buffer();
}
void BufferOperation::init_execution()
{
if (buffer_->is_a_single_elem()) {
init_mutex();
}
}
void *BufferOperation::initialize_tile_data(rcti * /*rect*/)
{
if (buffer_->is_a_single_elem() == false) {
return buffer_;
}
lock_mutex();
if (!inflated_buffer_) {
inflated_buffer_ = buffer_->inflate();
}
unlock_mutex();
return inflated_buffer_;
}
void BufferOperation::deinit_execution()
{
if (buffer_->is_a_single_elem()) {
deinit_mutex();
}
delete inflated_buffer_;
}
void BufferOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
switch (sampler) {
case PixelSampler::Nearest:
buffer_->read(output, x, y);
break;
case PixelSampler::Bilinear:
default:
buffer_->read_bilinear(output, x, y);
break;
case PixelSampler::Bicubic:
/* No bicubic. Same implementation as ReadBufferOperation. */
buffer_->read_bilinear(output, x, y);
break;
}
}
void BufferOperation::execute_pixel_filtered(
float output[4], float x, float y, float dx[2], float dy[2])
{
const float uv[2] = {x, y};
const float deriv[2][2] = {{dx[0], dx[1]}, {dy[0], dy[1]}};
buffer_->readEWA(output, uv, deriv);
}
} // namespace blender::compositor

28
source/blender/compositor/intern/COM_BufferOperation.h
View File

@ -1,28 +0,0 @@
/* SPDX-FileCopyrightText: 2021 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#pragma once
#include "COM_ConstantOperation.h"
namespace blender::compositor {
class BufferOperation : public ConstantOperation {
private:
MemoryBuffer *buffer_;
MemoryBuffer *inflated_buffer_;
public:
BufferOperation(MemoryBuffer *buffer, DataType data_type);
const float *get_constant_elem() override;
void *initialize_tile_data(rcti *rect) override;
void init_execution() override;
void deinit_execution() override;
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
void execute_pixel_filtered(
float output[4], float x, float y, float dx[2], float dy[2]) override;
};
} // namespace blender::compositor

19
source/blender/compositor/intern/COM_CPUDevice.cc
View File

@ -4,8 +4,7 @@
#include "COM_CPUDevice.h"
#include "COM_ExecutionGroup.h"
#include "COM_NodeOperation.h"
#include "COM_WorkPackage.h"
namespace blender::compositor {
@ -13,21 +12,7 @@ CPUDevice::CPUDevice(int thread_id) : thread_id_(thread_id) {}
void CPUDevice::execute(WorkPackage *work_package)
{
switch (work_package->type) {
case eWorkPackageType::Tile: {
const uint chunk_number = work_package->chunk_number;
ExecutionGroup *execution_group = work_package->execution_group;
execution_group->get_output_operation()->execute_region(&work_package->rect, chunk_number);
execution_group->finalize_chunk_execution(chunk_number, nullptr);
break;
}
case eWorkPackageType::CustomFunction: {
work_package->execute_fn();
break;
}
}
work_package->execute_fn();
if (work_package->executed_fn) {
work_package->executed_fn();
}

28
source/blender/compositor/intern/COM_ChunkOrder.cc
View File

@ -1,28 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#include <cfloat>
#include "COM_ChunkOrder.h"
namespace blender::compositor {
void ChunkOrder::update_distance(ChunkOrderHotspot *hotspots, uint len_hotspots)
{
double new_distance = DBL_MAX;
for (int index = 0; index < len_hotspots; index++) {
double distance_to_hotspot = hotspots[index].calc_distance(x, y);
if (distance_to_hotspot < new_distance) {
new_distance = distance_to_hotspot;
}
}
this->distance = new_distance;
}
bool operator<(const ChunkOrder &a, const ChunkOrder &b)
{
return a.distance < b.distance;
}
} // namespace blender::compositor

33
source/blender/compositor/intern/COM_ChunkOrder.h
View File

@ -1,33 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#pragma once
#ifdef WITH_CXX_GUARDEDALLOC
# include "MEM_guardedalloc.h"
#endif
#include "BLI_sys_types.h"
#include "COM_ChunkOrderHotspot.h"
namespace blender::compositor {
/** Helper to determine the order how chunks are prioritized during execution. */
struct ChunkOrder {
uint index = 0;
int x = 0;
int y = 0;
double distance = 0.0;
friend bool operator<(const ChunkOrder &a, const ChunkOrder &b);
void update_distance(ChunkOrderHotspot *hotspots, uint len_hotspots);
#ifdef WITH_CXX_GUARDEDALLOC
MEM_CXX_CLASS_ALLOC_FUNCS("COM:ChunkOrderHotspot")
#endif
};
} // namespace blender::compositor

19
source/blender/compositor/intern/COM_ChunkOrderHotspot.cc
View File

@ -1,19 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#include "COM_ChunkOrderHotspot.h"
#include <cmath>
namespace blender::compositor {
double ChunkOrderHotspot::calc_distance(int x, int y)
{
int dx = this->x - x;
int dy = this->y - y;
double result = sqrt(double(dx * dx + dy * dy));
result += double(this->addition);
return result;
}
} // namespace blender::compositor

27
source/blender/compositor/intern/COM_ChunkOrderHotspot.h
View File

@ -1,27 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#pragma once
#ifdef WITH_CXX_GUARDEDALLOC
# include "MEM_guardedalloc.h"
#endif
namespace blender::compositor {
struct ChunkOrderHotspot {
int x;
int y;
float addition;
ChunkOrderHotspot(int x, int y, float addition) : x(x), y(y), addition(addition) {}
double calc_distance(int x, int y);
#ifdef WITH_CXX_GUARDEDALLOC
MEM_CXX_CLASS_ALLOC_FUNCS("COM:ChunkOrderHotspot")
#endif
};
} // namespace blender::compositor

6
source/blender/compositor/intern/COM_CompositorContext.cc
View File

@ -11,7 +11,6 @@ CompositorContext::CompositorContext()
scene_ = nullptr;
rd_ = nullptr;
quality_ = eCompositorQuality::High;
hasActiveOpenCLDevices_ = false;
fast_calculation_ = false;
bnodetree_ = nullptr;
}
@ -28,9 +27,4 @@ Size2f CompositorContext::get_render_size() const
get_render_data()->ysch * get_render_percentage_as_factor()};
}
eExecutionModel CompositorContext::get_execution_model() const
{
return eExecutionModel::FullFrame;
}
} // namespace blender::compositor

35
source/blender/compositor/intern/COM_CompositorContext.h
View File

@ -59,11 +59,6 @@ class CompositorContext {
*/
bNodeInstanceHash *previews_;
/**
* \brief does this system have active opencl devices?
*/
bool hasActiveOpenCLDevices_;
/**
* \brief Skip slow nodes
*/
@ -180,22 +175,6 @@ class CompositorContext {
*/
int get_framenumber() const;
/**
* \brief has this system active opencl_devices?
*/
bool get_has_active_opencl_devices() const
{
return hasActiveOpenCLDevices_;
}
/**
* \brief set has this system active opencl_devices?
*/
void setHasActiveOpenCLDevices(bool hasAvtiveOpenCLDevices)
{
hasActiveOpenCLDevices_ = hasAvtiveOpenCLDevices;
}
/** Whether it has a view with a specific name and not the default one. */
bool has_explicit_view() const
{
@ -234,11 +213,6 @@ class CompositorContext {
view_name_ = view_name;
}
int get_chunksize() const
{
return 256;
}
void set_fast_calculation(bool fast_calculation)
{
fast_calculation_ = fast_calculation;
@ -247,10 +221,6 @@ class CompositorContext {
{
return fast_calculation_;
}
bool is_groupnode_buffer_enabled() const
{
return false;
}
/**
* \brief Get the render percentage as a factor.
@ -262,11 +232,6 @@ class CompositorContext {
}
Size2f get_render_size() const;
/**
* Get active execution model.
*/
eExecutionModel get_execution_model() const;
};
} // namespace blender::compositor

2
source/blender/compositor/intern/COM_ConstantFolder.cc
View File

@ -133,7 +133,7 @@ Vector<ConstantOperation *> ConstantFolder::try_fold_operations(Span<NodeOperati
int ConstantFolder::fold_operations()
{
WorkScheduler::start(operations_builder_.context());
WorkScheduler::start();
Vector<ConstantOperation *> last_folds = try_fold_operations(
operations_builder_.get_operations());
int folds_count = last_folds.size();

12
source/blender/compositor/intern/COM_Converter.cc
View File

@ -550,13 +550,11 @@ void COM_convert_canvas(NodeOperationBuilder &builder,
builder.add_operation(syop);
rcti scale_canvas = from_operation->get_canvas();
if (builder.context().get_execution_model() == eExecutionModel::FullFrame) {
ScaleOperation::scale_area(scale_canvas, scaleX, scaleY);
scale_canvas.xmax = scale_canvas.xmin + to_operation->get_width();
scale_canvas.ymax = scale_canvas.ymin + to_operation->get_height();
addX = 0;
addY = 0;
}
ScaleOperation::scale_area(scale_canvas, scaleX, scaleY);
scale_canvas.xmax = scale_canvas.xmin + to_operation->get_width();
scale_canvas.ymax = scale_canvas.ymin + to_operation->get_height();
addX = 0;
addY = 0;
scale_operation->set_canvas(scale_canvas);
sxop->set_canvas(scale_canvas);
syop->set_canvas(scale_canvas);

101
source/blender/compositor/intern/COM_Debug.cc
View File

@ -12,11 +12,8 @@
#include "IMB_imbuf.hh"
#include "IMB_imbuf_types.hh"
#include "COM_ExecutionGroup.h"
#include "COM_ReadBufferOperation.h"
#include "COM_SetValueOperation.h"
#include "COM_ViewerOperation.h"
#include "COM_WriteBufferOperation.h"
namespace blender::compositor {
@ -25,7 +22,6 @@ DebugInfo::NodeNameMap DebugInfo::node_names_;
DebugInfo::OpNameMap DebugInfo::op_names_;
std::string DebugInfo::current_node_name_;
std::string DebugInfo::current_op_name_;
DebugInfo::GroupStateMap DebugInfo::group_states_;
static std::string operation_class_name(const NodeOperation *op)
{
@ -56,7 +52,6 @@ std::string DebugInfo::operation_name(const NodeOperation *op)
int DebugInfo::graphviz_operation(const ExecutionSystem *system,
NodeOperation *operation,
const ExecutionGroup *group,
char *str,
int maxlen)
{
@ -78,20 +73,9 @@ int DebugInfo::graphviz_operation(const ExecutionSystem *system,
else if (operation->get_flags().is_set_operation) {
fillcolor = "khaki1";
}
else if (operation->get_flags().is_read_buffer_operation) {
fillcolor = "darkolivegreen3";
}
else if (operation->get_flags().is_write_buffer_operation) {
fillcolor = "darkorange";
}
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "// OPERATION: %p\r\n", operation);
if (group) {
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "\"O_%p_%p\"", operation, group);
}
else {
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "\"O_%p\"", operation);
}
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "\"O_%p\"", operation);
len += snprintf(str + len,
maxlen > len ? maxlen - len : 0,
" [fillcolor=%s,style=filled,shape=record,label=\"{",
@ -220,14 +204,11 @@ int DebugInfo::graphviz_legend_group(
return len;
}
int DebugInfo::graphviz_legend(char *str, int maxlen, const bool has_execution_groups)
int DebugInfo::graphviz_legend(char *str, int maxlen)
{
int len = 0;
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "{\r\n");
if (has_execution_groups) {
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "rank = sink;\r\n");
}
len += snprintf(
str + len, maxlen > len ? maxlen - len : 0, "Legend [shape=none, margin=0, label=<\r\n");
@ -247,25 +228,9 @@ int DebugInfo::graphviz_legend(char *str, int maxlen, const bool has_execution_g
"Viewer", "lightskyblue3", str + len, maxlen > len ? maxlen - len : 0);
len += graphviz_legend_color(
"Active Viewer", "lightskyblue1", str + len, maxlen > len ? maxlen - len : 0);
if (has_execution_groups) {
len += graphviz_legend_color(
"Write Buffer", "darkorange", str + len, maxlen > len ? maxlen - len : 0);
len += graphviz_legend_color(
"Read Buffer", "darkolivegreen3", str + len, maxlen > len ? maxlen - len : 0);
}
len += graphviz_legend_color(
"Input Value", "khaki1", str + len, maxlen > len ? maxlen - len : 0);
if (has_execution_groups) {
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "<TR><TD></TD></TR>\r\n");
len += graphviz_legend_group(
"Group Waiting", "white", "dashed", str + len, maxlen > len ? maxlen - len : 0);
len += graphviz_legend_group(
"Group Running", "firebrick1", "solid", str + len, maxlen > len ? maxlen - len : 0);
len += graphviz_legend_group(
"Group Finished", "chartreuse4", "solid", str + len, maxlen > len ? maxlen - len : 0);
}
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "</TABLE>\r\n");
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, ">];\r\n");
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "}\r\n");
@ -275,7 +240,6 @@ int DebugInfo::graphviz_legend(char *str, int maxlen, const bool has_execution_g
bool DebugInfo::graphviz_system(const ExecutionSystem *system, char *str, int maxlen)
{
char strbuf[64];
int len = 0;
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "digraph compositorexecution {\r\n");
@ -284,39 +248,7 @@ bool DebugInfo::graphviz_system(const ExecutionSystem *system, char *str, int ma
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "splines=false\r\n");
std::map<NodeOperation *, std::vector<std::string>> op_groups;
int index = 0;
for (const ExecutionGroup *group : system->groups_) {
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "// GROUP: %d\r\n", index);
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "subgraph cluster_%d{\r\n", index);
/* used as a check for executing group */
if (group_states_[group] == EG_WAIT) {
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "style=dashed\r\n");
}
else if (group_states_[group] == EG_RUNNING) {
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "style=filled\r\n");
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "color=black\r\n");
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "fillcolor=firebrick1\r\n");
}
else if (group_states_[group] == EG_FINISHED) {
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "style=filled\r\n");
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "color=black\r\n");
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "fillcolor=chartreuse4\r\n");
}
for (NodeOperation *operation : group->operations_) {
SNPRINTF(strbuf, "_%p", group);
op_groups[operation].push_back(std::string(strbuf));
len += graphviz_operation(
system, operation, group, str + len, maxlen > len ? maxlen - len : 0);
}
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "}\r\n");
index++;
}
/* operations not included in any group */
for (NodeOperation *operation : system->operations_) {
if (op_groups.find(operation) != op_groups.end()) {
continue;
@ -324,29 +256,7 @@ bool DebugInfo::graphviz_system(const ExecutionSystem *system, char *str, int ma
op_groups[operation].push_back(std::string(""));
len += graphviz_operation(
system, operation, nullptr, str + len, maxlen > len ? maxlen - len : 0);
}
for (NodeOperation *operation : system->operations_) {
if (operation->get_flags().is_read_buffer_operation) {
ReadBufferOperation *read = (ReadBufferOperation *)operation;
WriteBufferOperation *write = read->get_memory_proxy()->get_write_buffer_operation();
std::vector<std::string> &read_groups = op_groups[read];
std::vector<std::string> &write_groups = op_groups[write];
for (int k = 0; k < write_groups.size(); k++) {
for (int l = 0; l < read_groups.size(); l++) {
len += snprintf(str + len,
maxlen > len ? maxlen - len : 0,
"\"O_%p%s\" -> \"O_%p%s\" [style=dotted]\r\n",
write,
write_groups[k].c_str(),
read,
read_groups[l].c_str());
}
}
}
len += graphviz_operation(system, operation, str + len, maxlen > len ? maxlen - len : 0);
}
for (NodeOperation *op : system->operations_) {
@ -401,10 +311,7 @@ bool DebugInfo::graphviz_system(const ExecutionSystem *system, char *str, int ma
}
}
const bool has_execution_groups = system->get_context().get_execution_model() ==
eExecutionModel::Tiled &&
system->groups_.size() > 0;
len += graphviz_legend(str + len, maxlen > len ? maxlen - len : 0, has_execution_groups);
len += graphviz_legend(str + len, maxlen > len ? maxlen - len : 0);
len += snprintf(str + len, maxlen > len ? maxlen - len : 0, "}\r\n");

28
source/blender/compositor/intern/COM_Debug.h
View File

@ -24,15 +24,11 @@ static constexpr bool COM_EXPORT_OPERATION_BUFFERS = false;
class Node;
class NodeOperation;
class ExecutionSystem;
class ExecutionGroup;
class DebugInfo {
public:
typedef enum { EG_WAIT, EG_RUNNING, EG_FINISHED } GroupState;
typedef std::map<const Node *, std::string> NodeNameMap;
typedef std::map<const NodeOperation *, std::string> OpNameMap;
typedef std::map<const ExecutionGroup *, GroupState> GroupStateMap;
static std::string node_name(const Node *node);
static std::string operation_name(const NodeOperation *op);
@ -47,8 +43,6 @@ class DebugInfo {
static std::string current_node_name_;
/** Base name for automatic sub-operations. */
static std::string current_op_name_;
/** For visualizing group states. */
static GroupStateMap group_states_;
public:
static void convert_started()
@ -58,14 +52,10 @@ class DebugInfo {
}
}
static void execute_started(const ExecutionSystem *system)
static void execute_started()
{
if (COM_EXPORT_GRAPHVIZ) {
file_index_ = 1;
group_states_.clear();
for (ExecutionGroup *execution_group : system->groups_) {
group_states_[execution_group] = EG_WAIT;
}
}
if (COM_EXPORT_OPERATION_BUFFERS) {
delete_operation_exports();
@ -100,19 +90,6 @@ class DebugInfo {
}
};
static void execution_group_started(const ExecutionGroup *group)
{
if (COM_EXPORT_GRAPHVIZ) {
group_states_[group] = EG_RUNNING;
}
};
static void execution_group_finished(const ExecutionGroup *group)
{
if (COM_EXPORT_GRAPHVIZ) {
group_states_[group] = EG_FINISHED;
}
};
static void operation_rendered(const NodeOperation *op, MemoryBuffer *render)
{
/* Don't export constant operations as there are too many and it's rarely useful. */
@ -126,7 +103,6 @@ class DebugInfo {
protected:
static int graphviz_operation(const ExecutionSystem *system,
NodeOperation *operation,
const ExecutionGroup *group,
char *str,
int maxlen);
static int graphviz_legend_color(const char *name, const char *color, char *str, int maxlen);
@ -134,7 +110,7 @@ class DebugInfo {
const char *name, const char *color, const char *style, char *str, int maxlen);
static int graphviz_legend_group(
const char *name, const char *color, const char *style, char *str, int maxlen);
static int graphviz_legend(char *str, int maxlen, bool has_execution_groups);
static int graphviz_legend(char *str, int maxlen);
static bool graphviz_system(const ExecutionSystem *system, char *str, int maxlen);
static void export_operation(const NodeOperation *op, MemoryBuffer *render);

19
source/blender/compositor/intern/COM_Enums.cc
View File

@ -43,23 +43,4 @@ std::ostream &operator<<(std::ostream &os, const eCompositorPriority &priority)
return os;
}
std::ostream &operator<<(std::ostream &os, const eWorkPackageState &execution_state)
{
switch (execution_state) {
case eWorkPackageState::NotScheduled: {
os << "ExecutionState::NotScheduled";
break;
}
case eWorkPackageState::Scheduled: {
os << "ExecutionState::Scheduled";
break;
}
case eWorkPackageState::Executed: {
os << "ExecutionState::Executed";
break;
}
}
return os;
}
} // namespace blender::compositor

38
source/blender/compositor/intern/COM_Enums.h
View File

@ -31,48 +31,11 @@ enum class eCompositorQuality {
* \ingroup Execution
*/
enum class eCompositorPriority {
/** \brief High quality setting */
High = 2,
/** \brief Medium quality setting */
Medium = 1,
/** \brief Low quality setting */
Low = 0,
};
/**
* \brief the execution state of a chunk in an ExecutionGroup
* \ingroup Execution
*/
enum class eWorkPackageState {
/**
* \brief chunk is not yet scheduled
*/
NotScheduled = 0,
/**
* \brief chunk is scheduled, but not yet executed
*/
Scheduled = 1,
/**
* \brief chunk is executed.
*/
Executed = 2,
};
/**
* \brief Work type to execute.
* \ingroup Execution
*/
enum class eWorkPackageType {
/**
* \brief Executes an execution group tile.
*/
Tile = 0,
/**
* \brief Executes a custom function.
*/
CustomFunction = 1
};
enum class PixelSampler {
Nearest = 0,
Bilinear = 1,
@ -81,6 +44,5 @@ enum class PixelSampler {
void expand_area_for_sampler(rcti &area, PixelSampler sampler);
std::ostream &operator<<(std::ostream &os, const eCompositorPriority &priority);
std::ostream &operator<<(std::ostream &os, const eWorkPackageState &execution_state);
} // namespace blender::compositor

583
source/blender/compositor/intern/COM_ExecutionGroup.cc
View File

@ -1,583 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#include "COM_ExecutionGroup.h"
#include "COM_ChunkOrder.h"
#include "COM_Debug.h"
#include "COM_ReadBufferOperation.h"
#include "COM_ViewerOperation.h"
#include "COM_WorkScheduler.h"
#include "COM_WriteBufferOperation.h"
#include "COM_defines.h"
#include "BLI_rand.hh"
#include "BLI_string.h"
#include "BLI_time.h"
#include "BLT_translation.hh"
namespace blender::compositor {
std::ostream &operator<<(std::ostream &os, const ExecutionGroupFlags &flags)
{
if (flags.initialized) {
os << "init,";
}
if (flags.is_output) {
os << "output,";
}
if (flags.complex) {
os << "complex,";
}
if (flags.open_cl) {
os << "open_cl,";
}
if (flags.single_threaded) {
os << "single_threaded,";
}
return os;
}
ExecutionGroup::ExecutionGroup(int id)
{
id_ = id;
bTree_ = nullptr;
height_ = 0;
width_ = 0;
max_read_buffer_offset_ = 0;
x_chunks_len_ = 0;
y_chunks_len_ = 0;
chunks_len_ = 0;
chunks_finished_ = 0;
BLI_rcti_init(&viewer_border_, 0, 0, 0, 0);
execution_start_time_ = 0;
}
std::ostream &operator<<(std::ostream &os, const ExecutionGroup &execution_group)
{
os << "ExecutionGroup(id=" << execution_group.get_id();
os << ",flags={" << execution_group.get_flags() << "}";
os << ",operation=" << *execution_group.get_output_operation() << "";
os << ")";
return os;
}
eCompositorPriority ExecutionGroup::get_render_priority()
{
return this->get_output_operation()->get_render_priority();
}
bool ExecutionGroup::can_contain(NodeOperation &operation)
{
if (!flags_.initialized) {
return true;
}
if (operation.get_flags().is_read_buffer_operation) {
return true;
}
if (operation.get_flags().is_write_buffer_operation) {
return false;
}
if (operation.get_flags().is_set_operation) {
return true;
}
/* complex groups don't allow further ops (except read buffer and values, see above) */
if (flags_.complex) {
return false;
}
/* complex ops can't be added to other groups (except their own, which they initialize, see
* above) */
if (operation.get_flags().complex) {
return false;
}
return true;
}
bool ExecutionGroup::add_operation(NodeOperation *operation)
{
if (!can_contain(*operation)) {
return false;
}
if (!operation->get_flags().is_read_buffer_operation &&
!operation->get_flags().is_write_buffer_operation)
{
flags_.complex = operation->get_flags().complex;
flags_.open_cl = operation->get_flags().open_cl;
flags_.single_threaded = operation->get_flags().single_threaded;
flags_.initialized = true;
}
operations_.append(operation);
return true;
}
NodeOperation *ExecutionGroup::get_output_operation() const
{
/* The first operation of the group is always the output operation. */
return this->operations_[0];
}
void ExecutionGroup::init_work_packages()
{
work_packages_.clear();
if (chunks_len_ != 0) {
work_packages_.resize(chunks_len_);
for (uint index = 0; index < chunks_len_; index++) {
work_packages_[index].type = eWorkPackageType::Tile;
work_packages_[index].state = eWorkPackageState::NotScheduled;
work_packages_[index].execution_group = this;
work_packages_[index].chunk_number = index;
determine_chunk_rect(&work_packages_[index].rect, index);
}
}
}
void ExecutionGroup::init_read_buffer_operations()
{
uint max_offset = 0;
for (NodeOperation *operation : operations_) {
if (operation->get_flags().is_read_buffer_operation) {
ReadBufferOperation *read_operation = static_cast<ReadBufferOperation *>(operation);
read_operations_.append(read_operation);
max_offset = std::max(max_offset, read_operation->get_offset());
}
}
max_offset++;
max_read_buffer_offset_ = max_offset;
}
void ExecutionGroup::init_execution()
{
init_number_of_chunks();
init_work_packages();
init_read_buffer_operations();
}
void ExecutionGroup::deinit_execution()
{
work_packages_.clear();
chunks_len_ = 0;
x_chunks_len_ = 0;
y_chunks_len_ = 0;
read_operations_.clear();
bTree_ = nullptr;
}
void ExecutionGroup::determine_resolution(uint resolution[2])
{
NodeOperation *operation = this->get_output_operation();
resolution[0] = operation->get_width();
resolution[1] = operation->get_height();
this->set_resolution(resolution);
BLI_rcti_init(&viewer_border_, 0, width_, 0, height_);
}
void ExecutionGroup::init_number_of_chunks()
{
if (flags_.single_threaded) {
x_chunks_len_ = 1;
y_chunks_len_ = 1;
chunks_len_ = 1;
}
else {
const float chunk_sizef = chunk_size_;
const int border_width = BLI_rcti_size_x(&viewer_border_);
const int border_height = BLI_rcti_size_y(&viewer_border_);
x_chunks_len_ = ceil(border_width / chunk_sizef);
y_chunks_len_ = ceil(border_height / chunk_sizef);
chunks_len_ = x_chunks_len_ * y_chunks_len_;
}
}
blender::Array<uint> ExecutionGroup::get_execution_order() const
{
blender::Array<uint> chunk_order(chunks_len_);
for (int chunk_index = 0; chunk_index < chunks_len_; chunk_index++) {
chunk_order[chunk_index] = chunk_index;
}
NodeOperation *operation = this->get_output_operation();
float centerX = 0.5f;
float centerY = 0.5f;
ChunkOrdering order_type = ChunkOrdering::Default;
if (operation->get_flags().is_viewer_operation) {
ViewerOperation *viewer = (ViewerOperation *)operation;
centerX = viewer->getCenterX();
centerY = viewer->getCenterY();
order_type = viewer->get_chunk_order();
}
const int border_width = BLI_rcti_size_x(&viewer_border_);
const int border_height = BLI_rcti_size_y(&viewer_border_);
int index;
switch (order_type) {
case ChunkOrdering::Random: {
static blender::RandomNumberGenerator rng;
blender::MutableSpan<uint> span = chunk_order.as_mutable_span();
/* Shuffle twice to make it more random. */
rng.shuffle(span);
rng.shuffle(span);
break;
}
case ChunkOrdering::CenterOut: {
ChunkOrderHotspot hotspot(border_width * centerX, border_height * centerY, 0.0f);
blender::Array<ChunkOrder> chunk_orders(chunks_len_);
for (index = 0; index < chunks_len_; index++) {
const WorkPackage &work_package = work_packages_[index];
chunk_orders[index].index = index;
chunk_orders[index].x = work_package.rect.xmin - viewer_border_.xmin;
chunk_orders[index].y = work_package.rect.ymin - viewer_border_.ymin;
chunk_orders[index].update_distance(&hotspot, 1);
}
std::sort(&chunk_orders[0], &chunk_orders[chunks_len_ - 1]);
for (index = 0; index < chunks_len_; index++) {
chunk_order[index] = chunk_orders[index].index;
}
break;
}
case ChunkOrdering::RuleOfThirds: {
uint tx = border_width / 6;
uint ty = border_height / 6;
uint mx = border_width / 2;
uint my = border_height / 2;
uint bx = mx + 2 * tx;
uint by = my + 2 * ty;
float addition = chunks_len_ / COM_RULE_OF_THIRDS_DIVIDER;
ChunkOrderHotspot hotspots[9] = {
ChunkOrderHotspot(mx, my, addition * 0),
ChunkOrderHotspot(tx, my, addition * 1),
ChunkOrderHotspot(bx, my, addition * 2),
ChunkOrderHotspot(bx, by, addition * 3),
ChunkOrderHotspot(tx, ty, addition * 4),
ChunkOrderHotspot(bx, ty, addition * 5),
ChunkOrderHotspot(tx, by, addition * 6),
ChunkOrderHotspot(mx, ty, addition * 7),
ChunkOrderHotspot(mx, by, addition * 8),
};
blender::Array<ChunkOrder> chunk_orders(chunks_len_);
for (index = 0; index < chunks_len_; index++) {
const WorkPackage &work_package = work_packages_[index];
chunk_orders[index].index = index;
chunk_orders[index].x = work_package.rect.xmin - viewer_border_.xmin;
chunk_orders[index].y = work_package.rect.ymin - viewer_border_.ymin;
chunk_orders[index].update_distance(hotspots, 9);
}
std::sort(&chunk_orders[0], &chunk_orders[chunks_len_]);
for (index = 0; index < chunks_len_; index++) {
chunk_order[index] = chunk_orders[index].index;
}
break;
}
case ChunkOrdering::TopDown:
default:
break;
}
return chunk_order;
}
void ExecutionGroup::execute(ExecutionSystem *graph)
{
const CompositorContext &context = graph->get_context();
const bNodeTree *bTree = context.get_bnodetree();
if (width_ == 0 || height_ == 0) {
return;
} /** \note Break out... no pixels to calculate. */
if (bTree->runtime->test_break && bTree->runtime->test_break(bTree->runtime->tbh)) {
return;
} /** \note Early break out for blur and preview nodes. */
if (chunks_len_ == 0) {
return;
} /** \note Early break out. */
uint chunk_index;
execution_start_time_ = BLI_time_now_seconds();
chunks_finished_ = 0;
bTree_ = bTree;
blender::Array<uint> chunk_order = get_execution_order();
DebugInfo::execution_group_started(this);
DebugInfo::graphviz(graph);
bool breaked = false;
bool finished = false;
uint start_index = 0;
const int max_number_evaluated = BLI_system_thread_count() * 2;
while (!finished && !breaked) {
bool start_evaluated = false;
finished = true;
int number_evaluated = 0;
for (int index = start_index; index < chunks_len_ && number_evaluated < max_number_evaluated;
index++)
{
chunk_index = chunk_order[index];
int y_chunk = chunk_index / x_chunks_len_;
int x_chunk = chunk_index - (y_chunk * x_chunks_len_);
const WorkPackage &work_package = work_packages_[chunk_index];
switch (work_package.state) {
case eWorkPackageState::NotScheduled: {
schedule_chunk_when_possible(graph, x_chunk, y_chunk);
finished = false;
start_evaluated = true;
number_evaluated++;
if (bTree->runtime->update_draw) {
bTree->runtime->update_draw(bTree->runtime->udh);
}
break;
}
case eWorkPackageState::Scheduled: {
finished = false;
start_evaluated = true;
number_evaluated++;
break;
}
case eWorkPackageState::Executed: {
if (!start_evaluated) {
start_index = index + 1;
}
}
};
}
WorkScheduler::finish();
if (bTree->runtime->test_break && bTree->runtime->test_break(bTree->runtime->tbh)) {
breaked = true;
}
}
DebugInfo::execution_group_finished(this);
DebugInfo::graphviz(graph);
}
MemoryBuffer **ExecutionGroup::get_input_buffers_opencl(int chunk_number)
{
WorkPackage &work_package = work_packages_[chunk_number];
MemoryBuffer **memory_buffers = (MemoryBuffer **)MEM_callocN(
sizeof(MemoryBuffer *) * max_read_buffer_offset_, __func__);
rcti output;
for (ReadBufferOperation *read_operation : read_operations_) {
MemoryProxy *memory_proxy = read_operation->get_memory_proxy();
this->determine_depending_area_of_interest(&work_package.rect, read_operation, &output);
MemoryBuffer *memory_buffer =
memory_proxy->get_executor()->construct_consolidated_memory_buffer(*memory_proxy, output);
memory_buffers[read_operation->get_offset()] = memory_buffer;
}
return memory_buffers;
}
MemoryBuffer *ExecutionGroup::construct_consolidated_memory_buffer(MemoryProxy &memory_proxy,
rcti &rect)
{
MemoryBuffer *image_buffer = memory_proxy.get_buffer();
MemoryBuffer *result = new MemoryBuffer(&memory_proxy, rect, MemoryBufferState::Temporary);
result->fill_from(*image_buffer);
return result;
}
void ExecutionGroup::finalize_chunk_execution(int chunk_number, MemoryBuffer **memory_buffers)
{
WorkPackage &work_package = work_packages_[chunk_number];
if (work_package.state == eWorkPackageState::Scheduled) {
work_package.state = eWorkPackageState::Executed;
}
atomic_add_and_fetch_u(&chunks_finished_, 1);
if (memory_buffers) {
for (uint index = 0; index < max_read_buffer_offset_; index++) {
MemoryBuffer *buffer = memory_buffers[index];
if (buffer) {
if (buffer->is_temporarily()) {
memory_buffers[index] = nullptr;
delete buffer;
}
}
}
MEM_freeN(memory_buffers);
}
if (bTree_) {
/* Status report is only performed for top level Execution Groups. */
float progress = chunks_finished_;
progress /= chunks_len_;
bTree_->runtime->progress(bTree_->runtime->prh, progress);
char buf[128];
SNPRINTF(buf, RPT_("Compositing | Tile %u-%u"), chunks_finished_, chunks_len_);
bTree_->runtime->stats_draw(bTree_->runtime->sdh, buf);
}
}
inline void ExecutionGroup::determine_chunk_rect(rcti *r_rect,
const uint x_chunk,
const uint y_chunk) const
{
const int border_width = BLI_rcti_size_x(&viewer_border_);
const int border_height = BLI_rcti_size_y(&viewer_border_);
if (flags_.single_threaded) {
BLI_rcti_init(r_rect, viewer_border_.xmin, border_width, viewer_border_.ymin, border_height);
}
else {
const uint minx = x_chunk * chunk_size_ + viewer_border_.xmin;
const uint miny = y_chunk * chunk_size_ + viewer_border_.ymin;
const uint width = std::min(uint(viewer_border_.xmax), width_);
const uint height = std::min(uint(viewer_border_.ymax), height_);
BLI_rcti_init(r_rect,
std::min(minx, width_),
std::min(minx + chunk_size_, width),
std::min(miny, height_),
std::min(miny + chunk_size_, height));
}
}
void ExecutionGroup::determine_chunk_rect(rcti *r_rect, const uint chunk_number) const
{
const uint y_chunk = chunk_number / x_chunks_len_;
const uint x_chunk = chunk_number - (y_chunk * x_chunks_len_);
determine_chunk_rect(r_rect, x_chunk, y_chunk);
}
MemoryBuffer *ExecutionGroup::allocate_output_buffer(rcti &rect)
{
/* We assume that this method is only called from complex execution groups. */
NodeOperation *operation = this->get_output_operation();
if (operation->get_flags().is_write_buffer_operation) {
WriteBufferOperation *write_operation = (WriteBufferOperation *)operation;
MemoryBuffer *buffer = new MemoryBuffer(
write_operation->get_memory_proxy(), rect, MemoryBufferState::Temporary);
return buffer;
}
return nullptr;
}
bool ExecutionGroup::schedule_area_when_possible(ExecutionSystem *graph, rcti *area)
{
if (flags_.single_threaded) {
return schedule_chunk_when_possible(graph, 0, 0);
}
/* Find all chunks inside the rect
* determine `minxchunk`, `minychunk`, `maxxchunk`, `maxychunk`
* where x and y are chunk-numbers. */
int indexx, indexy;
int minx = max_ii(area->xmin - viewer_border_.xmin, 0);
int maxx = min_ii(area->xmax - viewer_border_.xmin, viewer_border_.xmax - viewer_border_.xmin);
int miny = max_ii(area->ymin - viewer_border_.ymin, 0);
int maxy = min_ii(area->ymax - viewer_border_.ymin, viewer_border_.ymax - viewer_border_.ymin);
int minxchunk = minx / int(chunk_size_);
int maxxchunk = (maxx + int(chunk_size_) - 1) / int(chunk_size_);
int minychunk = miny / int(chunk_size_);
int maxychunk = (maxy + int(chunk_size_) - 1) / int(chunk_size_);
minxchunk = max_ii(minxchunk, 0);
minychunk = max_ii(minychunk, 0);
maxxchunk = min_ii(maxxchunk, int(x_chunks_len_));
maxychunk = min_ii(maxychunk, int(y_chunks_len_));
bool result = true;
for (indexx = minxchunk; indexx < maxxchunk; indexx++) {
for (indexy = minychunk; indexy < maxychunk; indexy++) {
if (!schedule_chunk_when_possible(graph, indexx, indexy)) {
result = false;
}
}
}
return result;
}
bool ExecutionGroup::schedule_chunk(uint chunk_number)
{
WorkPackage &work_package = work_packages_[chunk_number];
if (work_package.state == eWorkPackageState::NotScheduled) {
work_package.state = eWorkPackageState::Scheduled;
WorkScheduler::schedule(&work_package);
return true;
}
return false;
}
bool ExecutionGroup::schedule_chunk_when_possible(ExecutionSystem *graph,
const int chunk_x,
const int chunk_y)
{
if (chunk_x < 0 || chunk_x >= int(x_chunks_len_)) {
return true;
}
if (chunk_y < 0 || chunk_y >= int(y_chunks_len_)) {
return true;
}
/* Check if chunk is already executed or scheduled and not yet executed. */
const int chunk_index = chunk_y * x_chunks_len_ + chunk_x;
WorkPackage &work_package = work_packages_[chunk_index];
if (work_package.state == eWorkPackageState::Executed) {
return true;
}
if (work_package.state == eWorkPackageState::Scheduled) {
return false;
}
bool can_be_executed = true;
rcti area;
for (ReadBufferOperation *read_operation : read_operations_) {
BLI_rcti_init(&area, 0, 0, 0, 0);
MemoryProxy *memory_proxy = read_operation->get_memory_proxy();
determine_depending_area_of_interest(&work_package.rect, read_operation, &area);
ExecutionGroup *group = memory_proxy->get_executor();
if (!group->schedule_area_when_possible(graph, &area)) {
can_be_executed = false;
}
}
if (can_be_executed) {
schedule_chunk(chunk_index);
}
return false;
}
void ExecutionGroup::determine_depending_area_of_interest(rcti *input,
ReadBufferOperation *read_operation,
rcti *output)
{
this->get_output_operation()->determine_depending_area_of_interest(
input, read_operation, output);
}
void ExecutionGroup::set_viewer_border(float xmin, float xmax, float ymin, float ymax)
{
const NodeOperation &operation = *this->get_output_operation();
if (operation.get_flags().use_viewer_border) {
BLI_rcti_init(&viewer_border_, xmin * width_, xmax * width_, ymin * height_, ymax * height_);
}
}
void ExecutionGroup::set_render_border(float xmin, float xmax, float ymin, float ymax)
{
const NodeOperation &operation = *this->get_output_operation();
if (operation.is_output_operation(true) && operation.get_flags().use_render_border) {
BLI_rcti_init(&viewer_border_, xmin * width_, xmax * width_, ymin * height_, ymax * height_);
}
}
} // namespace blender::compositor

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

@ -1,392 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#pragma once
#ifdef WITH_CXX_GUARDEDALLOC
# include "MEM_guardedalloc.h"
#endif
#include <iostream>
#include "BLI_array.hh"
#include "BLI_vector.hh"
#include "COM_Enums.h"
#include "COM_WorkPackage.h"
#include "DNA_node_types.h"
#include "DNA_vec_types.h"
namespace blender::compositor {
class ExecutionSystem;
class NodeOperation;
class MemoryProxy;
class MemoryBuffer;
class ReadBufferOperation;
struct ExecutionGroupFlags {
bool initialized : 1;
/**
* Is this ExecutionGroup an output ExecutionGroup
* An OutputExecution group are groups containing a
* ViewerOperation, CompositeOperation, PreviewOperation.
*/
bool is_output : 1;
bool complex : 1;
/**
* Can this ExecutionGroup be scheduled on an OpenCLDevice.
*/
bool open_cl : 1;
/**
* Schedule this execution group as a single chunk. This
* chunk will be executed by a single thread.
*/
bool single_threaded : 1;
ExecutionGroupFlags()
{
initialized = false;
is_output = false;
complex = false;
open_cl = false;
single_threaded = false;
}
};
std::ostream &operator<<(std::ostream &os, const ExecutionGroupFlags &flags);
/**
* \brief Class ExecutionGroup is a group of Operations that are executed as one.
* This grouping is used to combine Operations that can be executed as one whole when
* multi-processing.
* \ingroup Execution
*/
class ExecutionGroup {
private:
// fields
/**
* Id of the execution group. For debugging purposes.
*/
int id_;
/**
* \brief list of operations in this ExecutionGroup
*/
Vector<NodeOperation *> operations_;
ExecutionGroupFlags flags_;
/**
* \brief Width of the output
*/
unsigned int width_;
/**
* \brief Height of the output
*/
unsigned int height_;
/**
* \brief size of a single chunk, being Width or of height
* a chunk is always a square, except at the edges of the MemoryBuffer
*/
unsigned int chunk_size_;
/**
* \brief number of chunks in the x-axis
*/
unsigned int x_chunks_len_;
/**
* \brief number of chunks in the y-axis
*/
unsigned int y_chunks_len_;
/**
* \brief total number of chunks
*/
unsigned int chunks_len_;
/**
* \brief what is the maximum number field of all ReadBufferOperation in this ExecutionGroup.
* \note this is used to construct the MemoryBuffers that will be passed during execution.
*/
unsigned int max_read_buffer_offset_;
/**
* \brief All read operations of this execution group.
*/
Vector<ReadBufferOperation *> read_operations_;
/**
* \brief reference to the original bNodeTree,
* this field is only set for the 'top' execution group.
* \note can only be used to call the callbacks for progress, status and break.
*/
const bNodeTree *bTree_;
/**
* \brief total number of chunks that have been calculated for this ExecutionGroup
*/
unsigned int chunks_finished_;
/**
* \brief work_packages_ holds all unit of work.
*/
Vector<WorkPackage> work_packages_;
/**
* \brief denotes boundary for border compositing
* \note measured in pixel space
*/
rcti viewer_border_;
/**
* \brief start time of execution
*/
double execution_start_time_;
// methods
/**
* \brief check whether parameter operation can be added to the execution group
* \param operation: the operation to be added
*/
bool can_contain(NodeOperation &operation);
/**
* \brief Determine the rect (minx, maxx, miny, maxy) of a chunk at a position.
*/
void determine_chunk_rect(rcti *r_rect, unsigned int x_chunk, unsigned int y_chunk) const;
/**
* \brief determine the number of chunks, based on the chunk_size, width and height.
* \note The result are stored in the fields number_of_chunks, number_of_xchunks,
* number_of_ychunks
*/
void init_number_of_chunks();
/**
* \brief try to schedule a specific chunk.
* \note scheduling succeeds when all input requirements are met and the chunks hasn't been
* scheduled yet.
* \param graph:
* \param x_chunk:
* \param y_chunk:
* \return [true:false]
* true: package(s) are scheduled
* false: scheduling is deferred (depending workpackages are scheduled)
*/
bool schedule_chunk_when_possible(ExecutionSystem *graph, int chunk_x, int chunk_y);
/**
* \brief try to schedule a specific area.
* \note Check if a certain area is available, when not available this are will be checked.
* \note This method is called from other ExecutionGroup's.
* \param graph:
* \param area:
* \return [true:false]
* true: package(s) are scheduled
* false: scheduling is deferred (depending workpackages are scheduled)
*/
bool schedule_area_when_possible(ExecutionSystem *graph, rcti *area);
/**
* \brief add a chunk to the WorkScheduler.
* \param chunknumber:
*/
bool schedule_chunk(unsigned int chunk_number);
/**
* \brief determine the area of interest of a certain input area
* \note This method only evaluates a single ReadBufferOperation
* \param input: the input area
* \param read_operation: The ReadBufferOperation where the area needs to be evaluated
* \param output: the area needed of the ReadBufferOperation. Result
*/
void determine_depending_area_of_interest(rcti *input,
ReadBufferOperation *read_operation,
rcti *output);
/**
* Return the execution order of the user visible chunks.
*/
blender::Array<unsigned int> get_execution_order() const;
void init_read_buffer_operations();
void init_work_packages();
public:
// constructors
ExecutionGroup(int id);
int get_id() const
{
return id_;
}
const ExecutionGroupFlags get_flags() const
{
return flags_;
}
// methods
/**
* \brief add an operation to this ExecutionGroup
* \note this method will add input of the operations recursively
* \note this method can create multiple ExecutionGroup's
* \param system:
* \param operation:
* \return True if the operation was successfully added
*/
bool add_operation(NodeOperation *operation);
/**
* \brief set whether this ExecutionGroup is an output
* \param is_output:
*/
void set_output_execution_group(bool is_output)
{
flags_.is_output = is_output;
}
/**
* \brief determine the resolution of this ExecutionGroup
* \param resolution:
*/
void determine_resolution(unsigned int resolution[2]);
/**
* \brief set the resolution of this executiongroup
* \param resolution:
*/
void set_resolution(unsigned int resolution[2])
{
width_ = resolution[0];
height_ = resolution[1];
}
/**
* \brief get the width of this execution group
*/
unsigned int get_width() const
{
return width_;
}
/**
* \brief get the height of this execution group
*/
unsigned int get_height() const
{
return height_;
}
/**
* \brief get the output operation of this ExecutionGroup
* \return NodeOperation *output operation
*/
NodeOperation *get_output_operation() const;
/**
* \brief compose multiple chunks into a single chunk
* \return `(Memorybuffer *)` consolidated chunk
*/
MemoryBuffer *construct_consolidated_memory_buffer(MemoryProxy &memory_proxy, rcti &rect);
/**
* \brief init_execution is called just before the execution of the whole graph will be done.
* \note The implementation will calculate the chunk_size of this execution group.
*/
void init_execution();
/**
* \brief get all input-buffers needed to calculate an chunk
* \note all input-buffers must be executed
* \param chunk_number: the chunk to be calculated
* \return `(MemoryBuffer **)` the input-buffers.
*/
MemoryBuffer **get_input_buffers_opencl(int chunk_number);
/**
* \brief allocate the output-buffer of a chunk
* \param chunk_number: the number of the chunk in the ExecutionGroup
* \param rect: the rect of that chunk
* \see determine_chunk_rect
*/
MemoryBuffer *allocate_output_buffer(rcti &rect);
/**
* \brief after a chunk is executed the needed resources can be freed or unlocked.
* \param chunknumber:
* \param memorybuffers:
*/
void finalize_chunk_execution(int chunk_number, MemoryBuffer **memory_buffers);
/**
* \brief deinit_execution is called just after execution the whole graph.
* \note It will release all needed resources
*/
void deinit_execution();
/**
* \brief schedule an ExecutionGroup
* \note this method will return when all chunks have been calculated, or the execution has
* breaked (by user)
*
* first the order of the chunks will be determined. This is determined by finding the
* ViewerOperation and get the relevant information from it.
* - ChunkOrdering
* - CenterX
* - CenterY
*
* After determining the order of the chunks the chunks will be scheduled
*
* \see ViewerOperation
* \param graph:
*/
/**
* This method is called for the top execution groups. containing the compositor node or the
* preview node or the viewer node).
*/
void execute(ExecutionSystem *graph);
/**
* \brief Determine the rect (minx, maxx, miny, maxy) of a chunk.
*/
void determine_chunk_rect(rcti *r_rect, unsigned int chunk_number) const;
void set_chunksize(int chunksize)
{
chunk_size_ = chunksize;
}
/**
* \brief get the Render priority of this ExecutionGroup
* \see ExecutionSystem.execute
*/
eCompositorPriority get_render_priority();
/**
* \brief set border for viewer operation
* \note all the coordinates are assumed to be in normalized space
*/
void set_viewer_border(float xmin, float xmax, float ymin, float ymax);
void set_render_border(float xmin, float xmax, float ymin, float ymax);
/* allow the DebugInfo class to look at internals */
friend class DebugInfo;
#ifdef WITH_CXX_GUARDEDALLOC
MEM_CXX_CLASS_ALLOC_FUNCS("COM:ExecutionGroup")
#endif
};
std::ostream &operator<<(std::ostream &os, const ExecutionGroup &execution_group);
} // namespace blender::compositor

27
source/blender/compositor/intern/COM_ExecutionSystem.cc
View File

@ -5,11 +5,9 @@
#include "COM_ExecutionSystem.h"
#include "COM_Debug.h"
#include "COM_ExecutionGroup.h"
#include "COM_FullFrameExecutionModel.h"
#include "COM_NodeOperation.h"
#include "COM_NodeOperationBuilder.h"
#include "COM_TiledExecutionModel.h"
#include "COM_WorkPackage.h"
#include "COM_WorkScheduler.h"
@ -44,7 +42,6 @@ ExecutionSystem::ExecutionSystem(RenderData *rd,
context_.set_quality((eCompositorQuality)editingtree->edit_quality);
}
context_.set_rendering(rendering);
context_.setHasActiveOpenCLDevices(WorkScheduler::has_gpu_devices() && false);
context_.set_render_data(rd);
@ -56,17 +53,7 @@ ExecutionSystem::ExecutionSystem(RenderData *rd,
builder.convert_to_operations(this);
}
switch (context_.get_execution_model()) {
case eExecutionModel::Tiled:
execution_model_ = new TiledExecutionModel(context_, operations_, groups_);
break;
case eExecutionModel::FullFrame:
execution_model_ = new FullFrameExecutionModel(context_, active_buffers_, operations_);
break;
default:
BLI_assert_msg(0, "Non implemented execution model");
break;
}
execution_model_ = new FullFrameExecutionModel(context_, active_buffers_, operations_);
}
ExecutionSystem::~ExecutionSystem()
@ -80,23 +67,16 @@ ExecutionSystem::~ExecutionSystem()
delete operation;
}
operations_.clear();
for (ExecutionGroup *group : groups_) {
delete group;
}
groups_.clear();
}
void ExecutionSystem::set_operations(const Span<NodeOperation *> operations,
const Span<ExecutionGroup *> groups)
void ExecutionSystem::set_operations(const Span<NodeOperation *> operations)
{
operations_ = operations;
groups_ = groups;
}
void ExecutionSystem::execute()
{
DebugInfo::execute_started(this);
DebugInfo::execute_started();
for (NodeOperation *op : operations_) {
op->init_data();
}
@ -131,7 +111,6 @@ void ExecutionSystem::execute_work(const rcti &work_rect,
}
WorkPackage &sub_work = sub_works[i];
sub_work.type = eWorkPackageType::CustomFunction;
sub_work.execute_fn = [=, &work_func, &work_rect]() {
if (is_breaked()) {
return;

45
source/blender/compositor/intern/COM_ExecutionSystem.h
View File

@ -82,36 +82,9 @@ class ProfilerData;
*
* \see COM_convert_data_type Datatype conversions
* \see Converter.convert_resolution Image size conversions
*
* \section EM_Step4 Step4: group operations in executions groups
* ExecutionGroup are groups of operations that are calculated as being one bigger operation.
* All operations will be part of an ExecutionGroup.
* Complex nodes will be added to separate groups. Between ExecutionGroup's the data will be stored
* in MemoryBuffers. ReadBufferOperations and WriteBufferOperations are added where needed.
*
* <pre>
*
* +------------------------------+ +----------------+
* | ExecutionGroup A | |ExecutionGroup B| ExecutionGroup
* | +----------+ +----------+| |+----------+ |
* /----->| Operation|---->| Operation|-\ /--->| Operation|-\ | NodeOperation
* | | | A | | B ||| | || C | | |
* | | | cFFA | /->| cFFA ||| | || cFFA | | |
* | | +----------+ | +----------+|| | |+----------+ | |
* | +---------------|--------------+v | +-------------v--+
* +-*----+ +---*--+ +--*-*--+ +--*----+
* |inputA| |inputB| |outputA| |outputB| MemoryBuffer
* |cFAA | |cFAA | |cFAA | |cFAA |
* +------+ +------+ +-------+ +-------+
* </pre>
* \see ExecutionSystem.group_operations method doing this step
* \see ExecutionSystem.add_read_write_buffer_operations
* \see NodeOperation.is_complex
* \see ExecutionGroup class representing the ExecutionGroup
*/
/* Forward declarations. */
class ExecutionGroup;
class ExecutionModel;
class NodeOperation;
@ -136,11 +109,6 @@ class ExecutionSystem {
*/
Vector<NodeOperation *> operations_;
/**
* \brief vector of groups
*/
Vector<ExecutionGroup *> groups_;
/**
* Active execution model implementation.
*/
@ -178,13 +146,13 @@ class ExecutionSystem {
*/
~ExecutionSystem();
void set_operations(Span<NodeOperation *> operations, Span<ExecutionGroup *> groups);
void set_operations(Span<NodeOperation *> operations);
/**
* \brief execute this system
* - initialize the NodeOperation's and ExecutionGroup's
* - schedule the output ExecutionGroup's based on their priority
* - deinitialize the ExecutionGroup's and NodeOperation's
* - initialize the NodeOperation's
* - schedule the outputs based on their priority
* - deinitialize the NodeOperation's
*/
void execute();
@ -196,11 +164,6 @@ class ExecutionSystem {
return context_;
}
SharedOperationBuffers &get_active_buffers()
{
return active_buffers_;
}
/**
* Multi-threadedly execute given work function passing work_rect splits as argument.
*/

2
source/blender/compositor/intern/COM_FullFrameExecutionModel.cc
View File

@ -134,7 +134,7 @@ void FullFrameExecutionModel::render_operations()
{
const bool is_rendering = context_.is_rendering();
WorkScheduler::start(this->context_);
WorkScheduler::start();
for (eCompositorPriority priority : priorities_) {
for (NodeOperation *op : operations_) {
const bool has_size = op->get_width() > 0 && op->get_height() > 0;

60
source/blender/compositor/intern/COM_MemoryBuffer.cc
View File

@ -4,8 +4,6 @@
#include "COM_MemoryBuffer.h"
#include "COM_MemoryProxy.h"
#include "IMB_colormanagement.hh"
#include "IMB_imbuf_types.hh"
@ -30,31 +28,14 @@ static rcti create_rect(const int width, const int height)
return rect;
}
MemoryBuffer::MemoryBuffer(MemoryProxy *memory_proxy, const rcti &rect, MemoryBufferState state)
{
rect_ = rect;
is_a_single_elem_ = false;
memory_proxy_ = memory_proxy;
num_channels_ = COM_data_type_num_channels(memory_proxy->get_data_type());
buffer_ = (float *)MEM_mallocN_aligned(
sizeof(float) * buffer_len() * num_channels_, 16, "COM_MemoryBuffer");
owns_data_ = true;
state_ = state;
datatype_ = memory_proxy->get_data_type();
set_strides();
}
MemoryBuffer::MemoryBuffer(DataType data_type, const rcti &rect, bool is_a_single_elem)
{
rect_ = rect;
is_a_single_elem_ = is_a_single_elem;
memory_proxy_ = nullptr;
num_channels_ = COM_data_type_num_channels(data_type);
buffer_ = (float *)MEM_mallocN_aligned(
sizeof(float) * buffer_len() * num_channels_, 16, "COM_MemoryBuffer");
owns_data_ = true;
state_ = MemoryBufferState::Temporary;
datatype_ = data_type;
set_strides();
@ -73,19 +54,16 @@ MemoryBuffer::MemoryBuffer(float *buffer,
{
rect_ = rect;
is_a_single_elem_ = is_a_single_elem;
memory_proxy_ = nullptr;
num_channels_ = num_channels;
datatype_ = COM_num_channels_data_type(num_channels);
buffer_ = buffer;
owns_data_ = false;
state_ = MemoryBufferState::Temporary;
set_strides();
}
MemoryBuffer::MemoryBuffer(const MemoryBuffer &src) : MemoryBuffer(src.datatype_, src.rect_, false)
{
memory_proxy_ = src.memory_proxy_;
/* src may be single elem buffer */
fill_from(src);
}
@ -468,44 +446,6 @@ void MemoryBuffer::read_elem_filtered(
out);
}
/* TODO(manzanilla): to be removed with tiled implementation. */
static void read_ewa_pixel_sampled(void *userdata, int x, int y, float result[4])
{
MemoryBuffer *buffer = (MemoryBuffer *)userdata;
buffer->read(result, x, y);
}
/* TODO(manzanilla): to be removed with tiled implementation. */
void MemoryBuffer::readEWA(float *result, const float uv[2], const float derivatives[2][2])
{
if (is_a_single_elem_) {
memcpy(result, buffer_, sizeof(float) * num_channels_);
}
else {
BLI_assert(datatype_ == DataType::Color);
float inv_width = 1.0f / float(this->get_width()),
inv_height = 1.0f / float(this->get_height());
/* TODO(sergey): Render pipeline uses normalized coordinates and derivatives,
* but compositor uses pixel space. For now let's just divide the values and
* switch compositor to normalized space for EWA later.
*/
float uv_normal[2] = {uv[0] * inv_width, uv[1] * inv_height};
float du_normal[2] = {derivatives[0][0] * inv_width, derivatives[0][1] * inv_height};
float dv_normal[2] = {derivatives[1][0] * inv_width, derivatives[1][1] * inv_height};
BLI_ewa_filter(this->get_width(),
this->get_height(),
false,
true,
uv_normal,
du_normal,
dv_normal,
read_ewa_pixel_sampled,
this,
result);
}
}
void MemoryBuffer::copy_single_elem_from(const MemoryBuffer *src,
const int channel_offset,
const int elem_size,

73
source/blender/compositor/intern/COM_MemoryBuffer.h
View File

@ -21,28 +21,14 @@ struct ImBuf;
namespace blender::compositor {
/**
* \brief state of a memory buffer
* \ingroup Memory
*/
enum class MemoryBufferState {
/** \brief memory has been allocated on creator device and CPU machine,
* but kernel has not been executed */
Default = 0,
/** \brief chunk is consolidated from other chunks. special state. */
Temporary = 6,
};
enum class MemoryBufferExtend {
Clip,
Extend,
Repeat,
};
class MemoryProxy;
/**
* \brief a MemoryBuffer contains access to the data of a chunk
* \brief a MemoryBuffer contains access to the data
*/
class MemoryBuffer {
public:
@ -65,26 +51,16 @@ class MemoryBuffer {
int row_stride;
private:
/**
* \brief proxy of the memory (same for all chunks in the same buffer)
*/
MemoryProxy *memory_proxy_;
/**
* \brief the type of buffer DataType::Value, DataType::Vector, DataType::Color
*/
DataType datatype_;
/**
* \brief region of this buffer inside relative to the MemoryProxy
* \brief region of this buffer inside
*/
rcti rect_;
/**
* \brief state of the buffer
*/
MemoryBufferState state_;
/**
* \brief the actual float buffer/data
*/
@ -113,11 +89,6 @@ class MemoryBuffer {
int to_positive_y_stride_;
public:
/**
* \brief construct new temporarily MemoryBuffer for an area
*/
MemoryBuffer(MemoryProxy *memory_proxy, const rcti &rect, MemoryBufferState state);
/**
* \brief construct new temporarily MemoryBuffer for an area
*/
@ -384,12 +355,6 @@ class MemoryBuffer {
return buffer_;
}
float *release_ownership_buffer()
{
owns_data_ = false;
return buffer_;
}
/**
* Converts a single elem buffer to a full size buffer (allocates memory for all
* elements in resolution).
@ -491,8 +456,6 @@ class MemoryBuffer {
y = y + rect_.ymin;
}
/* TODO(manzanilla): to be removed with tiled implementation. For applying #MemoryBufferExtend
* use #wrap_pixel. */
inline void read(float *result,
int x,
int y,
@ -514,28 +477,6 @@ class MemoryBuffer {
memcpy(result, buffer, sizeof(float) * num_channels_);
}
}
/* TODO(manzanilla): to be removed with tiled implementation. */
inline void read_no_check(float *result,
int x,
int y,
MemoryBufferExtend extend_x = MemoryBufferExtend::Clip,
MemoryBufferExtend extend_y = MemoryBufferExtend::Clip)
{
int u = x;
int v = y;
this->wrap_pixel(u, v, extend_x, extend_y);
const int offset = get_coords_offset(u, v);
BLI_assert(offset >= 0);
BLI_assert(offset < this->buffer_len() * num_channels_);
BLI_assert(!(extend_x == MemoryBufferExtend::Clip && (u < rect_.xmin || u >= rect_.xmax)) &&
!(extend_y == MemoryBufferExtend::Clip && (v < rect_.ymin || v >= rect_.ymax)));
float *buffer = &buffer_[offset];
memcpy(result, buffer, sizeof(float) * num_channels_);
}
void write_pixel(int x, int y, const float color[4]);
void add_pixel(int x, int y, const float color[4]);
inline void read_bilinear(float *result,
@ -569,16 +510,6 @@ class MemoryBuffer {
}
}
void readEWA(float *result, const float uv[2], const float derivatives[2][2]);
/**
* \brief is this MemoryBuffer a temporarily buffer (based on an area, not on a chunk)
*/
inline bool is_temporarily() const
{
return state_ == MemoryBufferState::Temporary;
}
/**
* \brief Apply a color processor on the given area.
*/

37
source/blender/compositor/intern/COM_MemoryProxy.cc
View File

@ -1,37 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#include "COM_MemoryProxy.h"
#include "COM_MemoryBuffer.h"
namespace blender::compositor {
MemoryProxy::MemoryProxy(DataType datatype)
{
write_buffer_operation_ = nullptr;
executor_ = nullptr;
buffer_ = nullptr;
datatype_ = datatype;
}
void MemoryProxy::allocate(uint width, uint height)
{
rcti result;
result.xmin = 0;
result.xmax = width;
result.ymin = 0;
result.ymax = height;
buffer_ = new MemoryBuffer(this, result, MemoryBufferState::Default);
}
void MemoryProxy::free()
{
if (buffer_) {
delete buffer_;
buffer_ = nullptr;
}
}
} // namespace blender::compositor

114
source/blender/compositor/intern/COM_MemoryProxy.h
View File

@ -1,114 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#pragma once
#ifdef WITH_CXX_GUARDEDALLOC
# include "MEM_guardedalloc.h"
#endif
#include "COM_defines.h"
namespace blender::compositor {
/* Forward declarations. */
class MemoryBuffer;
class ExecutionGroup;
class WriteBufferOperation;
/**
* \brief A MemoryProxy is a unique identifier for a memory buffer.
* A single MemoryProxy is used among all chunks of the same buffer,
* the MemoryBuffer only stores the data of a single chunk.
* \ingroup Memory
*/
class MemoryProxy {
private:
/**
* \brief reference to the output operation of the executiongroup
*/
WriteBufferOperation *write_buffer_operation_;
/**
* \brief reference to the executor. the Execution group that can fill a chunk
*/
ExecutionGroup *executor_;
/**
* \brief the allocated memory
*/
MemoryBuffer *buffer_;
/**
* \brief datatype of this MemoryProxy
*/
DataType datatype_;
public:
MemoryProxy(DataType type);
/**
* \brief set the ExecutionGroup that can be scheduled to calculate a certain chunk.
* \param group: the ExecutionGroup to set
*/
void set_executor(ExecutionGroup *executor)
{
executor_ = executor;
}
/**
* \brief get the ExecutionGroup that can be scheduled to calculate a certain chunk.
*/
ExecutionGroup *get_executor() const
{
return executor_;
}
/**
* \brief set the WriteBufferOperation that is responsible for writing to this MemoryProxy
* \param operation:
*/
void set_write_buffer_operation(WriteBufferOperation *operation)
{
write_buffer_operation_ = operation;
}
/**
* \brief get the WriteBufferOperation that is responsible for writing to this MemoryProxy
* \return WriteBufferOperation
*/
WriteBufferOperation *get_write_buffer_operation() const
{
return write_buffer_operation_;
}
/**
* \brief allocate memory of size width x height
*/
void allocate(unsigned int width, unsigned int height);
/**
* \brief free the allocated memory
*/
void free();
/**
* \brief get the allocated memory
*/
inline MemoryBuffer *get_buffer()
{
return buffer_;
}
inline DataType get_data_type()
{
return datatype_;
}
#ifdef WITH_CXX_GUARDEDALLOC
MEM_CXX_CLASS_ALLOC_FUNCS("COM:MemoryProxy")
#endif
};
} // namespace blender::compositor

1
source/blender/compositor/intern/COM_MultiThreadedOperation.cc
View File

@ -11,7 +11,6 @@ MultiThreadedOperation::MultiThreadedOperation()
{
num_passes_ = 1;
current_pass_ = 0;
flags_.is_fullframe_operation = true;
}
void MultiThreadedOperation::update_memory_buffer(MemoryBuffer *output,

14
source/blender/compositor/intern/COM_NodeGraph.cc
View File

@ -231,7 +231,7 @@ void NodeGraph::add_proxies_group_inputs(bNode *b_node, bNode *b_node_io)
}
}
void NodeGraph::add_proxies_group_outputs(const CompositorContext &context,
void NodeGraph::add_proxies_group_outputs(const CompositorContext & /*context*/,
bNode *b_node,
bNode *b_node_io)
{
@ -247,16 +247,8 @@ void NodeGraph::add_proxies_group_outputs(const CompositorContext &context,
{
bNodeSocket *b_sock_group = find_b_node_output(b_node, b_sock_io->identifier);
if (b_sock_group) {
if (context.is_groupnode_buffer_enabled() &&
context.get_execution_model() == eExecutionModel::Tiled)
{
SocketBufferNode *buffer = new SocketBufferNode(b_node_io, b_sock_io, b_sock_group);
add_node(buffer, b_group_tree, key, is_active_group);
}
else {
SocketProxyNode *proxy = new SocketProxyNode(b_node_io, b_sock_io, b_sock_group, true);
add_node(proxy, b_group_tree, key, is_active_group);
}
SocketProxyNode *proxy = new SocketProxyNode(b_node_io, b_sock_io, b_sock_group, true);
add_node(proxy, b_group_tree, key, is_active_group);
}
}
}

189
source/blender/compositor/intern/COM_NodeOperation.cc
View File

@ -4,9 +4,7 @@
#include <cstdio>
#include "COM_BufferOperation.h"
#include "COM_ExecutionSystem.h"
#include "COM_ReadBufferOperation.h"
#include "COM_NodeOperation.h" /* own include */
@ -156,26 +154,6 @@ void NodeOperation::init_execution()
/* pass */
}
void NodeOperation::init_mutex()
{
BLI_mutex_init(&mutex_);
}
void NodeOperation::lock_mutex()
{
BLI_mutex_lock(&mutex_);
}
void NodeOperation::unlock_mutex()
{
BLI_mutex_unlock(&mutex_);
}
void NodeOperation::deinit_mutex()
{
BLI_mutex_end(&mutex_);
}
void NodeOperation::deinit_execution()
{
/* pass */
@ -213,57 +191,15 @@ NodeOperation *NodeOperation::get_input_operation(int index)
return nullptr;
}
bool NodeOperation::determine_depending_area_of_interest(rcti *input,
ReadBufferOperation *read_operation,
rcti *output)
{
if (inputs_.is_empty()) {
BLI_rcti_init(output, input->xmin, input->xmax, input->ymin, input->ymax);
return false;
}
rcti temp_output;
bool first = true;
for (int i = 0; i < get_number_of_input_sockets(); i++) {
NodeOperation *input_operation = this->get_input_operation(i);
if (input_operation &&
input_operation->determine_depending_area_of_interest(input, read_operation, &temp_output))
{
if (first) {
output->xmin = temp_output.xmin;
output->ymin = temp_output.ymin;
output->xmax = temp_output.xmax;
output->ymax = temp_output.ymax;
first = false;
}
else {
output->xmin = std::min(output->xmin, temp_output.xmin);
output->ymin = std::min(output->ymin, temp_output.ymin);
output->xmax = std::max(output->xmax, temp_output.xmax);
output->ymax = std::max(output->ymax, temp_output.ymax);
}
}
}
return !first;
}
/* -------------------------------------------------------------------- */
/** \name Full Frame Methods
* \{ */
void NodeOperation::get_area_of_interest(const int input_idx,
void NodeOperation::get_area_of_interest(const int /*input_idx*/,
const rcti &output_area,
rcti &r_input_area)
{
if (get_flags().is_fullframe_operation) {
r_input_area = output_area;
}
else {
/* Non full-frame operations never implement this method. To ensure correctness assume
* whole area is used. */
NodeOperation *input_op = get_input_operation(input_idx);
r_input_area = input_op->get_canvas();
}
r_input_area = output_area;
}
void NodeOperation::get_area_of_interest(NodeOperation *input_op,
@ -283,12 +219,7 @@ void NodeOperation::render(MemoryBuffer *output_buf,
Span<rcti> areas,
Span<MemoryBuffer *> inputs_bufs)
{
if (get_flags().is_fullframe_operation) {
render_full_frame(output_buf, areas, inputs_bufs);
}
else {
render_full_frame_fallback(output_buf, areas, inputs_bufs);
}
render_full_frame(output_buf, areas, inputs_bufs);
}
void NodeOperation::render_full_frame(MemoryBuffer *output_buf,
@ -302,92 +233,6 @@ void NodeOperation::render_full_frame(MemoryBuffer *output_buf,
deinit_execution();
}
void NodeOperation::render_full_frame_fallback(MemoryBuffer *output_buf,
Span<rcti> areas,
Span<MemoryBuffer *> inputs_bufs)
{
Vector<NodeOperationOutput *> orig_input_links = replace_inputs_with_buffers(inputs_bufs);
init_execution();
const bool is_output_operation = get_number_of_output_sockets() == 0;
if (!is_output_operation && output_buf->is_a_single_elem()) {
float *output_elem = output_buf->get_elem(0, 0);
read_sampled(output_elem, 0, 0, PixelSampler::Nearest);
}
else {
for (const rcti &rect : areas) {
exec_system_->execute_work(rect, [=](const rcti &split_rect) {
rcti tile_rect = split_rect;
if (is_output_operation) {
execute_region(&tile_rect, 0);
}
else {
render_tile(output_buf, &tile_rect);
}
});
}
}
deinit_execution();
remove_buffers_and_restore_original_inputs(orig_input_links);
}
void NodeOperation::render_tile(MemoryBuffer *output_buf, rcti *tile_rect)
{
const bool is_complex = get_flags().complex;
void *tile_data = is_complex ? initialize_tile_data(tile_rect) : nullptr;
const int elem_stride = output_buf->elem_stride;
for (int y = tile_rect->ymin; y < tile_rect->ymax; y++) {
float *output_elem = output_buf->get_elem(tile_rect->xmin, y);
if (is_complex) {
for (int x = tile_rect->xmin; x < tile_rect->xmax; x++) {
read(output_elem, x, y, tile_data);
output_elem += elem_stride;
}
}
else {
for (int x = tile_rect->xmin; x < tile_rect->xmax; x++) {
read_sampled(output_elem, x, y, PixelSampler::Nearest);
output_elem += elem_stride;
}
}
}
if (tile_data) {
deinitialize_tile_data(tile_rect, tile_data);
}
}
Vector<NodeOperationOutput *> NodeOperation::replace_inputs_with_buffers(
Span<MemoryBuffer *> inputs_bufs)
{
BLI_assert(inputs_bufs.size() == get_number_of_input_sockets());
Vector<NodeOperationOutput *> orig_links(inputs_bufs.size());
for (int i = 0; i < inputs_bufs.size(); i++) {
NodeOperationInput *input_socket = get_input_socket(i);
BufferOperation *buffer_op = new BufferOperation(inputs_bufs[i],
input_socket->get_data_type());
orig_links[i] = input_socket->get_link();
input_socket->set_link(buffer_op->get_output_socket());
buffer_op->init_execution();
}
return orig_links;
}
void NodeOperation::remove_buffers_and_restore_original_inputs(
Span<NodeOperationOutput *> original_inputs_links)
{
BLI_assert(original_inputs_links.size() == get_number_of_input_sockets());
for (int i = 0; i < original_inputs_links.size(); i++) {
NodeOperation *buffer_op = get_input_operation(i);
BLI_assert(buffer_op != nullptr);
BLI_assert(typeid(*buffer_op) == typeid(BufferOperation));
buffer_op->deinit_execution();
NodeOperationInput *input_socket = get_input_socket(i);
input_socket->set_link(original_inputs_links[i]);
delete buffer_op;
}
}
/** \} */
/*****************
@ -444,15 +289,6 @@ void NodeOperationOutput::determine_canvas(const rcti &preferred_area, rcti &r_a
std::ostream &operator<<(std::ostream &os, const NodeOperationFlags &node_operation_flags)
{
if (node_operation_flags.complex) {
os << "complex,";
}
if (node_operation_flags.open_cl) {
os << "open_cl,";
}
if (node_operation_flags.single_threaded) {
os << "single_threaded,";
}
if (node_operation_flags.use_render_border) {
os << "render_border,";
}
@ -465,12 +301,6 @@ std::ostream &operator<<(std::ostream &os, const NodeOperationFlags &node_operat
if (node_operation_flags.is_set_operation) {
os << "set_operation,";
}
if (node_operation_flags.is_write_buffer_operation) {
os << "write_buffer,";
}
if (node_operation_flags.is_read_buffer_operation) {
os << "read_buffer,";
}
if (node_operation_flags.is_proxy_operation) {
os << "proxy,";
}
@ -483,9 +313,6 @@ std::ostream &operator<<(std::ostream &os, const NodeOperationFlags &node_operat
if (!node_operation_flags.use_datatype_conversion) {
os << "no_conversion,";
}
if (node_operation_flags.is_fullframe_operation) {
os << "full_frame,";
}
if (node_operation_flags.is_constant_operation) {
os << "contant_operation,";
}
@ -505,16 +332,6 @@ std::ostream &operator<<(std::ostream &os, const NodeOperation &node_operation)
os << ",name=" << node_operation.get_name();
}
os << ",flags={" << flags << "}";
if (flags.is_read_buffer_operation) {
const ReadBufferOperation *read_operation = (const ReadBufferOperation *)&node_operation;
const MemoryProxy *proxy = read_operation->get_memory_proxy();
if (proxy) {
const WriteBufferOperation *write_operation = proxy->get_write_buffer_operation();
if (write_operation) {
os << ",write=" << (NodeOperation &)*write_operation;
}
}
}
os << ")";
return os;

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

@ -22,14 +22,10 @@
#include "BKE_node.hh"
#include "BKE_node_runtime.hh"
#include "clew.h"
#include "DNA_node_types.h"
namespace blender::compositor {
class OpenCLDevice;
class ReadBufferOperation;
class ExecutionSystem;
class NodeOperation;
class NodeOperationOutput;
@ -165,28 +161,6 @@ class NodeOperationOutput {
};
struct NodeOperationFlags {
/**
* Is this an complex operation.
*
* The input and output buffers of Complex operations are stored in buffers. It allows
* sequential and read/write.
*
* Complex operations are typically doing many reads to calculate the output of a single pixel.
* Mostly Filter types (Blurs, Convolution, Defocus etc) need this to be set to true.
*/
bool complex : 1;
/**
* Does this operation support OpenCL.
*/
bool open_cl : 1;
/**
* TODO: Remove this flag and #SingleThreadedOperation if tiled implementation is removed.
* Full-frame implementation doesn't need it.
*/
bool single_threaded : 1;
/**
* Does the operation needs a viewer border.
* Basically, setting border need to happen for only operations
@ -214,8 +188,6 @@ struct NodeOperationFlags {
* TODO: To be replaced by is_constant_operation flag once tiled implementation is removed.
*/
bool is_set_operation : 1;
bool is_write_buffer_operation : 1;
bool is_read_buffer_operation : 1;
bool is_proxy_operation : 1;
bool is_viewer_operation : 1;
bool is_preview_operation : 1;
@ -228,11 +200,6 @@ struct NodeOperationFlags {
*/
bool use_datatype_conversion : 1;
/**
* Has this operation fullframe implementation.
*/
bool is_fullframe_operation : 1;
/**
* Whether operation is a primitive constant operation (Color/Vector/Value).
*/
@ -246,20 +213,14 @@ struct NodeOperationFlags {
NodeOperationFlags()
{
complex = false;
single_threaded = false;
open_cl = false;
use_render_border = false;
use_viewer_border = false;
is_canvas_set = false;
is_set_operation = false;
is_read_buffer_operation = false;
is_write_buffer_operation = false;
is_proxy_operation = false;
is_viewer_operation = false;
is_preview_operation = false;
use_datatype_conversion = true;
is_fullframe_operation = false;
is_constant_operation = false;
can_be_constant = false;
}
@ -326,28 +287,12 @@ class NodeOperation {
std::function<void(rcti &canvas)> modify_determined_canvas_fn_;
/**
* \brief mutex reference for very special node initializations
* \note only use when you really know what you are doing.
* this mutex is used to share data among chunks in the same operation
* \see TonemapOperation for an example of usage
* \see NodeOperation.init_mutex initializes this mutex
* \see NodeOperation.deinit_mutex deinitializes this mutex
* \see NodeOperation.get_mutex retrieve a pointer to this mutex.
*/
ThreadMutex mutex_;
/**
* \brief reference to the editing bNodeTree, used for break and update callback
*/
const bNodeTree *btree_;
protected:
/**
* Compositor execution model.
*/
eExecutionModel execution_model_;
rcti canvas_ = COM_AREA_NONE;
/**
@ -441,11 +386,6 @@ class NodeOperation {
return false;
}
void set_execution_model(const eExecutionModel model)
{
execution_model_ = model;
}
void set_bnodetree(const bNodeTree *tree)
{
btree_ = tree;
@ -464,58 +404,6 @@ class NodeOperation {
virtual void init_execution();
/**
* \brief when a chunk is executed by a CPUDevice, this method is called
* \ingroup execution
* \param rect: the rectangle of the chunk (location and size)
* \param chunk_number: the chunk_number to be calculated
* \param memory_buffers: all input MemoryBuffer's needed
*/
virtual void execute_region(rcti * /*rect*/, unsigned int /*chunk_number*/) {}
/**
* \brief when a chunk is executed by an OpenCLDevice, this method is called
* \ingroup execution
* \note this method is only implemented in WriteBufferOperation
* \param context: the OpenCL context
* \param program: the OpenCL program containing all compositor kernels
* \param queue: the OpenCL command queue of the device the chunk is executed on
* \param rect: the rectangle of the chunk (location and size)
* \param chunk_number: the chunk_number to be calculated
* \param memory_buffers: all input MemoryBuffer's needed
* \param output_buffer: the output-buffer to write to
*/
virtual void execute_opencl_region(OpenCLDevice * /*device*/,
rcti * /*rect*/,
unsigned int /*chunk_number*/,
MemoryBuffer ** /*memory_buffers*/,
MemoryBuffer * /*output_buffer*/)
{
}
/**
* \brief custom handle to add new tasks to the OpenCL command queue
* in order to execute a chunk on an GPUDevice.
* \ingroup execution
* \param context: the OpenCL context
* \param program: the OpenCL program containing all compositor kernels
* \param queue: the OpenCL command queue of the device the chunk is executed on
* \param output_memory_buffer: the allocated memory buffer in main CPU memory
* \param cl_output_buffer: the allocated memory buffer in OpenCLDevice memory
* \param input_memory_buffers: all input MemoryBuffer's needed
* \param cl_mem_to_clean_up: all created cl_mem references must be added to this list.
* Framework will clean this after execution
* \param cl_kernels_to_clean_up: all created cl_kernel references must be added to this list.
* Framework will clean this after execution
*/
virtual void execute_opencl(OpenCLDevice * /*device*/,
MemoryBuffer * /*output_memory_buffer*/,
cl_mem /*cl_output_buffer*/,
MemoryBuffer ** /*input_memory_buffers*/,
std::list<cl_mem> * /*cl_mem_to_clean_up*/,
std::list<cl_kernel> * /*cl_kernels_to_clean_up*/)
{
}
virtual void deinit_execution();
void set_canvas(const rcti &canvas_area);
@ -538,10 +426,6 @@ class NodeOperation {
return false;
}
virtual bool determine_depending_area_of_interest(rcti *input,
ReadBufferOperation *read_operation,
rcti *output);
/**
* \brief set the index of the input socket that will determine the canvas of this
* operation \param index: the index to set
@ -589,36 +473,6 @@ class NodeOperation {
return BLI_rcti_size_y(&get_canvas());
}
inline void read_sampled(float result[4], float x, float y, PixelSampler sampler)
{
execute_pixel_sampled(result, x, y, sampler);
}
inline void read_filtered(float result[4], float x, float y, float dx[2], float dy[2])
{
execute_pixel_filtered(result, x, y, dx, dy);
}
inline void read(float result[4], int x, int y, void *chunk_data)
{
execute_pixel(result, x, y, chunk_data);
}
inline void read_clamped(float result[4], int x, int y, void *chunk_data)
{
execute_pixel(result,
math::clamp(x, 0, int(this->get_width()) - 1),
math::clamp(y, 0, int(this->get_height()) - 1),
chunk_data);
}
virtual void *initialize_tile_data(rcti * /*rect*/)
{
return 0;
}
virtual void deinitialize_tile_data(rcti * /*rect*/, void * /*data*/) {}
virtual MemoryBuffer *get_input_memory_buffer(MemoryBuffer ** /*memory_buffers*/)
{
return 0;
@ -704,85 +558,9 @@ class NodeOperation {
void add_input_socket(DataType datatype, ResizeMode resize_mode = ResizeMode::Center);
void add_output_socket(DataType datatype);
/* TODO(manzanilla): to be removed with tiled implementation. */
void set_width(unsigned int width)
{
canvas_.xmax = canvas_.xmin + width;
flags_.is_canvas_set = true;
}
void set_height(unsigned int height)
{
canvas_.ymax = canvas_.ymin + height;
flags_.is_canvas_set = true;
}
SocketReader *get_input_socket_reader(unsigned int index);
void deinit_mutex();
void init_mutex();
void lock_mutex();
void unlock_mutex();
/**
* \brief set whether this operation is complex
*
* Complex operations are typically doing many reads to calculate the output of a single pixel.
* Mostly Filter types (Blurs, Convolution, Defocus etc) need this to be set to true.
*/
void set_complex(bool complex)
{
flags_.complex = complex;
}
/**
* \brief calculate a single pixel
* \note this method is called for non-complex
* \param result: is a float[4] array to store the result
* \param x: the x-coordinate of the pixel to calculate in image space
* \param y: the y-coordinate of the pixel to calculate in image space
* \param input_buffers: chunks that can be read by their ReadBufferOperation.
*/
virtual void execute_pixel_sampled(float /*output*/[4],
float /*x*/,
float /*y*/,
PixelSampler /*sampler*/)
{
}
/**
* \brief calculate a single pixel
* \note this method is called for complex
* \param result: is a float[4] array to store the result
* \param x: the x-coordinate of the pixel to calculate in image space
* \param y: the y-coordinate of the pixel to calculate in image space
* \param input_buffers: chunks that can be read by their ReadBufferOperation.
* \param chunk_data: chunk specific data a during execution time.
*/
virtual void execute_pixel(float output[4], int x, int y, void * /*chunk_data*/)
{
execute_pixel_sampled(output, x, y, PixelSampler::Nearest);
}
/**
* \brief calculate a single pixel using an EWA filter
* \note this method is called for complex
* \param result: is a float[4] array to store the result
* \param x: the x-coordinate of the pixel to calculate in image space
* \param y: the y-coordinate of the pixel to calculate in image space
* \param dx:
* \param dy:
* \param input_buffers: chunks that can be read by their ReadBufferOperation.
*/
virtual void execute_pixel_filtered(
float /*output*/[4], float /*x*/, float /*y*/, float /*dx*/[2], float /*dy*/[2])
{
}
private:
/* -------------------------------------------------------------------- */
/** \name Full Frame Methods
* \{ */
/**
* Renders given areas using operations full frame implementation.
*/
@ -790,22 +568,6 @@ class NodeOperation {
Span<rcti> areas,
Span<MemoryBuffer *> inputs_bufs);
/**
* Renders given areas using operations tiled implementation.
*/
void render_full_frame_fallback(MemoryBuffer *output_buf,
Span<rcti> areas,
Span<MemoryBuffer *> inputs);
void render_tile(MemoryBuffer *output_buf, rcti *tile_rect);
/**
* \return Replaced inputs links.
*/
Vector<NodeOperationOutput *> replace_inputs_with_buffers(Span<MemoryBuffer *> inputs_bufs);
void remove_buffers_and_restore_original_inputs(
Span<NodeOperationOutput *> original_inputs_links);
/** \} */
/* allow the DebugInfo class to look at internals */
friend class DebugInfo;

221
source/blender/compositor/intern/COM_NodeOperationBuilder.cc
View File

@ -11,14 +11,11 @@
#include "COM_Converter.h"
#include "COM_Debug.h"
#include "COM_ExecutionGroup.h"
#include "COM_PreviewOperation.h"
#include "COM_ReadBufferOperation.h"
#include "COM_SetColorOperation.h"
#include "COM_SetValueOperation.h"
#include "COM_SetVectorOperation.h"
#include "COM_ViewerOperation.h"
#include "COM_WriteBufferOperation.h"
#include "COM_ConstantFolder.h"
#include "COM_NodeOperationBuilder.h" /* own include */
@ -84,22 +81,15 @@ void NodeOperationBuilder::convert_to_operations(ExecutionSystem *system)
add_datatype_conversions();
if (context_->get_execution_model() == eExecutionModel::FullFrame) {
save_graphviz("compositor_prior_folding");
ConstantFolder folder(*this);
folder.fold_operations();
}
save_graphviz("compositor_prior_folding");
ConstantFolder folder(*this);
folder.fold_operations();
determine_canvases();
save_graphviz("compositor_prior_merging");
merge_equal_operations();
if (context_->get_execution_model() == eExecutionModel::Tiled) {
/* surround complex ops with read/write buffer */
add_complex_operation_buffers();
}
/* links not available from here on */
/* XXX make links_ a local variable to avoid confusion! */
links_.clear();
@ -109,13 +99,8 @@ void NodeOperationBuilder::convert_to_operations(ExecutionSystem *system)
/* ensure topological (link-based) order of nodes */
// sort_operations(); /* not needed yet. */
if (context_->get_execution_model() == eExecutionModel::Tiled) {
/* create execution groups */
group_operations();
}
/* transfer resulting operations to the system */
system->set_operations(operations_, groups_);
system->set_operations(operations_);
}
void NodeOperationBuilder::add_operation(NodeOperation *operation)
@ -126,7 +111,6 @@ void NodeOperationBuilder::add_operation(NodeOperation *operation)
operation->set_name(current_node_->get_bnode()->name);
operation->set_node_instance_key(current_node_->get_instance_key());
}
operation->set_execution_model(context_->get_execution_model());
operation->set_execution_system(exec_system_);
}
@ -517,133 +501,6 @@ Vector<NodeOperationInput *> NodeOperationBuilder::cache_output_links(
return inputs;
}
WriteBufferOperation *NodeOperationBuilder::find_attached_write_buffer_operation(
NodeOperationOutput *output) const
{
for (const Link &link : links_) {
if (link.from() == output) {
NodeOperation &op = link.to()->get_operation();
if (op.get_flags().is_write_buffer_operation) {
return (WriteBufferOperation *)(&op);
}
}
}
return nullptr;
}
void NodeOperationBuilder::add_input_buffers(NodeOperation * /*operation*/,
NodeOperationInput *input)
{
if (!input->is_connected()) {
return;
}
NodeOperationOutput *output = input->get_link();
if (output->get_operation().get_flags().is_read_buffer_operation) {
/* input is already buffered, no need to add another */
return;
}
/* this link will be replaced below */
remove_input_link(input);
/* check of other end already has write operation, otherwise add a new one */
WriteBufferOperation *writeoperation = find_attached_write_buffer_operation(output);
if (!writeoperation) {
writeoperation = new WriteBufferOperation(output->get_data_type());
writeoperation->set_bnodetree(context_->get_bnodetree());
add_operation(writeoperation);
add_link(output, writeoperation->get_input_socket(0));
writeoperation->read_resolution_from_input_socket();
}
/* add readbuffer op for the input */
ReadBufferOperation *readoperation = new ReadBufferOperation(output->get_data_type());
readoperation->set_memory_proxy(writeoperation->get_memory_proxy());
this->add_operation(readoperation);
add_link(readoperation->get_output_socket(), input);
readoperation->read_resolution_from_write_buffer();
}
void NodeOperationBuilder::add_output_buffers(NodeOperation *operation,
NodeOperationOutput *output)
{
/* cache connected sockets, so we can safely remove links first before replacing them */
Vector<NodeOperationInput *> targets = cache_output_links(output);
if (targets.is_empty()) {
return;
}
WriteBufferOperation *write_operation = nullptr;
for (NodeOperationInput *target : targets) {
/* try to find existing write buffer operation */
if (target->get_operation().get_flags().is_write_buffer_operation) {
BLI_assert(write_operation == nullptr); /* there should only be one write op connected */
write_operation = (WriteBufferOperation *)&target->get_operation();
}
else {
/* remove all links to other nodes */
remove_input_link(target);
}
}
/* if no write buffer operation exists yet, create a new one */
if (!write_operation) {
write_operation = new WriteBufferOperation(operation->get_output_socket()->get_data_type());
write_operation->set_bnodetree(context_->get_bnodetree());
add_operation(write_operation);
add_link(output, write_operation->get_input_socket(0));
}
write_operation->read_resolution_from_input_socket();
/* add readbuffer op for every former connected input */
for (NodeOperationInput *target : targets) {
if (&target->get_operation() == write_operation) {
continue; /* skip existing write op links */
}
ReadBufferOperation *readoperation = new ReadBufferOperation(
operation->get_output_socket()->get_data_type());
readoperation->set_memory_proxy(write_operation->get_memory_proxy());
add_operation(readoperation);
add_link(readoperation->get_output_socket(), target);
readoperation->read_resolution_from_write_buffer();
}
}
void NodeOperationBuilder::add_complex_operation_buffers()
{
/* NOTE: complex ops and get cached here first, since adding operations
* will invalidate iterators over the main operations_
*/
Vector<NodeOperation *> complex_ops;
for (NodeOperation *operation : operations_) {
if (operation->get_flags().complex) {
complex_ops.append(operation);
}
}
for (NodeOperation *op : complex_ops) {
DebugInfo::operation_read_write_buffer(op);
for (int index = 0; index < op->get_number_of_input_sockets(); index++) {
add_input_buffers(op, op->get_input_socket(index));
}
for (int index = 0; index < op->get_number_of_output_sockets(); index++) {
add_output_buffers(op, op->get_output_socket(index));
}
}
}
using Tags = std::set<NodeOperation *>;
static void find_reachable_operations_recursive(Tags &reachable, NodeOperation *op)
@ -659,13 +516,6 @@ static void find_reachable_operations_recursive(Tags &reachable, NodeOperation *
find_reachable_operations_recursive(reachable, &input->get_link()->get_operation());
}
}
/* associated write-buffer operations are executed as well */
if (op->get_flags().is_read_buffer_operation) {
ReadBufferOperation *read_op = (ReadBufferOperation *)op;
MemoryProxy *memproxy = read_op->get_memory_proxy();
find_reachable_operations_recursive(reachable, memproxy->get_write_buffer_operation());
}
}
void NodeOperationBuilder::prune_operations()
@ -725,62 +575,10 @@ void NodeOperationBuilder::sort_operations()
operations_ = sorted;
}
static void add_group_operations_recursive(Tags &visited, NodeOperation *op, ExecutionGroup *group)
{
if (visited.find(op) != visited.end()) {
return;
}
visited.insert(op);
if (!group->add_operation(op)) {
return;
}
/* add all eligible input ops to the group */
for (int i = 0; i < op->get_number_of_input_sockets(); i++) {
NodeOperationInput *input = op->get_input_socket(i);
if (input->is_connected()) {
add_group_operations_recursive(visited, &input->get_link()->get_operation(), group);
}
}
}
ExecutionGroup *NodeOperationBuilder::make_group(NodeOperation *op)
{
ExecutionGroup *group = new ExecutionGroup(groups_.size());
groups_.append(group);
Tags visited;
add_group_operations_recursive(visited, op, group);
return group;
}
void NodeOperationBuilder::group_operations()
{
for (NodeOperation *op : operations_) {
if (op->is_output_operation(context_->is_rendering())) {
ExecutionGroup *group = make_group(op);
group->set_output_execution_group(true);
}
/* add new groups for associated memory proxies where needed */
if (op->get_flags().is_read_buffer_operation) {
ReadBufferOperation *read_op = (ReadBufferOperation *)op;
MemoryProxy *memproxy = read_op->get_memory_proxy();
if (memproxy->get_executor() == nullptr) {
ExecutionGroup *group = make_group(memproxy->get_write_buffer_operation());
memproxy->set_executor(group);
}
}
}
}
void NodeOperationBuilder::save_graphviz(StringRefNull name)
{
if (COM_EXPORT_GRAPHVIZ) {
exec_system_->set_operations(operations_, groups_);
exec_system_->set_operations(operations_);
DebugInfo::graphviz(exec_system_, name);
}
}
@ -800,15 +598,6 @@ std::ostream &operator<<(std::ostream &os, const NodeOperationBuilder &builder)
os << " op" << link.from()->get_operation().get_id() << " -> op"
<< link.to()->get_operation().get_id() << ";\n";
}
for (const NodeOperation *operation : builder.get_operations()) {
if (operation->get_flags().is_read_buffer_operation) {
const ReadBufferOperation &read_operation = static_cast<const ReadBufferOperation &>(
*operation);
const WriteBufferOperation &write_operation =
*read_operation.get_memory_proxy()->get_write_buffer_operation();
os << " op" << write_operation.get_id() << " -> op" << read_operation.get_id() << ";\n";
}
}
os << "}\n";
os << "# Builder end\n";

13
source/blender/compositor/intern/COM_NodeOperationBuilder.h
View File

@ -18,13 +18,11 @@ class NodeInput;
class NodeOutput;
class ExecutionSystem;
class ExecutionGroup;
class NodeOperation;
class NodeOperationInput;
class NodeOperationOutput;
class PreviewOperation;
class WriteBufferOperation;
class ViewerOperation;
class ConstantOperation;
@ -55,7 +53,6 @@ class NodeOperationBuilder {
Vector<NodeOperation *> operations_;
Vector<Link> links_;
Vector<ExecutionGroup *> groups_;
/** Maps operation inputs to node inputs */
Map<NodeOperationInput *, NodeInput *> input_map_;
@ -133,12 +130,6 @@ class NodeOperationBuilder {
/** Helper function to store connected inputs for replacement */
Vector<NodeOperationInput *> cache_output_links(NodeOperationOutput *output) const;
/** Find a connected write buffer operation to an OpOutput */
WriteBufferOperation *find_attached_write_buffer_operation(NodeOperationOutput *output) const;
/** Add read/write buffer operations around complex operations */
void add_complex_operation_buffers();
void add_input_buffers(NodeOperation *operation, NodeOperationInput *input);
void add_output_buffers(NodeOperation *operation, NodeOperationOutput *output);
/** Remove unreachable operations */
void prune_operations();
@ -146,10 +137,6 @@ class NodeOperationBuilder {
/** Sort operations by link dependencies */
void sort_operations();
/** Create execution groups */
void group_operations();
ExecutionGroup *make_group(NodeOperation *op);
private:
PreviewOperation *make_preview_operation() const;
void unlink_inputs_and_relink_outputs(NodeOperation *unlinked_op, NodeOperation *linked_op);

271
source/blender/compositor/intern/COM_OpenCLDevice.cc
View File

@ -1,271 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#include "COM_OpenCLDevice.h"
#include "COM_ExecutionGroup.h"
#include "COM_ReadBufferOperation.h"
namespace blender::compositor {
enum COM_VendorID { NVIDIA = 0x10DE, AMD = 0x1002 };
const cl_image_format IMAGE_FORMAT_COLOR = {
CL_RGBA,
CL_FLOAT,
};
const cl_image_format IMAGE_FORMAT_VECTOR = {
CL_RGB,
CL_FLOAT,
};
const cl_image_format IMAGE_FORMAT_VALUE = {
CL_R,
CL_FLOAT,
};
OpenCLDevice::OpenCLDevice(cl_context context,
cl_device_id device,
cl_program program,
cl_int vendor_id)
{
device_ = device;
context_ = context;
program_ = program;
queue_ = nullptr;
vendor_id_ = vendor_id;
cl_int error;
queue_ = clCreateCommandQueue(context_, device_, 0, &error);
}
OpenCLDevice::OpenCLDevice(OpenCLDevice &&other) noexcept
: context_(other.context_),
device_(other.device_),
program_(other.program_),
queue_(other.queue_),
vendor_id_(other.vendor_id_)
{
other.queue_ = nullptr;
}
OpenCLDevice::~OpenCLDevice()
{
if (queue_) {
clReleaseCommandQueue(queue_);
}
}
void OpenCLDevice::execute(WorkPackage *work_package)
{
const uint chunk_number = work_package->chunk_number;
ExecutionGroup *execution_group = work_package->execution_group;
MemoryBuffer **input_buffers = execution_group->get_input_buffers_opencl(chunk_number);
MemoryBuffer *output_buffer = execution_group->allocate_output_buffer(work_package->rect);
execution_group->get_output_operation()->execute_opencl_region(
this, &work_package->rect, chunk_number, input_buffers, output_buffer);
delete output_buffer;
execution_group->finalize_chunk_execution(chunk_number, input_buffers);
}
cl_mem OpenCLDevice::COM_cl_attach_memory_buffer_to_kernel_parameter(
cl_kernel kernel,
int parameter_index,
int offset_index,
std::list<cl_mem> *cleanup,
MemoryBuffer **input_memory_buffers,
SocketReader *reader)
{
return COM_cl_attach_memory_buffer_to_kernel_parameter(kernel,
parameter_index,
offset_index,
cleanup,
input_memory_buffers,
(ReadBufferOperation *)reader);
}
const cl_image_format *OpenCLDevice::determine_image_format(MemoryBuffer *memory_buffer)
{
switch (memory_buffer->get_num_channels()) {
case 1:
return &IMAGE_FORMAT_VALUE;
break;
case 3:
return &IMAGE_FORMAT_VECTOR;
break;
case 4:
return &IMAGE_FORMAT_COLOR;
break;
default:
BLI_assert_msg(0, "Unsupported num_channels.");
}
return &IMAGE_FORMAT_COLOR;
}
cl_mem OpenCLDevice::COM_cl_attach_memory_buffer_to_kernel_parameter(
cl_kernel kernel,
int parameter_index,
int offset_index,
std::list<cl_mem> *cleanup,
MemoryBuffer **input_memory_buffers,
ReadBufferOperation *reader)
{
cl_int error;
MemoryBuffer *result = reader->get_input_memory_buffer(input_memory_buffers);
const cl_image_format *image_format = determine_image_format(result);
cl_mem cl_buffer = clCreateImage2D(context_,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
image_format,
result->get_width(),
result->get_height(),
0,
result->get_buffer(),
&error);
if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
if (error == CL_SUCCESS) {
cleanup->push_back(cl_buffer);
}
error = clSetKernelArg(kernel, parameter_index, sizeof(cl_mem), &cl_buffer);
if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
COM_cl_attach_memory_buffer_offset_to_kernel_parameter(kernel, offset_index, result);
return cl_buffer;
}
void OpenCLDevice::COM_cl_attach_memory_buffer_offset_to_kernel_parameter(
cl_kernel kernel, int offset_index, MemoryBuffer *memory_buffer)
{
if (offset_index != -1) {
cl_int error;
const rcti &rect = memory_buffer->get_rect();
cl_int2 offset = {{rect.xmin, rect.ymin}};
error = clSetKernelArg(kernel, offset_index, sizeof(cl_int2), &offset);
if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
}
}
void OpenCLDevice::COM_cl_attach_size_to_kernel_parameter(cl_kernel kernel,
int offset_index,
NodeOperation *operation)
{
if (offset_index != -1) {
cl_int error;
cl_int2 offset = {{(cl_int)operation->get_width(), (cl_int)operation->get_height()}};
error = clSetKernelArg(kernel, offset_index, sizeof(cl_int2), &offset);
if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
}
}
void OpenCLDevice::COM_cl_attach_output_memory_buffer_to_kernel_parameter(
cl_kernel kernel, int parameter_index, cl_mem cl_output_memory_buffer)
{
cl_int error;
error = clSetKernelArg(kernel, parameter_index, sizeof(cl_mem), &cl_output_memory_buffer);
if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
}
void OpenCLDevice::COM_cl_enqueue_range(cl_kernel kernel, MemoryBuffer *output_memory_buffer)
{
cl_int error;
const size_t size[] = {
size_t(output_memory_buffer->get_width()),
size_t(output_memory_buffer->get_height()),
};
error = clEnqueueNDRangeKernel(queue_, kernel, 2, nullptr, size, nullptr, 0, nullptr, nullptr);
if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
}
void OpenCLDevice::COM_cl_enqueue_range(cl_kernel kernel,
MemoryBuffer *output_memory_buffer,
int offset_index,
NodeOperation *operation)
{
cl_int error;
const int width = output_memory_buffer->get_width();
const int height = output_memory_buffer->get_height();
int offsetx;
int offsety;
int local_size = 1024;
size_t size[2];
cl_int2 offset;
if (vendor_id_ == NVIDIA) {
local_size = 32;
}
bool breaked = false;
for (offsety = 0; offsety < height && (!breaked); offsety += local_size) {
offset.s[1] = offsety;
if (offsety + local_size < height) {
size[1] = local_size;
}
else {
size[1] = height - offsety;
}
for (offsetx = 0; offsetx < width && (!breaked); offsetx += local_size) {
if (offsetx + local_size < width) {
size[0] = local_size;
}
else {
size[0] = width - offsetx;
}
offset.s[0] = offsetx;
error = clSetKernelArg(kernel, offset_index, sizeof(cl_int2), &offset);
if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
error = clEnqueueNDRangeKernel(
queue_, kernel, 2, nullptr, size, nullptr, 0, nullptr, nullptr);
if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
clFlush(queue_);
if (operation->is_braked()) {
breaked = false;
}
}
}
}
cl_kernel OpenCLDevice::COM_cl_create_kernel(const char *kernelname,
std::list<cl_kernel> *cl_kernels_to_clean_up)
{
cl_int error;
cl_kernel kernel = clCreateKernel(program_, kernelname, &error);
if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
else {
if (cl_kernels_to_clean_up) {
cl_kernels_to_clean_up->push_back(kernel);
}
}
return kernel;
}
} // namespace blender::compositor

120
source/blender/compositor/intern/COM_OpenCLDevice.h
View File

@ -1,120 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
class OpenCLDevice;
#pragma once
#include <list>
#include "COM_Device.h"
#include "clew.h"
namespace blender::compositor {
class NodeOperation;
class MemoryBuffer;
class ReadBufferOperation;
typedef NodeOperation SocketReader;
/**
* \brief device representing an GPU OpenCL device.
* an instance of this class represents a single cl_device
*/
class OpenCLDevice : public Device {
private:
/**
* \brief OPENCL context
*/
cl_context context_;
/**
* \brief OPENCL device
*/
cl_device_id device_;
/**
* \brief OPENCL program
*/
cl_program program_;
/**
* \brief OPENCL command queue
*/
cl_command_queue queue_;
/**
* \brief OPENCL vendor ID
*/
cl_int vendor_id_;
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, cl_int vendor_id);
OpenCLDevice(OpenCLDevice &&other) noexcept;
~OpenCLDevice();
/**
* \brief execute a WorkPackage
* \param work: the WorkPackage to execute
*/
void execute(WorkPackage *work) override;
/**
* \brief determine an image format
* \param memorybuffer:
*/
static const cl_image_format *determine_image_format(MemoryBuffer *memory_buffer);
cl_context get_context()
{
return context_;
}
cl_command_queue get_queue()
{
return queue_;
}
cl_mem COM_cl_attach_memory_buffer_to_kernel_parameter(cl_kernel kernel,
int parameter_index,
int offset_index,
std::list<cl_mem> *cleanup,
MemoryBuffer **input_memory_buffers,
SocketReader *reader);
cl_mem COM_cl_attach_memory_buffer_to_kernel_parameter(cl_kernel kernel,
int parameter_index,
int offset_index,
std::list<cl_mem> *cleanup,
MemoryBuffer **input_memory_buffers,
ReadBufferOperation *reader);
void COM_cl_attach_memory_buffer_offset_to_kernel_parameter(cl_kernel kernel,
int offset_index,
MemoryBuffer *memory_buffers);
void COM_cl_attach_output_memory_buffer_to_kernel_parameter(cl_kernel kernel,
int parameter_index,
cl_mem cl_output_memory_buffer);
void COM_cl_attach_size_to_kernel_parameter(cl_kernel kernel,
int offset_index,
NodeOperation *operation);
void COM_cl_enqueue_range(cl_kernel kernel, MemoryBuffer *output_memory_buffer);
void COM_cl_enqueue_range(cl_kernel kernel,
MemoryBuffer *output_memory_buffer,
int offset_index,
NodeOperation *operation);
cl_kernel COM_cl_create_kernel(const char *kernelname,
std::list<cl_kernel> *cl_kernels_to_clean_up);
};
} // namespace blender::compositor

49
source/blender/compositor/intern/COM_SingleThreadedOperation.cc
View File

@ -1,49 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#include "COM_SingleThreadedOperation.h"
namespace blender::compositor {
SingleThreadedOperation::SingleThreadedOperation()
{
cached_instance_ = nullptr;
flags_.complex = true;
flags_.single_threaded = true;
}
void SingleThreadedOperation::init_execution()
{
init_mutex();
}
void SingleThreadedOperation::execute_pixel(float output[4], int x, int y, void * /*data*/)
{
cached_instance_->read_no_check(output, x, y);
}
void SingleThreadedOperation::deinit_execution()
{
deinit_mutex();
if (cached_instance_) {
delete cached_instance_;
cached_instance_ = nullptr;
}
}
void *SingleThreadedOperation::initialize_tile_data(rcti *rect)
{
if (cached_instance_) {
return cached_instance_;
}
lock_mutex();
if (cached_instance_ == nullptr) {
//
cached_instance_ = create_memory_buffer(rect);
}
unlock_mutex();
return cached_instance_;
}
} // namespace blender::compositor

44
source/blender/compositor/intern/COM_SingleThreadedOperation.h
View File

@ -1,44 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#pragma once
#include "COM_NodeOperation.h"
namespace blender::compositor {
class SingleThreadedOperation : public NodeOperation {
private:
MemoryBuffer *cached_instance_;
protected:
inline bool is_cached()
{
return cached_instance_ != nullptr;
}
public:
SingleThreadedOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel(float output[4], int x, int y, void *data) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
void *initialize_tile_data(rcti *rect) override;
virtual MemoryBuffer *create_memory_buffer(rcti *rect) = 0;
};
} // namespace blender::compositor

148
source/blender/compositor/intern/COM_TiledExecutionModel.cc
View File

@ -1,148 +0,0 @@
/* SPDX-FileCopyrightText: 2021 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#include "COM_TiledExecutionModel.h"
#include "COM_Debug.h"
#include "COM_ExecutionGroup.h"
#include "COM_ReadBufferOperation.h"
#include "COM_WorkScheduler.h"
#include "BLT_translation.hh"
#ifdef WITH_CXX_GUARDEDALLOC
# include "MEM_guardedalloc.h"
#endif
namespace blender::compositor {
TiledExecutionModel::TiledExecutionModel(CompositorContext &context,
Span<NodeOperation *> operations,
Span<ExecutionGroup *> groups)
: ExecutionModel(context, operations), groups_(groups)
{
const bNodeTree *node_tree = context.get_bnodetree();
node_tree->runtime->stats_draw(node_tree->runtime->sdh,
RPT_("Compositing | Determining resolution"));
uint resolution[2];
for (ExecutionGroup *group : groups_) {
resolution[0] = 0;
resolution[1] = 0;
group->determine_resolution(resolution);
if (border_.use_render_border) {
const rctf *render_border = border_.render_border;
group->set_render_border(
render_border->xmin, render_border->xmax, render_border->ymin, render_border->ymax);
}
if (border_.use_viewer_border) {
const rctf *viewer_border = border_.viewer_border;
group->set_viewer_border(
viewer_border->xmin, viewer_border->xmax, viewer_border->ymin, viewer_border->ymax);
}
}
}
static void update_read_buffer_offset(Span<NodeOperation *> operations)
{
uint order = 0;
for (NodeOperation *operation : operations) {
if (operation->get_flags().is_read_buffer_operation) {
ReadBufferOperation *read_operation = (ReadBufferOperation *)operation;
read_operation->set_offset(order);
order++;
}
}
}
static void init_write_operations_for_execution(Span<NodeOperation *> operations,
const bNodeTree *bTree)
{
for (NodeOperation *operation : operations) {
if (operation->get_flags().is_write_buffer_operation) {
operation->set_bnodetree(bTree);
operation->init_execution();
}
}
}
static void link_write_buffers(Span<NodeOperation *> operations)
{
for (NodeOperation *operation : operations) {
if (operation->get_flags().is_read_buffer_operation) {
ReadBufferOperation *read_operation = static_cast<ReadBufferOperation *>(operation);
read_operation->update_memory_buffer();
}
}
}
static void init_non_write_operations_for_execution(Span<NodeOperation *> operations,
const bNodeTree *bTree)
{
for (NodeOperation *operation : operations) {
if (!operation->get_flags().is_write_buffer_operation) {
operation->set_bnodetree(bTree);
operation->init_execution();
}
}
}
static void init_execution_groups_for_execution(Span<ExecutionGroup *> groups,
const int chunk_size)
{
for (ExecutionGroup *execution_group : groups) {
execution_group->set_chunksize(chunk_size);
execution_group->init_execution();
}
}
void TiledExecutionModel::execute(ExecutionSystem &exec_system)
{
const bNodeTree *editingtree = this->context_.get_bnodetree();
editingtree->runtime->stats_draw(editingtree->runtime->sdh,
RPT_("Compositing | Initializing execution"));
update_read_buffer_offset(operations_);
init_write_operations_for_execution(operations_, context_.get_bnodetree());
link_write_buffers(operations_);
init_non_write_operations_for_execution(operations_, context_.get_bnodetree());
init_execution_groups_for_execution(groups_, context_.get_chunksize());
WorkScheduler::start(context_);
execute_groups(eCompositorPriority::High, exec_system);
if (!context_.is_fast_calculation()) {
execute_groups(eCompositorPriority::Medium, exec_system);
execute_groups(eCompositorPriority::Low, exec_system);
}
WorkScheduler::finish();
WorkScheduler::stop();
editingtree->runtime->stats_draw(editingtree->runtime->sdh,
RPT_("Compositing | De-initializing execution"));
for (NodeOperation *operation : operations_) {
operation->deinit_execution();
}
for (ExecutionGroup *execution_group : groups_) {
execution_group->deinit_execution();
}
}
void TiledExecutionModel::execute_groups(eCompositorPriority priority,
ExecutionSystem &exec_system)
{
for (ExecutionGroup *execution_group : groups_) {
if (execution_group->get_flags().is_output &&
execution_group->get_render_priority() == priority)
{
execution_group->execute(&exec_system);
}
}
}
} // namespace blender::compositor

41
source/blender/compositor/intern/COM_TiledExecutionModel.h
View File

@ -1,41 +0,0 @@
/* SPDX-FileCopyrightText: 2021 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#pragma once
#include "COM_Enums.h"
#include "COM_ExecutionModel.h"
#ifdef WITH_CXX_GUARDEDALLOC
# include "MEM_guardedalloc.h"
#endif
namespace blender::compositor {
class ExecutionGroup;
/**
* Operations are executed from outputs to inputs grouped in execution groups and rendered in
* tiles.
*/
class TiledExecutionModel : public ExecutionModel {
private:
Span<ExecutionGroup *> groups_;
public:
TiledExecutionModel(CompositorContext &context,
Span<NodeOperation *> operations,
Span<ExecutionGroup *> groups);
void execute(ExecutionSystem &exec_system) override;
private:
void execute_groups(eCompositorPriority priority, ExecutionSystem &exec_system);
#ifdef WITH_CXX_GUARDEDALLOC
MEM_CXX_CLASS_ALLOC_FUNCS("COM:TiledExecutionModel")
#endif
};
} // namespace blender::compositor

22
source/blender/compositor/intern/COM_WorkPackage.cc
View File

@ -1,22 +0,0 @@
/* SPDX-FileCopyrightText: 2011 Blender Authors
*
* SPDX-License-Identifier: GPL-2.0-or-later */
#include "COM_WorkPackage.h"
#include "COM_ExecutionGroup.h"
namespace blender::compositor {
std::ostream &operator<<(std::ostream &os, const WorkPackage &work_package)
{
os << "WorkPackage(execution_group=" << *work_package.execution_group;
os << ",chunk=" << work_package.chunk_number;
os << ",state=" << work_package.state;
os << ",rect=(" << work_package.rect.xmin << "," << work_package.rect.ymin << ")-("
<< work_package.rect.xmax << "," << work_package.rect.ymax << ")";
os << ")";
return os;
}
} // namespace blender::compositor

30
source/blender/compositor/intern/COM_WorkPackage.h
View File

@ -8,43 +8,17 @@
# include "MEM_guardedalloc.h"
#endif
#include "COM_Enums.h"
#include "DNA_vec_types.h"
#include <functional>
#include <ostream>
namespace blender::compositor {
/* Forward Declarations. */
class ExecutionGroup;
/**
* \brief contains data about work that can be scheduled
* \see WorkScheduler
*/
struct WorkPackage {
eWorkPackageType type;
eWorkPackageState state = eWorkPackageState::NotScheduled;
/**
* \brief execution_group with the operations-setup to be evaluated
*/
ExecutionGroup *execution_group;
/**
* \brief number of the chunk to be executed
*/
unsigned int chunk_number;
/**
* Area of the execution group that the work package calculates.
*/
rcti rect;
/**
* Custom function to execute when work package type is CustomFunction.
* Called to execute work.
*/
std::function<void()> execute_fn;
@ -58,6 +32,4 @@ struct WorkPackage {
#endif
};
std::ostream &operator<<(std::ostream &os, const WorkPackage &work_package);
} // namespace blender::compositor

266
source/blender/compositor/intern/COM_WorkScheduler.cc
View File

@ -5,13 +5,6 @@
#include "COM_WorkScheduler.h"
#include "COM_CPUDevice.h"
#include "COM_CompositorContext.h"
#include "COM_ExecutionGroup.h"
#include "COM_OpenCLDevice.h"
#include "COM_OpenCLKernels.cl.h"
#include "COM_WriteBufferOperation.h"
#include "clew.h"
#include "MEM_guardedalloc.h"
@ -42,14 +35,6 @@ constexpr ThreadingModel COM_threading_model()
return ThreadingModel::Queue;
}
/**
* Does the active threading model support opencl?
*/
constexpr bool COM_is_opencl_enabled()
{
return COM_threading_model() != ThreadingModel::SingleThreaded;
}
static ThreadLocal(CPUDevice *) g_thread_device;
static struct {
struct {
@ -69,222 +54,9 @@ static struct {
TaskPool *pool;
} task;
struct {
ThreadQueue *queue;
cl_context context;
cl_program program;
/** \brief list of all OpenCLDevices. for every OpenCL GPU device an instance of OpenCLDevice
* is created. */
Vector<OpenCLDevice> devices;
/** \brief list of all thread for every GPUDevice in cpudevices a thread exists. */
ListBase threads;
/** \brief all scheduled work for the GPU. */
bool active = false;
bool initialized = false;
} opencl;
int num_cpu_threads;
} g_work_scheduler;
/* -------------------------------------------------------------------- */
/** \name OpenCL Scheduling
* \{ */
static void CL_CALLBACK cl_context_error(const char *errinfo,
const void * /*private_info*/,
size_t /*cb*/,
void * /*user_data*/)
{
printf("OPENCL error: %s\n", errinfo);
}
static void *thread_execute_gpu(void *data)
{
Device *device = (Device *)data;
WorkPackage *work;
while ((work = (WorkPackage *)BLI_thread_queue_pop(g_work_scheduler.opencl.queue))) {
device->execute(work);
}
return nullptr;
}
static void opencl_start(const CompositorContext &context)
{
if (context.get_has_active_opencl_devices()) {
g_work_scheduler.opencl.queue = BLI_thread_queue_init();
BLI_threadpool_init(&g_work_scheduler.opencl.threads,
thread_execute_gpu,
g_work_scheduler.opencl.devices.size());
for (Device &device : g_work_scheduler.opencl.devices) {
BLI_threadpool_insert(&g_work_scheduler.opencl.threads, &device);
}
g_work_scheduler.opencl.active = true;
}
else {
g_work_scheduler.opencl.active = false;
}
}
static bool opencl_schedule(WorkPackage *package)
{
if (package->type == eWorkPackageType::Tile && package->execution_group->get_flags().open_cl &&
g_work_scheduler.opencl.active)
{
BLI_thread_queue_push(g_work_scheduler.opencl.queue, package);
return true;
}
return false;
}
static void opencl_finish()
{
if (g_work_scheduler.opencl.active) {
BLI_thread_queue_wait_finish(g_work_scheduler.opencl.queue);
}
}
static void opencl_stop()
{
if (g_work_scheduler.opencl.active) {
BLI_thread_queue_nowait(g_work_scheduler.opencl.queue);
BLI_threadpool_end(&g_work_scheduler.opencl.threads);
BLI_thread_queue_free(g_work_scheduler.opencl.queue);
g_work_scheduler.opencl.queue = nullptr;
}
}
static bool opencl_has_gpu_devices()
{
return !g_work_scheduler.opencl.devices.is_empty();
}
static void opencl_initialize(const bool use_opencl)
{
/* deinitialize OpenCL GPU's */
if (use_opencl && !g_work_scheduler.opencl.initialized) {
g_work_scheduler.opencl.context = nullptr;
g_work_scheduler.opencl.program = nullptr;
/* This will check for errors and skip if already initialized. */
if (clewInit() != CLEW_SUCCESS) {
return;
}
if (clCreateContextFromType) {
cl_uint number_of_platforms = 0;
cl_int error;
error = clGetPlatformIDs(0, nullptr, &number_of_platforms);
if (error == -1001) {
} /* GPU not supported */
else if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
if (G.f & G_DEBUG) {
printf("%u number of platforms\n", number_of_platforms);
}
cl_platform_id *platforms = (cl_platform_id *)MEM_mallocN(
sizeof(cl_platform_id) * number_of_platforms, __func__);
error = clGetPlatformIDs(number_of_platforms, platforms, nullptr);
uint index_platform;
for (index_platform = 0; index_platform < number_of_platforms; index_platform++) {
cl_platform_id platform = platforms[index_platform];
cl_uint number_of_devices = 0;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, nullptr, &number_of_devices);
if (number_of_devices <= 0) {
continue;
}
cl_device_id *cldevices = (cl_device_id *)MEM_mallocN(
sizeof(cl_device_id) * number_of_devices, __func__);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, number_of_devices, cldevices, nullptr);
g_work_scheduler.opencl.context = clCreateContext(
nullptr, number_of_devices, cldevices, cl_context_error, nullptr, &error);
if (error != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
const char *cl_str[2] = {datatoc_COM_OpenCLKernels_cl, nullptr};
g_work_scheduler.opencl.program = clCreateProgramWithSource(
g_work_scheduler.opencl.context, 1, cl_str, nullptr, &error);
error = clBuildProgram(g_work_scheduler.opencl.program,
number_of_devices,
cldevices,
nullptr,
nullptr,
nullptr);
if (error != CL_SUCCESS) {
cl_int error2;
size_t ret_val_size = 0;
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
error2 = clGetProgramBuildInfo(g_work_scheduler.opencl.program,
cldevices[0],
CL_PROGRAM_BUILD_LOG,
0,
nullptr,
&ret_val_size);
if (error2 != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
char *build_log = (char *)MEM_mallocN(sizeof(char) * ret_val_size + 1, __func__);
error2 = clGetProgramBuildInfo(g_work_scheduler.opencl.program,
cldevices[0],
CL_PROGRAM_BUILD_LOG,
ret_val_size,
build_log,
nullptr);
if (error2 != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error, clewErrorString(error));
}
build_log[ret_val_size] = '\0';
printf("%s", build_log);
MEM_freeN(build_log);
}
else {
uint index_devices;
for (index_devices = 0; index_devices < number_of_devices; index_devices++) {
cl_device_id device = cldevices[index_devices];
cl_int vendorID = 0;
cl_int error2 = clGetDeviceInfo(
device, CL_DEVICE_VENDOR_ID, sizeof(cl_int), &vendorID, nullptr);
if (error2 != CL_SUCCESS) {
printf("CLERROR[%d]: %s\n", error2, clewErrorString(error2));
}
g_work_scheduler.opencl.devices.append_as(g_work_scheduler.opencl.context,
device,
g_work_scheduler.opencl.program,
vendorID);
}
}
MEM_freeN(cldevices);
}
MEM_freeN(platforms);
}
g_work_scheduler.opencl.initialized = true;
}
}
static void opencl_deinitialize()
{
g_work_scheduler.opencl.devices.clear_and_shrink();
if (g_work_scheduler.opencl.program) {
clReleaseProgram(g_work_scheduler.opencl.program);
g_work_scheduler.opencl.program = nullptr;
}
if (g_work_scheduler.opencl.context) {
clReleaseContext(g_work_scheduler.opencl.context);
g_work_scheduler.opencl.context = nullptr;
}
g_work_scheduler.opencl.initialized = false;
}
/** \} */
/* -------------------------------------------------------------------- */
/** \name Single threaded Scheduling
* \{ */
@ -419,12 +191,6 @@ static void threading_model_task_stop()
void WorkScheduler::schedule(WorkPackage *package)
{
if (COM_is_opencl_enabled()) {
if (opencl_schedule(package)) {
return;
}
}
switch (COM_threading_model()) {
case ThreadingModel::SingleThreaded: {
threading_model_single_thread_execute(package);
@ -443,12 +209,8 @@ void WorkScheduler::schedule(WorkPackage *package)
}
}
void WorkScheduler::start(const CompositorContext &context)
void WorkScheduler::start()
{
if (COM_is_opencl_enabled()) {
opencl_start(context);
}
switch (COM_threading_model()) {
case ThreadingModel::SingleThreaded:
/* Nothing to do. */
@ -466,10 +228,6 @@ void WorkScheduler::start(const CompositorContext &context)
void WorkScheduler::finish()
{
if (COM_is_opencl_enabled()) {
opencl_finish();
}
switch (COM_threading_model()) {
case ThreadingModel::SingleThreaded:
/* Nothing to do. */
@ -487,10 +245,6 @@ void WorkScheduler::finish()
void WorkScheduler::stop()
{
if (COM_is_opencl_enabled()) {
opencl_stop();
}
switch (COM_threading_model()) {
case ThreadingModel::SingleThreaded:
/* Nothing to do. */
@ -506,20 +260,8 @@ void WorkScheduler::stop()
}
}
bool WorkScheduler::has_gpu_devices()
void WorkScheduler::initialize(int num_cpu_threads)
{
if (COM_is_opencl_enabled()) {
return opencl_has_gpu_devices();
}
return false;
}
void WorkScheduler::initialize(bool use_opencl, int num_cpu_threads)
{
if (COM_is_opencl_enabled()) {
opencl_initialize(use_opencl);
}
g_work_scheduler.num_cpu_threads = num_cpu_threads;
switch (COM_threading_model()) {
case ThreadingModel::SingleThreaded:
@ -538,10 +280,6 @@ void WorkScheduler::initialize(bool use_opencl, int num_cpu_threads)
void WorkScheduler::deinitialize()
{
if (COM_is_opencl_enabled()) {
opencl_deinitialize();
}
switch (COM_threading_model()) {
case ThreadingModel::SingleThreaded:
/* Nothing to do. */

23
source/blender/compositor/intern/COM_WorkScheduler.h
View File

@ -12,8 +12,6 @@ namespace blender::compositor {
struct WorkPackage;
class CompositorContext;
/** \brief the workscheduler
* \ingroup execution
*/
@ -21,9 +19,6 @@ struct WorkScheduler {
/**
* \brief schedule a chunk of a group to be calculated.
* An execution group schedules a chunk in the WorkScheduler
* when ExecutionGroup.get_flags().open_cl is set the work will be handled by a OpenCLDevice
* otherwise the work is scheduled for an CPUDevice
* \see ExecutionGroup.execute
*/
static void schedule(WorkPackage *package);
@ -33,13 +28,9 @@ struct WorkScheduler {
* during initialization the mutexes are initialized.
* there are two mutexes (for every device type one)
* After mutex initialization the system is queried in order to count the number of CPUDevices
* and GPUDevices to be created. For every hardware thread a CPUDevice and for every OpenCL GPU
* device a OpenCLDevice is created. these devices are stored in a separate list (cpudevices &
* gpudevices)
*
* This function can be called multiple times to lazily initialize OpenCL.
* to be created. For every hardware thread a CPUDevice is created.
*/
static void initialize(bool use_opencl, int num_cpu_threads);
static void initialize(int num_cpu_threads);
/**
* \brief deinitialize the WorkScheduler
@ -53,7 +44,7 @@ struct WorkScheduler {
* for every device a thread is created.
* \see initialize Initialization and query of the number of devices
*/
static void start(const CompositorContext &context);
static void start();
/**
* \brief stop the execution
@ -67,14 +58,6 @@ struct WorkScheduler {
*/
static void finish();
/**
* \brief Are there OpenCL capable GPU devices initialized?
* the result of this method is stored in the CompositorContext
* A node can generate a different operation tree when OpenCLDevices exists.
* \see CompositorContext.get_has_active_opencl_devices
*/
static bool has_gpu_devices();
static int get_num_cpu_threads();
static int current_thread_id();

2
source/blender/compositor/intern/COM_compositor.cc
View File

@ -89,7 +89,7 @@ void COM_execute(Render *render,
/* CPU compositor. */
/* Initialize workscheduler. */
blender::compositor::WorkScheduler::initialize(false, BKE_render_num_threads(render_data));
blender::compositor::WorkScheduler::initialize(BKE_render_num_threads(render_data));
/* Execute. */
const bool twopass = (node_tree->flag & NTREE_TWO_PASS) && !rendering;

8
source/blender/compositor/nodes/COM_BlurNode.cc
View File

@ -5,11 +5,9 @@
#include "COM_BlurNode.h"
#include "COM_FastGaussianBlurOperation.h"
#include "COM_GammaCorrectOperation.h"
#include "COM_GaussianAlphaXBlurOperation.h"
#include "COM_GaussianAlphaYBlurOperation.h"
#include "COM_GaussianAlphaBlurBaseOperation.h"
#include "COM_GaussianBlurBaseOperation.h"
#include "COM_GaussianBokehBlurOperation.h"
#include "COM_GaussianXBlurOperation.h"
#include "COM_GaussianYBlurOperation.h"
#include "COM_MathBaseOperation.h"
#include "COM_SetValueOperation.h"
@ -93,7 +91,6 @@ void BlurNode::convert_to_operations(NodeConverter &converter,
GaussianXBlurOperation *operationx = new GaussianXBlurOperation();
operationx->set_data(data);
operationx->set_quality(quality);
operationx->check_opencl();
operationx->set_extend_bounds(extend_bounds);
converter.add_operation(operationx);
@ -102,7 +99,6 @@ void BlurNode::convert_to_operations(NodeConverter &converter,
GaussianYBlurOperation *operationy = new GaussianYBlurOperation();
operationy->set_data(data);
operationy->set_quality(quality);
operationy->check_opencl();
operationy->set_extend_bounds(extend_bounds);
converter.add_operation(operationy);

3
source/blender/compositor/nodes/COM_DilateErodeNode.cc
View File

@ -4,8 +4,7 @@
#include "COM_DilateErodeNode.h"
#include "COM_DilateErodeOperation.h"
#include "COM_GaussianAlphaXBlurOperation.h"
#include "COM_GaussianAlphaYBlurOperation.h"
#include "COM_GaussianAlphaBlurBaseOperation.h"
#include "COM_SMAAOperation.h"
namespace blender::compositor {

3
source/blender/compositor/nodes/COM_KeyingNode.cc
View File

@ -18,8 +18,7 @@
#include "COM_SetAlphaMultiplyOperation.h"
#include "COM_GaussianAlphaXBlurOperation.h"
#include "COM_GaussianAlphaYBlurOperation.h"
#include "COM_GaussianAlphaBlurBaseOperation.h"
#include "BLI_math_color.h"

3
source/blender/compositor/nodes/COM_KuwaharaNode.cc
View File

@ -9,8 +9,7 @@
#include "COM_KuwaharaNode.h"
#include "COM_GaussianXBlurOperation.h"
#include "COM_GaussianYBlurOperation.h"
#include "COM_GaussianBlurBaseOperation.h"
#include "COM_KuwaharaAnisotropicOperation.h"
#include "COM_KuwaharaAnisotropicStructureTensorOperation.h"
#include "COM_KuwaharaClassicOperation.h"

19
source/blender/compositor/nodes/COM_RotateNode.cc
View File

@ -15,7 +15,7 @@ RotateNode::RotateNode(bNode *editor_node) : Node(editor_node)
}
void RotateNode::convert_to_operations(NodeConverter &converter,
const CompositorContext &context) const
const CompositorContext & /*context*/) const
{
NodeInput *input_socket = this->get_input_socket(0);
NodeInput *input_degree_socket = this->get_input_socket(1);
@ -24,21 +24,8 @@ void RotateNode::convert_to_operations(NodeConverter &converter,
converter.add_operation(operation);
PixelSampler sampler = (PixelSampler)this->get_bnode()->custom1;
switch (context.get_execution_model()) {
case eExecutionModel::Tiled: {
SetSamplerOperation *sampler_op = new SetSamplerOperation();
sampler_op->set_sampler(sampler);
converter.add_operation(sampler_op);
converter.add_link(sampler_op->get_output_socket(), operation->get_input_socket(0));
converter.map_input_socket(input_socket, sampler_op->get_input_socket(0));
break;
}
case eExecutionModel::FullFrame: {
operation->set_sampler(sampler);
converter.map_input_socket(input_socket, operation->get_input_socket(0));
break;
}
}
operation->set_sampler(sampler);
converter.map_input_socket(input_socket, operation->get_input_socket(0));
converter.map_input_socket(input_degree_socket, operation->get_input_socket(1));
converter.map_output_socket(output_socket, operation->get_output_socket(0));

45
source/blender/compositor/nodes/COM_SocketProxyNode.cc
View File

@ -3,8 +3,6 @@
* SPDX-License-Identifier: GPL-2.0-or-later */
#include "COM_SocketProxyNode.h"
#include "COM_ReadBufferOperation.h"
#include "COM_WriteBufferOperation.h"
namespace blender::compositor {
@ -43,47 +41,4 @@ void SocketProxyNode::convert_to_operations(NodeConverter &converter,
converter.map_output_socket(get_output_socket(), proxy_output);
}
SocketBufferNode::SocketBufferNode(bNode *editor_node,
bNodeSocket *editor_input,
bNodeSocket *editor_output)
: Node(editor_node, false)
{
DataType dt;
dt = DataType::Value;
if (editor_input->type == SOCK_RGBA) {
dt = DataType::Color;
}
if (editor_input->type == SOCK_VECTOR) {
dt = DataType::Vector;
}
this->add_input_socket(dt, editor_input);
dt = DataType::Value;
if (editor_output->type == SOCK_RGBA) {
dt = DataType::Color;
}
if (editor_output->type == SOCK_VECTOR) {
dt = DataType::Vector;
}
this->add_output_socket(dt, editor_output);
}
void SocketBufferNode::convert_to_operations(NodeConverter &converter,
const CompositorContext & /*context*/) const
{
NodeOutput *output = this->get_output_socket(0);
NodeInput *input = this->get_input_socket(0);
DataType datatype = output->get_data_type();
WriteBufferOperation *write_operation = new WriteBufferOperation(datatype);
ReadBufferOperation *read_operation = new ReadBufferOperation(datatype);
read_operation->set_memory_proxy(write_operation->get_memory_proxy());
converter.add_operation(write_operation);
converter.add_operation(read_operation);
converter.map_input_socket(input, write_operation->get_input_socket(0));
converter.map_output_socket(output, read_operation->get_output_socket());
}
} // namespace blender::compositor

7
source/blender/compositor/nodes/COM_SocketProxyNode.h
View File

@ -35,11 +35,4 @@ class SocketProxyNode : public Node {
bool use_conversion_;
};
class SocketBufferNode : public Node {
public:
SocketBufferNode(bNode *editor_node, bNodeSocket *editor_input, bNodeSocket *editor_output);
void convert_to_operations(NodeConverter &converter,
const CompositorContext &context) const override;
};
} // namespace blender::compositor

146
source/blender/compositor/nodes/COM_Stabilize2dNode.cc
View File

@ -55,119 +55,55 @@ void Stabilize2dNode::convert_to_operations(NodeConverter &converter,
converter.add_operation(x_attribute);
converter.add_operation(y_attribute);
switch (context.get_execution_model()) {
case eExecutionModel::Tiled: {
ScaleRelativeOperation *scale_operation = new ScaleRelativeOperation();
scale_operation->set_sampler(sampler);
RotateOperation *rotate_operation = new RotateOperation();
rotate_operation->set_do_degree2_rad_conversion(false);
TranslateOperation *translate_operation = new TranslateOperation();
SetSamplerOperation *psoperation = new SetSamplerOperation();
psoperation->set_sampler(sampler);
ScaleRelativeOperation *scale_operation = new ScaleRelativeOperation();
scale_operation->set_sampler(sampler);
RotateOperation *rotate_operation = new RotateOperation();
rotate_operation->set_do_degree2_rad_conversion(false);
rotate_operation->set_sampler(sampler);
TranslateOperation *translate_operation = new TranslateCanvasOperation();
converter.add_operation(scale_operation);
converter.add_operation(translate_operation);
converter.add_operation(rotate_operation);
converter.add_operation(psoperation);
converter.add_operation(scale_operation);
converter.add_operation(translate_operation);
converter.add_operation(rotate_operation);
converter.add_link(scale_attribute->get_output_socket(),
scale_operation->get_input_socket(1));
converter.add_link(scale_attribute->get_output_socket(),
scale_operation->get_input_socket(2));
converter.add_link(scale_attribute->get_output_socket(), scale_operation->get_input_socket(1));
converter.add_link(scale_attribute->get_output_socket(), scale_operation->get_input_socket(2));
converter.add_link(angle_attribute->get_output_socket(),
rotate_operation->get_input_socket(1));
converter.add_link(angle_attribute->get_output_socket(), rotate_operation->get_input_socket(1));
converter.add_link(x_attribute->get_output_socket(),
translate_operation->get_input_socket(1));
converter.add_link(y_attribute->get_output_socket(),
translate_operation->get_input_socket(2));
converter.add_link(x_attribute->get_output_socket(), translate_operation->get_input_socket(1));
converter.add_link(y_attribute->get_output_socket(), translate_operation->get_input_socket(2));
converter.map_output_socket(get_output_socket(), psoperation->get_output_socket());
NodeOperationInput *stabilization_socket = nullptr;
if (invert) {
/* Translate -> Rotate -> Scale. */
stabilization_socket = translate_operation->get_input_socket(0);
converter.map_input_socket(image_input, translate_operation->get_input_socket(0));
if (invert) {
/* Translate -> Rotate -> Scale. */
converter.map_input_socket(image_input, translate_operation->get_input_socket(0));
converter.add_link(translate_operation->get_output_socket(),
rotate_operation->get_input_socket(0));
converter.add_link(rotate_operation->get_output_socket(),
scale_operation->get_input_socket(0));
converter.add_link(translate_operation->get_output_socket(),
rotate_operation->get_input_socket(0));
converter.add_link(rotate_operation->get_output_socket(),
scale_operation->get_input_socket(0));
converter.add_link(scale_operation->get_output_socket(), psoperation->get_input_socket(0));
}
else {
/* Scale -> Rotate -> Translate. */
converter.map_input_socket(image_input, scale_operation->get_input_socket(0));
converter.add_link(scale_operation->get_output_socket(),
rotate_operation->get_input_socket(0));
converter.add_link(rotate_operation->get_output_socket(),
translate_operation->get_input_socket(0));
converter.add_link(translate_operation->get_output_socket(),
psoperation->get_input_socket(0));
}
break;
}
case eExecutionModel::FullFrame: {
ScaleRelativeOperation *scale_operation = new ScaleRelativeOperation();
scale_operation->set_sampler(sampler);
RotateOperation *rotate_operation = new RotateOperation();
rotate_operation->set_do_degree2_rad_conversion(false);
rotate_operation->set_sampler(sampler);
TranslateOperation *translate_operation = new TranslateCanvasOperation();
converter.add_operation(scale_operation);
converter.add_operation(translate_operation);
converter.add_operation(rotate_operation);
converter.add_link(scale_attribute->get_output_socket(),
scale_operation->get_input_socket(1));
converter.add_link(scale_attribute->get_output_socket(),
scale_operation->get_input_socket(2));
converter.add_link(angle_attribute->get_output_socket(),
rotate_operation->get_input_socket(1));
converter.add_link(x_attribute->get_output_socket(),
translate_operation->get_input_socket(1));
converter.add_link(y_attribute->get_output_socket(),
translate_operation->get_input_socket(2));
NodeOperationInput *stabilization_socket = nullptr;
if (invert) {
/* Translate -> Rotate -> Scale. */
stabilization_socket = translate_operation->get_input_socket(0);
converter.map_input_socket(image_input, translate_operation->get_input_socket(0));
converter.add_link(translate_operation->get_output_socket(),
rotate_operation->get_input_socket(0));
converter.add_link(rotate_operation->get_output_socket(),
scale_operation->get_input_socket(0));
converter.map_output_socket(get_output_socket(), scale_operation->get_output_socket());
}
else {
/* Scale -> Rotate -> Translate. */
stabilization_socket = scale_operation->get_input_socket(0);
converter.map_input_socket(image_input, scale_operation->get_input_socket(0));
converter.add_link(scale_operation->get_output_socket(),
rotate_operation->get_input_socket(0));
converter.add_link(rotate_operation->get_output_socket(),
translate_operation->get_input_socket(0));
converter.map_output_socket(get_output_socket(), translate_operation->get_output_socket());
}
x_attribute->set_socket_input_resolution_for_stabilization(stabilization_socket);
y_attribute->set_socket_input_resolution_for_stabilization(stabilization_socket);
scale_attribute->set_socket_input_resolution_for_stabilization(stabilization_socket);
angle_attribute->set_socket_input_resolution_for_stabilization(stabilization_socket);
break;
}
converter.map_output_socket(get_output_socket(), scale_operation->get_output_socket());
}
else {
/* Scale -> Rotate -> Translate. */
stabilization_socket = scale_operation->get_input_socket(0);
converter.map_input_socket(image_input, scale_operation->get_input_socket(0));
converter.add_link(scale_operation->get_output_socket(),
rotate_operation->get_input_socket(0));
converter.add_link(rotate_operation->get_output_socket(),
translate_operation->get_input_socket(0));
converter.map_output_socket(get_output_socket(), translate_operation->get_output_socket());
}
x_attribute->set_socket_input_resolution_for_stabilization(stabilization_socket);
y_attribute->set_socket_input_resolution_for_stabilization(stabilization_socket);
scale_attribute->set_socket_input_resolution_for_stabilization(stabilization_socket);
angle_attribute->set_socket_input_resolution_for_stabilization(stabilization_socket);
}
} // namespace blender::compositor

83
source/blender/compositor/nodes/COM_TransformNode.cc
View File

@ -16,7 +16,7 @@ TransformNode::TransformNode(bNode *editor_node) : Node(editor_node)
}
void TransformNode::convert_to_operations(NodeConverter &converter,
const CompositorContext &context) const
const CompositorContext & /*context*/) const
{
NodeInput *image_input = this->get_input_socket(0);
NodeInput *x_input = this->get_input_socket(1);
@ -24,73 +24,34 @@ void TransformNode::convert_to_operations(NodeConverter &converter,
NodeInput *angle_input = this->get_input_socket(3);
NodeInput *scale_input = this->get_input_socket(4);
switch (context.get_execution_model()) {
case eExecutionModel::Tiled: {
ScaleRelativeOperation *scale_operation = new ScaleRelativeOperation();
converter.add_operation(scale_operation);
ScaleRelativeOperation *scale_operation = new ScaleRelativeOperation();
converter.add_operation(scale_operation);
RotateOperation *rotate_operation = new RotateOperation();
rotate_operation->set_do_degree2_rad_conversion(false);
converter.add_operation(rotate_operation);
RotateOperation *rotate_operation = new RotateOperation();
rotate_operation->set_do_degree2_rad_conversion(false);
converter.add_operation(rotate_operation);
TranslateOperation *translate_operation = new TranslateOperation();
converter.add_operation(translate_operation);
TranslateOperation *translate_operation = new TranslateCanvasOperation();
converter.add_operation(translate_operation);
SetSamplerOperation *sampler = new SetSamplerOperation();
sampler->set_sampler((PixelSampler)this->get_bnode()->custom1);
converter.add_operation(sampler);
PixelSampler sampler = (PixelSampler)this->get_bnode()->custom1;
scale_operation->set_sampler(sampler);
rotate_operation->set_sampler(sampler);
converter.map_input_socket(image_input, sampler->get_input_socket(0));
converter.add_link(sampler->get_output_socket(), scale_operation->get_input_socket(0));
converter.map_input_socket(scale_input, scale_operation->get_input_socket(1));
converter.map_input_socket(scale_input,
scale_operation->get_input_socket(2)); // xscale = yscale
converter.map_input_socket(image_input, scale_operation->get_input_socket(0));
converter.map_input_socket(scale_input, scale_operation->get_input_socket(1));
converter.map_input_socket(scale_input,
scale_operation->get_input_socket(2)); // xscale = yscale
converter.add_link(scale_operation->get_output_socket(),
rotate_operation->get_input_socket(0));
converter.map_input_socket(angle_input, rotate_operation->get_input_socket(1));
converter.add_link(scale_operation->get_output_socket(), rotate_operation->get_input_socket(0));
converter.map_input_socket(angle_input, rotate_operation->get_input_socket(1));
converter.add_link(rotate_operation->get_output_socket(),
translate_operation->get_input_socket(0));
converter.map_input_socket(x_input, translate_operation->get_input_socket(1));
converter.map_input_socket(y_input, translate_operation->get_input_socket(2));
converter.add_link(rotate_operation->get_output_socket(),
translate_operation->get_input_socket(0));
converter.map_input_socket(x_input, translate_operation->get_input_socket(1));
converter.map_input_socket(y_input, translate_operation->get_input_socket(2));
converter.map_output_socket(get_output_socket(), translate_operation->get_output_socket());
break;
}
case eExecutionModel::FullFrame: {
ScaleRelativeOperation *scale_operation = new ScaleRelativeOperation();
converter.add_operation(scale_operation);
RotateOperation *rotate_operation = new RotateOperation();
rotate_operation->set_do_degree2_rad_conversion(false);
converter.add_operation(rotate_operation);
TranslateOperation *translate_operation = new TranslateCanvasOperation();
converter.add_operation(translate_operation);
PixelSampler sampler = (PixelSampler)this->get_bnode()->custom1;
scale_operation->set_sampler(sampler);
rotate_operation->set_sampler(sampler);
converter.map_input_socket(image_input, scale_operation->get_input_socket(0));
converter.map_input_socket(scale_input, scale_operation->get_input_socket(1));
converter.map_input_socket(scale_input,
scale_operation->get_input_socket(2)); // xscale = yscale
converter.add_link(scale_operation->get_output_socket(),
rotate_operation->get_input_socket(0));
converter.map_input_socket(angle_input, rotate_operation->get_input_socket(1));
converter.add_link(rotate_operation->get_output_socket(),
translate_operation->get_input_socket(0));
converter.map_input_socket(x_input, translate_operation->get_input_socket(1));
converter.map_input_socket(y_input, translate_operation->get_input_socket(2));
converter.map_output_socket(get_output_socket(), translate_operation->get_output_socket());
break;
}
}
converter.map_output_socket(get_output_socket(), translate_operation->get_output_socket());
}
} // namespace blender::compositor

24
source/blender/compositor/nodes/COM_TranslateNode.cc
View File

@ -5,8 +5,6 @@
#include "COM_TranslateNode.h"
#include "COM_TranslateOperation.h"
#include "COM_WrapOperation.h"
#include "COM_WriteBufferOperation.h"
namespace blender::compositor {
@ -16,7 +14,7 @@ TranslateNode::TranslateNode(bNode *editor_node) : Node(editor_node)
}
void TranslateNode::convert_to_operations(NodeConverter &converter,
const CompositorContext &context) const
const CompositorContext & /*context*/) const
{
const bNode *bnode = this->get_bnode();
const NodeTranslateData *data = (const NodeTranslateData *)bnode->storage;
@ -26,9 +24,7 @@ void TranslateNode::convert_to_operations(NodeConverter &converter,
NodeInput *input_ysocket = this->get_input_socket(2);
NodeOutput *output_socket = this->get_output_socket(0);
TranslateOperation *operation = context.get_execution_model() == eExecutionModel::Tiled ?
new TranslateOperation() :
new TranslateCanvasOperation();
TranslateOperation *operation = new TranslateCanvasOperation();
operation->set_wrapping(data->wrap_axis);
operation->set_is_relative(data->relative);
@ -36,21 +32,7 @@ void TranslateNode::convert_to_operations(NodeConverter &converter,
converter.map_input_socket(input_xsocket, operation->get_input_socket(1));
converter.map_input_socket(input_ysocket, operation->get_input_socket(2));
converter.map_output_socket(output_socket, operation->get_output_socket(0));
if (data->wrap_axis && context.get_execution_model() != eExecutionModel::FullFrame) {
/* TODO: To be removed with tiled implementation. */
WriteBufferOperation *write_operation = new WriteBufferOperation(DataType::Color);
WrapOperation *wrap_operation = new WrapOperation(DataType::Color);
wrap_operation->set_memory_proxy(write_operation->get_memory_proxy());
wrap_operation->set_wrapping(data->wrap_axis);
converter.add_operation(write_operation);
converter.add_operation(wrap_operation);
converter.map_input_socket(input_socket, write_operation->get_input_socket(0));
converter.add_link(wrap_operation->get_output_socket(), operation->get_input_socket(0));
}
else {
converter.map_input_socket(input_socket, operation->get_input_socket(0));
}
converter.map_input_socket(input_socket, operation->get_input_socket(0));
}
} // namespace blender::compositor

1
source/blender/compositor/nodes/COM_ViewerNode.cc
View File

@ -29,7 +29,6 @@ void ViewerNode::convert_to_operations(NodeConverter &converter,
viewer_operation->set_bnodetree(context.get_bnodetree());
viewer_operation->set_image(image);
viewer_operation->set_image_user(image_user);
viewer_operation->set_chunk_order((ChunkOrdering)editor_node->custom1);
viewer_operation->setCenterX(editor_node->custom3);
viewer_operation->setCenterY(editor_node->custom4);
/* alpha socket gives either 1 or a custom alpha value if "use alpha" is enabled */

30
source/blender/compositor/operations/COM_AlphaOverKeyOperation.cc
View File

@ -11,36 +11,6 @@ AlphaOverKeyOperation::AlphaOverKeyOperation()
flags_.can_be_constant = true;
}
void AlphaOverKeyOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float input_color1[4];
float input_over_color[4];
float value[4];
input_value_operation_->read_sampled(value, x, y, sampler);
input_color1_operation_->read_sampled(input_color1, x, y, sampler);
input_color2_operation_->read_sampled(input_over_color, x, y, sampler);
if (input_over_color[3] <= 0.0f) {
copy_v4_v4(output, input_color1);
}
else if (value[0] == 1.0f && input_over_color[3] >= 1.0f) {
copy_v4_v4(output, input_over_color);
}
else {
float premul = value[0] * input_over_color[3];
float mul = 1.0f - premul;
output[0] = (mul * input_color1[0]) + premul * input_over_color[0];
output[1] = (mul * input_color1[1]) + premul * input_over_color[1];
output[2] = (mul * input_color1[2]) + premul * input_over_color[2];
output[3] = (mul * input_color1[3]) + value[0] * input_over_color[3];
}
}
void AlphaOverKeyOperation::update_memory_buffer_row(PixelCursor &p)
{
for (; p.out < p.row_end; p.next()) {

5
source/blender/compositor/operations/COM_AlphaOverKeyOperation.h
View File

@ -16,11 +16,6 @@ class AlphaOverKeyOperation : public MixBaseOperation {
public:
AlphaOverKeyOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
void update_memory_buffer_row(PixelCursor &p) override;
};

31
source/blender/compositor/operations/COM_AlphaOverMixedOperation.cc
View File

@ -12,37 +12,6 @@ AlphaOverMixedOperation::AlphaOverMixedOperation()
flags_.can_be_constant = true;
}
void AlphaOverMixedOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float input_color1[4];
float input_over_color[4];
float value[4];
input_value_operation_->read_sampled(value, x, y, sampler);
input_color1_operation_->read_sampled(input_color1, x, y, sampler);
input_color2_operation_->read_sampled(input_over_color, x, y, sampler);
if (input_over_color[3] <= 0.0f) {
copy_v4_v4(output, input_color1);
}
else if (value[0] == 1.0f && input_over_color[3] >= 1.0f) {
copy_v4_v4(output, input_over_color);
}
else {
float addfac = 1.0f - x_ + input_over_color[3] * x_;
float premul = value[0] * addfac;
float mul = 1.0f - value[0] * input_over_color[3];
output[0] = (mul * input_color1[0]) + premul * input_over_color[0];
output[1] = (mul * input_color1[1]) + premul * input_over_color[1];
output[2] = (mul * input_color1[2]) + premul * input_over_color[2];
output[3] = (mul * input_color1[3]) + value[0] * input_over_color[3];
}
}
void AlphaOverMixedOperation::update_memory_buffer_row(PixelCursor &p)
{
for (; p.out < p.row_end; p.next()) {

5
source/blender/compositor/operations/COM_AlphaOverMixedOperation.h
View File

@ -22,11 +22,6 @@ class AlphaOverMixedOperation : public MixBaseOperation {
*/
AlphaOverMixedOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
void setX(float x)
{
x_ = x;

30
source/blender/compositor/operations/COM_AlphaOverPremultiplyOperation.cc
View File

@ -11,36 +11,6 @@ AlphaOverPremultiplyOperation::AlphaOverPremultiplyOperation()
flags_.can_be_constant = true;
}
void AlphaOverPremultiplyOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float input_color1[4];
float input_over_color[4];
float value[4];
input_value_operation_->read_sampled(value, x, y, sampler);
input_color1_operation_->read_sampled(input_color1, x, y, sampler);
input_color2_operation_->read_sampled(input_over_color, x, y, sampler);
/* Zero alpha values should still permit an add of RGB data */
if (input_over_color[3] < 0.0f) {
copy_v4_v4(output, input_color1);
}
else if (value[0] == 1.0f && input_over_color[3] >= 1.0f) {
copy_v4_v4(output, input_over_color);
}
else {
float mul = 1.0f - value[0] * input_over_color[3];
output[0] = (mul * input_color1[0]) + value[0] * input_over_color[0];
output[1] = (mul * input_color1[1]) + value[0] * input_over_color[1];
output[2] = (mul * input_color1[2]) + value[0] * input_over_color[2];
output[3] = (mul * input_color1[3]) + value[0] * input_over_color[3];
}
}
void AlphaOverPremultiplyOperation::update_memory_buffer_row(PixelCursor &p)
{
for (; p.out < p.row_end; p.next()) {

5
source/blender/compositor/operations/COM_AlphaOverPremultiplyOperation.h
View File

@ -16,11 +16,6 @@ class AlphaOverPremultiplyOperation : public MixBaseOperation {
public:
AlphaOverPremultiplyOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
void update_memory_buffer_row(PixelCursor &p) override;
};

73
source/blender/compositor/operations/COM_BilateralBlurOperation.cc
View File

@ -11,87 +11,14 @@ BilateralBlurOperation::BilateralBlurOperation()
this->add_input_socket(DataType::Color);
this->add_input_socket(DataType::Color);
this->add_output_socket(DataType::Color);
flags_.complex = true;
flags_.can_be_constant = true;
input_color_program_ = nullptr;
input_determinator_program_ = nullptr;
}
void BilateralBlurOperation::init_execution()
{
input_color_program_ = get_input_socket_reader(0);
input_determinator_program_ = get_input_socket_reader(1);
QualityStepHelper::init_execution(COM_QH_INCREASE);
}
void BilateralBlurOperation::execute_pixel(float output[4], int x, int y, void *data)
{
/* Read the determinator color at x, y,
* this will be used as the reference color for the determinator. */
float determinator_reference_color[4];
float determinator[4];
float temp_color[4];
float blur_color[4];
float blur_divider;
float sigmacolor = data_->sigma_color;
float delta_color;
input_determinator_program_->read(determinator_reference_color, x, y, data);
zero_v4(blur_color);
blur_divider = 0.0f;
/* TODO(sergey): This isn't really good bilateral filter, it should be
* using gaussian bell for weights. Also sigma_color doesn't seem to be
* used correct at all.
*/
for (int yi = -radius_; yi <= radius_; yi += QualityStepHelper::get_step()) {
for (int xi = -radius_; xi <= radius_; xi += QualityStepHelper::get_step()) {
/* Read determinator. */
input_determinator_program_->read_clamped(determinator, x + xi, y + yi, data);
delta_color = (fabsf(determinator_reference_color[0] - determinator[0]) +
fabsf(determinator_reference_color[1] - determinator[1]) +
/* Do not take the alpha channel into account. */
fabsf(determinator_reference_color[2] - determinator[2]));
if (delta_color < sigmacolor) {
/* Add this to the blur. */
input_color_program_->read_clamped(temp_color, x + xi, y + yi, data);
add_v4_v4(blur_color, temp_color);
blur_divider += 1.0f;
}
}
}
if (blur_divider > 0.0f) {
mul_v4_v4fl(output, blur_color, 1.0f / blur_divider);
}
else {
output[0] = 0.0f;
output[1] = 0.0f;
output[2] = 0.0f;
output[3] = 1.0f;
}
}
void BilateralBlurOperation::deinit_execution()
{
input_color_program_ = nullptr;
input_determinator_program_ = nullptr;
}
bool BilateralBlurOperation::determine_depending_area_of_interest(
rcti *input, ReadBufferOperation *read_operation, rcti *output)
{
rcti new_input;
int add = radius_ + 1;
new_input.xmax = input->xmax + (add);
new_input.xmin = input->xmin - (add);
new_input.ymax = input->ymax + (add);
new_input.ymin = input->ymin - (add);
return NodeOperation::determine_depending_area_of_interest(&new_input, read_operation, output);
}
void BilateralBlurOperation::get_area_of_interest(const int /*input_idx*/,
const rcti &output_area,
rcti &r_input_area)

19
source/blender/compositor/operations/COM_BilateralBlurOperation.h
View File

@ -13,33 +13,14 @@ namespace blender::compositor {
class BilateralBlurOperation : public MultiThreadedOperation, public QualityStepHelper {
private:
SocketReader *input_color_program_;
SocketReader *input_determinator_program_;
NodeBilateralBlurData *data_;
int radius_;
public:
BilateralBlurOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel(float output[4], int x, int y, void *data) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
bool determine_depending_area_of_interest(rcti *input,
ReadBufferOperation *read_operation,
rcti *output) override;
void set_data(NodeBilateralBlurData *data)
{
data_ = data;

63
source/blender/compositor/operations/COM_BlurBaseOperation.cc
View File

@ -15,9 +15,7 @@ BlurBaseOperation::BlurBaseOperation(DataType data_type)
this->add_input_socket(data_type);
this->add_input_socket(DataType::Value);
this->add_output_socket(data_type);
flags_.complex = true;
flags_.can_be_constant = true;
input_program_ = nullptr;
memset(&data_, 0, sizeof(NodeBlurData));
size_ = 1.0f;
sizeavailable_ = false;
@ -27,9 +25,7 @@ BlurBaseOperation::BlurBaseOperation(DataType data_type)
void BlurBaseOperation::init_data()
{
if (execution_model_ == eExecutionModel::FullFrame) {
update_size();
}
update_size();
data_.image_in_width = this->get_width();
data_.image_in_height = this->get_height();
@ -55,9 +51,6 @@ void BlurBaseOperation::init_data()
void BlurBaseOperation::init_execution()
{
input_program_ = this->get_input_socket_reader(0);
input_size_ = this->get_input_socket_reader(1);
QualityStepHelper::init_execution(COM_QH_MULTIPLY);
}
@ -148,12 +141,6 @@ float *BlurBaseOperation::make_dist_fac_inverse(float rad, int size, int falloff
return dist_fac_invert;
}
void BlurBaseOperation::deinit_execution()
{
input_program_ = nullptr;
input_size_ = nullptr;
}
void BlurBaseOperation::set_data(const NodeBlurData *data)
{
memcpy(&data_, data, sizeof(NodeBlurData));
@ -176,21 +163,10 @@ void BlurBaseOperation::update_size()
return;
}
switch (execution_model_) {
case eExecutionModel::Tiled: {
float result[4];
this->get_input_socket_reader(1)->read_sampled(result, 0, 0, PixelSampler::Nearest);
size_ = result[0];
break;
}
case eExecutionModel::FullFrame: {
NodeOperation *size_input = get_input_operation(SIZE_INPUT_INDEX);
if (size_input->get_flags().is_constant_operation) {
size_ = *static_cast<ConstantOperation *>(size_input)->get_constant_elem();
} /* Else use default. */
break;
}
}
NodeOperation *size_input = get_input_operation(SIZE_INPUT_INDEX);
if (size_input->get_flags().is_constant_operation) {
size_ = *static_cast<ConstantOperation *>(size_input)->get_constant_elem();
} /* Else use default. */
sizeavailable_ = true;
}
@ -201,26 +177,15 @@ void BlurBaseOperation::determine_canvas(const rcti &preferred_area, rcti &r_are
return;
}
switch (execution_model_) {
case eExecutionModel::Tiled: {
NodeOperation::determine_canvas(preferred_area, r_area);
r_area.xmax += 2 * size_ * data_.sizex;
r_area.ymax += 2 * size_ * data_.sizey;
break;
}
case eExecutionModel::FullFrame: {
/* Setting a modifier ensures all non main inputs have extended bounds as preferred
* canvas, avoiding unnecessary canvas conversions that would hide constant
* operations. */
set_determined_canvas_modifier([=](rcti &canvas) {
/* Rounding to even prevents jiggling in backdrop while switching size values. */
canvas.xmax += round_to_even(2 * size_ * data_.sizex);
canvas.ymax += round_to_even(2 * size_ * data_.sizey);
});
NodeOperation::determine_canvas(preferred_area, r_area);
break;
}
}
/* Setting a modifier ensures all non main inputs have extended bounds as preferred
* canvas, avoiding unnecessary canvas conversions that would hide constant
* operations. */
set_determined_canvas_modifier([=](rcti &canvas) {
/* Rounding to even prevents jiggling in backdrop while switching size values. */
canvas.xmax += round_to_even(2 * size_ * data_.sizex);
canvas.ymax += round_to_even(2 * size_ * data_.sizey);
});
NodeOperation::determine_canvas(preferred_area, r_area);
}
void BlurBaseOperation::get_area_of_interest(const int input_idx,

13
source/blender/compositor/operations/COM_BlurBaseOperation.h
View File

@ -35,11 +35,6 @@ class BlurBaseOperation : public MultiThreadedOperation, public QualityStepHelpe
void update_size();
/**
* Cached reference to the input_program
*/
SocketReader *input_program_;
SocketReader *input_size_;
NodeBlurData data_;
float size_;
@ -50,16 +45,8 @@ class BlurBaseOperation : public MultiThreadedOperation, public QualityStepHelpe
public:
virtual void init_data() override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
void set_data(const NodeBlurData *data);
void set_size(float size)

203
source/blender/compositor/operations/COM_BokehBlurOperation.cc
View File

@ -8,8 +8,6 @@
#include "COM_BokehBlurOperation.h"
#include "COM_ConstantOperation.h"
#include "COM_OpenCLDevice.h"
namespace blender::compositor {
constexpr int IMAGE_INPUT_INDEX = 0;
@ -25,195 +23,35 @@ BokehBlurOperation::BokehBlurOperation()
this->add_input_socket(DataType::Value);
this->add_output_socket(DataType::Color);
flags_.complex = true;
flags_.open_cl = true;
flags_.can_be_constant = true;
size_ = 1.0f;
sizeavailable_ = false;
input_program_ = nullptr;
input_bokeh_program_ = nullptr;
input_bounding_box_reader_ = nullptr;
extend_bounds_ = false;
}
void BokehBlurOperation::init_data()
{
if (execution_model_ == eExecutionModel::FullFrame) {
update_size();
}
}
void *BokehBlurOperation::initialize_tile_data(rcti * /*rect*/)
{
lock_mutex();
if (!sizeavailable_) {
update_size();
}
void *buffer = get_input_operation(0)->initialize_tile_data(nullptr);
unlock_mutex();
return buffer;
update_size();
}
void BokehBlurOperation::init_execution()
{
init_mutex();
input_program_ = get_input_socket_reader(0);
input_bokeh_program_ = get_input_socket_reader(1);
input_bounding_box_reader_ = get_input_socket_reader(2);
QualityStepHelper::init_execution(COM_QH_INCREASE);
}
void BokehBlurOperation::execute_pixel(float output[4], int x, int y, void *data)
{
MemoryBuffer *input_buffer = (MemoryBuffer *)data;
float temp_bounding_box[4];
input_bounding_box_reader_->read_sampled(temp_bounding_box, x, y, PixelSampler::Nearest);
if (temp_bounding_box[0] <= 0.0f) {
copy_v4_v4(output, input_buffer->get_elem(x, y));
return;
}
const float max_dim = std::max(this->get_width(), this->get_height());
int radius = size_ * max_dim / 100.0f;
const int2 bokeh_size = int2(input_bokeh_program_->get_width(),
input_bokeh_program_->get_height());
float4 accumulated_color = float4(0.0f);
float4 accumulated_weight = float4(0.0f);
int step = get_step();
for (int yi = -radius; yi <= radius; yi += step) {
for (int xi = -radius; xi <= radius; xi += step) {
const float2 normalized_texel = (float2(xi, yi) + radius + 0.5f) / (radius * 2.0f + 1.0f);
const float2 weight_texel = (1.0f - normalized_texel) * float2(bokeh_size - 1);
float4 weight;
input_bokeh_program_->read(weight, int(weight_texel.x), int(weight_texel.y), nullptr);
const float4 color = float4(input_buffer->get_elem_clamped(x + xi, y + yi)) * weight;
accumulated_color += color;
accumulated_weight += weight;
}
}
const float4 final_color = math::safe_divide(accumulated_color, accumulated_weight);
copy_v4_v4(output, final_color);
}
void BokehBlurOperation::deinit_execution()
{
deinit_mutex();
input_program_ = nullptr;
input_bokeh_program_ = nullptr;
input_bounding_box_reader_ = nullptr;
}
bool BokehBlurOperation::determine_depending_area_of_interest(rcti *input,
ReadBufferOperation *read_operation,
rcti *output)
{
rcti new_input;
rcti bokeh_input;
const float max_dim = std::max(this->get_width(), this->get_height());
if (sizeavailable_) {
new_input.xmax = input->xmax + (size_ * max_dim / 100.0f);
new_input.xmin = input->xmin - (size_ * max_dim / 100.0f);
new_input.ymax = input->ymax + (size_ * max_dim / 100.0f);
new_input.ymin = input->ymin - (size_ * max_dim / 100.0f);
}
else {
new_input.xmax = input->xmax + (10.0f * max_dim / 100.0f);
new_input.xmin = input->xmin - (10.0f * max_dim / 100.0f);
new_input.ymax = input->ymax + (10.0f * max_dim / 100.0f);
new_input.ymin = input->ymin - (10.0f * max_dim / 100.0f);
}
NodeOperation *operation = get_input_operation(1);
bokeh_input.xmax = operation->get_width();
bokeh_input.xmin = 0;
bokeh_input.ymax = operation->get_height();
bokeh_input.ymin = 0;
if (operation->determine_depending_area_of_interest(&bokeh_input, read_operation, output)) {
return true;
}
operation = get_input_operation(0);
if (operation->determine_depending_area_of_interest(&new_input, read_operation, output)) {
return true;
}
operation = get_input_operation(2);
if (operation->determine_depending_area_of_interest(input, read_operation, output)) {
return true;
}
if (!sizeavailable_) {
rcti size_input;
size_input.xmin = 0;
size_input.ymin = 0;
size_input.xmax = 5;
size_input.ymax = 5;
operation = get_input_operation(3);
if (operation->determine_depending_area_of_interest(&size_input, read_operation, output)) {
return true;
}
}
return false;
}
void BokehBlurOperation::execute_opencl(OpenCLDevice *device,
MemoryBuffer *output_memory_buffer,
cl_mem cl_output_buffer,
MemoryBuffer **input_memory_buffers,
std::list<cl_mem> *cl_mem_to_clean_up,
std::list<cl_kernel> * /*cl_kernels_to_clean_up*/)
{
cl_kernel kernel = device->COM_cl_create_kernel("bokeh_blur_kernel", nullptr);
if (!sizeavailable_) {
update_size();
}
const float max_dim = std::max(this->get_width(), this->get_height());
cl_int radius = size_ * max_dim / 100.0f;
cl_int step = this->get_step();
device->COM_cl_attach_memory_buffer_to_kernel_parameter(
kernel, 0, -1, cl_mem_to_clean_up, input_memory_buffers, input_bounding_box_reader_);
device->COM_cl_attach_memory_buffer_to_kernel_parameter(
kernel, 1, 4, cl_mem_to_clean_up, input_memory_buffers, input_program_);
device->COM_cl_attach_memory_buffer_to_kernel_parameter(
kernel, 2, -1, cl_mem_to_clean_up, input_memory_buffers, input_bokeh_program_);
device->COM_cl_attach_output_memory_buffer_to_kernel_parameter(kernel, 3, cl_output_buffer);
device->COM_cl_attach_memory_buffer_offset_to_kernel_parameter(kernel, 5, output_memory_buffer);
clSetKernelArg(kernel, 6, sizeof(cl_int), &radius);
clSetKernelArg(kernel, 7, sizeof(cl_int), &step);
device->COM_cl_attach_size_to_kernel_parameter(kernel, 8, this);
device->COM_cl_enqueue_range(kernel, output_memory_buffer, 9, this);
}
void BokehBlurOperation::update_size()
{
if (sizeavailable_) {
return;
}
switch (execution_model_) {
case eExecutionModel::Tiled: {
float result[4];
this->get_input_socket_reader(3)->read_sampled(result, 0, 0, PixelSampler::Nearest);
size_ = result[0];
CLAMP(size_, 0.0f, 10.0f);
break;
}
case eExecutionModel::FullFrame: {
NodeOperation *size_input = get_input_operation(SIZE_INPUT_INDEX);
if (size_input->get_flags().is_constant_operation) {
size_ = *static_cast<ConstantOperation *>(size_input)->get_constant_elem();
CLAMP(size_, 0.0f, 10.0f);
} /* Else use default. */
break;
}
}
NodeOperation *size_input = get_input_operation(SIZE_INPUT_INDEX);
if (size_input->get_flags().is_constant_operation) {
size_ = *static_cast<ConstantOperation *>(size_input)->get_constant_elem();
CLAMP(size_, 0.0f, 10.0f);
} /* Else use default. */
sizeavailable_ = true;
}
@ -224,27 +62,14 @@ void BokehBlurOperation::determine_canvas(const rcti &preferred_area, rcti &r_ar
return;
}
switch (execution_model_) {
case eExecutionModel::Tiled: {
NodeOperation::determine_canvas(preferred_area, r_area);
const float max_dim = std::max(BLI_rcti_size_x(&r_area), BLI_rcti_size_y(&r_area));
float add_size = round_to_even(2 * size_ * max_dim / 100.0f);
r_area.xmax += add_size;
r_area.ymax += add_size;
break;
}
case eExecutionModel::FullFrame: {
set_determined_canvas_modifier([=](rcti &canvas) {
const float max_dim = std::max(BLI_rcti_size_x(&canvas), BLI_rcti_size_y(&canvas));
/* Rounding to even prevents image jiggling in backdrop while switching size values. */
float add_size = round_to_even(2 * size_ * max_dim / 100.0f);
canvas.xmax += add_size;
canvas.ymax += add_size;
});
NodeOperation::determine_canvas(preferred_area, r_area);
break;
}
}
set_determined_canvas_modifier([=](rcti &canvas) {
const float max_dim = std::max(BLI_rcti_size_x(&canvas), BLI_rcti_size_y(&canvas));
/* Rounding to even prevents image jiggling in backdrop while switching size values. */
float add_size = round_to_even(2 * size_ * max_dim / 100.0f);
canvas.xmax += add_size;
canvas.ymax += add_size;
});
NodeOperation::determine_canvas(preferred_area, r_area);
}
void BokehBlurOperation::get_area_of_interest(const int input_idx,

28
source/blender/compositor/operations/COM_BokehBlurOperation.h
View File

@ -11,9 +11,6 @@ namespace blender::compositor {
class BokehBlurOperation : public MultiThreadedOperation, public QualityStepHelper {
private:
SocketReader *input_program_;
SocketReader *input_bokeh_program_;
SocketReader *input_bounding_box_reader_;
void update_size();
float size_;
bool sizeavailable_;
@ -25,39 +22,14 @@ class BokehBlurOperation : public MultiThreadedOperation, public QualityStepHelp
void init_data() override;
void *initialize_tile_data(rcti *rect) override;
/**
* The inner loop of this operation.
*/
void execute_pixel(float output[4], int x, int y, void *data) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
bool determine_depending_area_of_interest(rcti *input,
ReadBufferOperation *read_operation,
rcti *output) override;
void set_size(float size)
{
size_ = size;
sizeavailable_ = true;
}
void execute_opencl(OpenCLDevice *device,
MemoryBuffer *output_memory_buffer,
cl_mem cl_output_buffer,
MemoryBuffer **input_memory_buffers,
std::list<cl_mem> *cl_mem_to_clean_up,
std::list<cl_kernel> *cl_kernels_to_clean_up) override;
void set_extend_bounds(bool extend_bounds)
{
extend_bounds_ = extend_bounds;

24
source/blender/compositor/operations/COM_BokehImageOperation.cc
View File

@ -75,30 +75,6 @@ float BokehImageOperation::is_inside_bokeh(float distance, float x, float y)
}
return inside_bokeh;
}
void BokehImageOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler /*sampler*/)
{
float shift = data_->lensshift;
float shift2 = shift / 2.0f;
float distance = circular_distance_;
float inside_bokeh_max = is_inside_bokeh(distance, x, y);
float inside_bokeh_med = is_inside_bokeh(distance - fabsf(shift2 * distance), x, y);
float inside_bokeh_min = is_inside_bokeh(distance - fabsf(shift * distance), x, y);
if (shift < 0) {
output[0] = inside_bokeh_max;
output[1] = inside_bokeh_med;
output[2] = inside_bokeh_min;
}
else {
output[0] = inside_bokeh_min;
output[1] = inside_bokeh_med;
output[2] = inside_bokeh_max;
}
output[3] = (inside_bokeh_max + inside_bokeh_med + inside_bokeh_min) / 3.0f;
}
void BokehImageOperation::update_memory_buffer_partial(MemoryBuffer *output,
const rcti &area,
Span<MemoryBuffer *> /*inputs*/)

12
source/blender/compositor/operations/COM_BokehImageOperation.h
View File

@ -95,19 +95,7 @@ class BokehImageOperation : public MultiThreadedOperation {
public:
BokehImageOperation();
/**
* \brief The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
/**
* \brief Initialize the execution
*/
void init_execution() override;
/**
* \brief De-initialize the execution
*/
void deinit_execution() override;
/**

76
source/blender/compositor/operations/COM_BoxMaskOperation.cc
View File

@ -11,87 +11,17 @@ BoxMaskOperation::BoxMaskOperation()
this->add_input_socket(DataType::Value);
this->add_input_socket(DataType::Value);
this->add_output_socket(DataType::Value);
input_mask_ = nullptr;
input_value_ = nullptr;
cosine_ = 0.0f;
sine_ = 0.0f;
}
void BoxMaskOperation::init_execution()
{
input_mask_ = this->get_input_socket_reader(0);
input_value_ = this->get_input_socket_reader(1);
const double rad = double(data_->rotation);
cosine_ = cos(rad);
sine_ = sin(rad);
aspect_ratio_ = float(this->get_width()) / this->get_height();
}
void BoxMaskOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float input_mask[4];
float input_value[4];
float rx = x / std::max(this->get_width() - 1.0f, FLT_EPSILON);
float ry = y / std::max(this->get_height() - 1.0f, FLT_EPSILON);
const float dy = (ry - data_->y) / aspect_ratio_;
const float dx = rx - data_->x;
rx = data_->x + (cosine_ * dx + sine_ * dy);
ry = data_->y + (-sine_ * dx + cosine_ * dy);
input_mask_->read_sampled(input_mask, x, y, sampler);
input_value_->read_sampled(input_value, x, y, sampler);
float half_height = data_->height / 2.0f + FLT_EPSILON;
float half_width = data_->width / 2.0f + FLT_EPSILON;
bool inside = (rx >= data_->x - half_width && rx <= data_->x + half_width &&
ry >= data_->y - half_height && ry <= data_->y + half_height);
switch (mask_type_) {
case CMP_NODE_MASKTYPE_ADD:
if (inside) {
output[0] = std::max(input_mask[0], input_value[0]);
}
else {
output[0] = input_mask[0];
}
break;
case CMP_NODE_MASKTYPE_SUBTRACT:
if (inside) {
output[0] = input_mask[0] - input_value[0];
CLAMP(output[0], 0, 1);
}
else {
output[0] = input_mask[0];
}
break;
case CMP_NODE_MASKTYPE_MULTIPLY:
if (inside) {
output[0] = input_mask[0] * input_value[0];
}
else {
output[0] = 0;
}
break;
case CMP_NODE_MASKTYPE_NOT:
if (inside) {
if (input_mask[0] > 0.0f) {
output[0] = 0;
}
else {
output[0] = input_value[0];
}
}
else {
output[0] = input_mask[0];
}
break;
}
}
void BoxMaskOperation::update_memory_buffer_partial(MemoryBuffer *output,
const rcti &area,
Span<MemoryBuffer *> inputs)
@ -150,10 +80,4 @@ void BoxMaskOperation::apply_mask(MemoryBuffer *output,
}
}
void BoxMaskOperation::deinit_execution()
{
input_mask_ = nullptr;
input_value_ = nullptr;
}
} // namespace blender::compositor

19
source/blender/compositor/operations/COM_BoxMaskOperation.h
View File

@ -12,12 +12,6 @@ class BoxMaskOperation : public MultiThreadedOperation {
private:
using MaskFunc = std::function<float(bool is_inside, const float *mask, const float *value)>;
/**
* Cached reference to the input_program
*/
SocketReader *input_mask_;
SocketReader *input_value_;
float sine_;
float cosine_;
float aspect_ratio_;
@ -28,21 +22,8 @@ class BoxMaskOperation : public MultiThreadedOperation {
public:
BoxMaskOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
void set_data(const NodeBoxMask *data)
{
data_ = data;

58
source/blender/compositor/operations/COM_BrightnessOperation.cc
View File

@ -14,7 +14,6 @@ BrightnessOperation::BrightnessOperation()
this->add_input_socket(DataType::Value);
this->add_input_socket(DataType::Value);
this->add_output_socket(DataType::Color);
input_program_ = nullptr;
use_premultiply_ = false;
flags_.can_be_constant = true;
}
@ -24,56 +23,6 @@ void BrightnessOperation::set_use_premultiply(bool use_premultiply)
use_premultiply_ = use_premultiply;
}
void BrightnessOperation::init_execution()
{
input_program_ = this->get_input_socket_reader(0);
input_brightness_program_ = this->get_input_socket_reader(1);
input_contrast_program_ = this->get_input_socket_reader(2);
}
void BrightnessOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float input_value[4];
float a, b;
float input_brightness[4];
float input_contrast[4];
input_program_->read_sampled(input_value, x, y, sampler);
input_brightness_program_->read_sampled(input_brightness, x, y, sampler);
input_contrast_program_->read_sampled(input_contrast, x, y, sampler);
float brightness = input_brightness[0];
float contrast = input_contrast[0];
brightness /= 100.0f;
float delta = contrast / 200.0f;
/*
* The algorithm is by Werner D. Streidt
* (http://visca.com/ffactory/archives/5-99/msg00021.html)
* Extracted of OpenCV `demhist.c`.
*/
if (contrast > 0) {
a = 1.0f - delta * 2.0f;
a = 1.0f / max_ff(a, FLT_EPSILON);
b = a * (brightness - delta);
}
else {
delta *= -1;
a = max_ff(1.0f - delta * 2.0f, 0.0f);
b = a * brightness + delta;
}
if (use_premultiply_) {
premul_to_straight_v4(input_value);
}
output[0] = a * input_value[0] + b;
output[1] = a * input_value[1] + b;
output[2] = a * input_value[2] + b;
output[3] = input_value[3];
if (use_premultiply_) {
straight_to_premul_v4(output);
}
}
void BrightnessOperation::update_memory_buffer_partial(MemoryBuffer *output,
const rcti &area,
Span<MemoryBuffer *> inputs)
@ -118,11 +67,4 @@ void BrightnessOperation::update_memory_buffer_partial(MemoryBuffer *output,
}
}
void BrightnessOperation::deinit_execution()
{
input_program_ = nullptr;
input_brightness_program_ = nullptr;
input_contrast_program_ = nullptr;
}
} // namespace blender::compositor

22
source/blender/compositor/operations/COM_BrightnessOperation.h
View File

@ -10,33 +10,11 @@ namespace blender::compositor {
class BrightnessOperation : public MultiThreadedOperation {
private:
/**
* Cached reference to the input_program
*/
SocketReader *input_program_;
SocketReader *input_brightness_program_;
SocketReader *input_contrast_program_;
bool use_premultiply_;
public:
BrightnessOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
void set_use_premultiply(bool use_premultiply);
void update_memory_buffer_partial(MemoryBuffer *output,

90
source/blender/compositor/operations/COM_CalculateMeanOperation.cc
View File

@ -14,10 +14,8 @@ CalculateMeanOperation::CalculateMeanOperation()
{
this->add_input_socket(DataType::Color, ResizeMode::Align);
this->add_output_socket(DataType::Value);
image_reader_ = nullptr;
is_calculated_ = false;
setting_ = 1;
flags_.complex = true;
flags_.is_constant_operation = true;
needs_canvas_to_get_constant_ = true;
@ -25,95 +23,7 @@ CalculateMeanOperation::CalculateMeanOperation()
void CalculateMeanOperation::init_execution()
{
image_reader_ = this->get_input_socket_reader(0);
is_calculated_ = false;
NodeOperation::init_mutex();
}
void CalculateMeanOperation::execute_pixel(float output[4], int /*x*/, int /*y*/, void * /*data*/)
{
output[0] = constant_value_;
}
void CalculateMeanOperation::deinit_execution()
{
image_reader_ = nullptr;
NodeOperation::deinit_mutex();
}
bool CalculateMeanOperation::determine_depending_area_of_interest(
rcti * /*input*/, ReadBufferOperation *read_operation, rcti *output)
{
rcti image_input;
if (is_calculated_) {
return false;
}
NodeOperation *operation = get_input_operation(0);
image_input.xmax = operation->get_width();
image_input.xmin = 0;
image_input.ymax = operation->get_height();
image_input.ymin = 0;
if (operation->determine_depending_area_of_interest(&image_input, read_operation, output)) {
return true;
}
return false;
}
void *CalculateMeanOperation::initialize_tile_data(rcti *rect)
{
lock_mutex();
if (!is_calculated_) {
MemoryBuffer *tile = (MemoryBuffer *)image_reader_->initialize_tile_data(rect);
constant_value_ = calculate_mean_tile(tile);
is_calculated_ = true;
}
unlock_mutex();
return nullptr;
}
float CalculateMeanOperation::calculate_mean_tile(MemoryBuffer *tile) const
{
float *buffer = tile->get_buffer();
int size = tile->get_width() * tile->get_height();
int pixels = 0;
float sum = 0.0f;
for (int i = 0, offset = 0; i < size; i++, offset += 4) {
if (buffer[offset + 3] > 0) {
pixels++;
switch (setting_) {
case 1: {
sum += IMB_colormanagement_get_luminance(&buffer[offset]);
break;
}
case 2: {
sum += buffer[offset];
break;
}
case 3: {
sum += buffer[offset + 1];
break;
}
case 4: {
sum += buffer[offset + 2];
break;
}
case 5: {
float yuv[3];
rgb_to_yuv(buffer[offset],
buffer[offset + 1],
buffer[offset + 2],
&yuv[0],
&yuv[1],
&yuv[2],
BLI_YUV_ITU_BT709);
sum += yuv[0];
break;
}
}
}
}
return sum / pixels;
}
void CalculateMeanOperation::set_setting(int setting)

24
source/blender/compositor/operations/COM_CalculateMeanOperation.h
View File

@ -22,11 +22,6 @@ class CalculateMeanOperation : public ConstantOperation {
};
protected:
/**
* \brief Cached reference to the reader
*/
SocketReader *image_reader_;
bool is_calculated_;
float constant_value_;
int setting_;
@ -35,26 +30,8 @@ class CalculateMeanOperation : public ConstantOperation {
public:
CalculateMeanOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel(float output[4], int x, int y, void *data) override;
/**
* Initialize the execution
*/
void init_execution() override;
void *initialize_tile_data(rcti *rect) override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
bool determine_depending_area_of_interest(rcti *input,
ReadBufferOperation *read_operation,
rcti *output) override;
void set_setting(int setting);
void get_area_of_interest(int input_idx, const rcti &output_area, rcti &r_input_area) override;
@ -71,7 +48,6 @@ class CalculateMeanOperation : public ConstantOperation {
* The caller takes care of checking the value is only calculated once. */
virtual float calculate_value(const MemoryBuffer *input) const;
float calculate_mean_tile(MemoryBuffer *tile) const;
float calculate_mean(const MemoryBuffer *input) const;
private:

71
source/blender/compositor/operations/COM_CalculateStandardDeviationOperation.cc
View File

@ -10,77 +10,6 @@
namespace blender::compositor {
void CalculateStandardDeviationOperation::execute_pixel(float output[4],
int /*x*/,
int /*y*/,
void * /*data*/)
{
output[0] = standard_deviation_;
}
void *CalculateStandardDeviationOperation::initialize_tile_data(rcti *rect)
{
lock_mutex();
if (!is_calculated_) {
MemoryBuffer *tile = (MemoryBuffer *)image_reader_->initialize_tile_data(rect);
standard_deviation_ = 0.0f;
float *buffer = tile->get_buffer();
int size = tile->get_width() * tile->get_height();
int pixels = 0;
float sum = 0.0f;
const float mean = this->calculate_mean_tile(tile);
for (int i = 0, offset = 0; i < size; i++, offset += 4) {
if (buffer[offset + 3] > 0) {
pixels++;
switch (setting_) {
case 1: /* rgb combined */
{
float value = IMB_colormanagement_get_luminance(&buffer[offset]);
sum += (value - mean) * (value - mean);
break;
}
case 2: /* red */
{
float value = buffer[offset];
sum += (value - mean) * (value - mean);
break;
}
case 3: /* green */
{
float value = buffer[offset + 1];
sum += (value - mean) * (value - mean);
break;
}
case 4: /* blue */
{
float value = buffer[offset + 2];
sum += (value - mean) * (value - mean);
break;
}
case 5: /* luminance */
{
float yuv[3];
rgb_to_yuv(buffer[offset],
buffer[offset + 1],
buffer[offset + 2],
&yuv[0],
&yuv[1],
&yuv[2],
BLI_YUV_ITU_BT709);
sum += (yuv[0] - mean) * (yuv[0] - mean);
break;
}
}
}
}
standard_deviation_ = sqrt(sum / float(pixels - 1));
is_calculated_ = true;
}
unlock_mutex();
return nullptr;
}
float CalculateStandardDeviationOperation::calculate_value(const MemoryBuffer *input) const
{
const float mean = this->calculate_mean(input);

9
source/blender/compositor/operations/COM_CalculateStandardDeviationOperation.h
View File

@ -19,15 +19,6 @@ class CalculateStandardDeviationOperation : public CalculateMeanOperation {
protected:
float standard_deviation_;
public:
/**
* The inner loop of this operation.
*/
void execute_pixel(float output[4], int x, int y, void *data) override;
void *initialize_tile_data(rcti *rect) override;
protected:
float calculate_value(const MemoryBuffer *input) const override;
private:

42
source/blender/compositor/operations/COM_ChangeHSVOperation.cc
View File

@ -13,51 +13,9 @@ ChangeHSVOperation::ChangeHSVOperation()
this->add_input_socket(DataType::Value);
this->add_input_socket(DataType::Value);
this->add_output_socket(DataType::Color);
input_operation_ = nullptr;
flags_.can_be_constant = true;
}
void ChangeHSVOperation::init_execution()
{
input_operation_ = get_input_socket_reader(0);
hue_operation_ = get_input_socket_reader(1);
saturation_operation_ = get_input_socket_reader(2);
value_operation_ = get_input_socket_reader(3);
}
void ChangeHSVOperation::deinit_execution()
{
input_operation_ = nullptr;
hue_operation_ = nullptr;
saturation_operation_ = nullptr;
value_operation_ = nullptr;
}
void ChangeHSVOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float input_color1[4];
float hue[4], saturation[4], value[4];
input_operation_->read_sampled(input_color1, x, y, sampler);
hue_operation_->read_sampled(hue, x, y, sampler);
saturation_operation_->read_sampled(saturation, x, y, sampler);
value_operation_->read_sampled(value, x, y, sampler);
output[0] = input_color1[0] + (hue[0] - 0.5f);
if (output[0] > 1.0f) {
output[0] -= 1.0f;
}
else if (output[0] < 0.0f) {
output[0] += 1.0f;
}
output[1] = input_color1[1] * saturation[0];
output[2] = input_color1[2] * value[0];
output[3] = input_color1[3];
}
void ChangeHSVOperation::update_memory_buffer_partial(MemoryBuffer *output,
const rcti &area,
Span<MemoryBuffer *> inputs)

17
source/blender/compositor/operations/COM_ChangeHSVOperation.h
View File

@ -13,26 +13,9 @@ namespace blender::compositor {
* it assumes we are in sRGB color space.
*/
class ChangeHSVOperation : public MultiThreadedOperation {
private:
SocketReader *input_operation_;
SocketReader *hue_operation_;
SocketReader *saturation_operation_;
SocketReader *value_operation_;
public:
/**
* Default constructor
*/
ChangeHSVOperation();
void init_execution() override;
void deinit_execution() override;
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
void update_memory_buffer_partial(MemoryBuffer *output,
const rcti &area,
Span<MemoryBuffer *> inputs) override;

47
source/blender/compositor/operations/COM_ChannelMatteOperation.cc
View File

@ -11,14 +11,11 @@ ChannelMatteOperation::ChannelMatteOperation()
add_input_socket(DataType::Color);
add_output_socket(DataType::Value);
input_image_program_ = nullptr;
flags_.can_be_constant = true;
}
void ChannelMatteOperation::init_execution()
{
input_image_program_ = this->get_input_socket_reader(0);
limit_range_ = limit_max_ - limit_min_;
switch (limit_method_) {
@ -63,50 +60,6 @@ void ChannelMatteOperation::init_execution()
}
}
void ChannelMatteOperation::deinit_execution()
{
input_image_program_ = nullptr;
}
void ChannelMatteOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float in_color[4];
float alpha;
const float limit_max = limit_max_;
const float limit_min = limit_min_;
const float limit_range = limit_range_;
input_image_program_->read_sampled(in_color, x, y, sampler);
/* matte operation */
alpha = in_color[ids_[0]] - std::max(in_color[ids_[1]], in_color[ids_[2]]);
/* flip because 0.0 is transparent, not 1.0 */
alpha = 1.0f - alpha;
/* test range */
if (alpha > limit_max) {
alpha = in_color[3]; /* Whatever it was prior. */
}
else if (alpha < limit_min) {
alpha = 0.0f;
}
else { /* Blend. */
alpha = (alpha - limit_min) / limit_range;
}
/* Store matte(alpha) value in [0] to go with
* COM_SetAlphaMultiplyOperation and the Value output.
*/
/* Don't make something that was more transparent less transparent. */
output[0] = std::min(alpha, in_color[3]);
}
void ChannelMatteOperation::update_memory_buffer_partial(MemoryBuffer *output,
const rcti &area,
Span<MemoryBuffer *> inputs)

11
source/blender/compositor/operations/COM_ChannelMatteOperation.h
View File

@ -14,8 +14,6 @@ namespace blender::compositor {
*/
class ChannelMatteOperation : public MultiThreadedOperation {
private:
SocketReader *input_image_program_;
// int color_space_; /* node->custom1 */ /* UNUSED */ /* TODO? */
int matte_channel_; /* node->custom2 */
int limit_method_; /* node->algorithm */
@ -37,18 +35,9 @@ class ChannelMatteOperation : public MultiThreadedOperation {
int ids_[3];
public:
/**
* Default constructor
*/
ChannelMatteOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
void init_execution() override;
void deinit_execution() override;
void set_settings(NodeChroma *node_chroma, const int custom2)
{

81
source/blender/compositor/operations/COM_ChromaMatteOperation.cc
View File

@ -12,90 +12,9 @@ ChromaMatteOperation::ChromaMatteOperation()
add_input_socket(DataType::Color);
add_output_socket(DataType::Value);
input_image_program_ = nullptr;
input_key_program_ = nullptr;
flags_.can_be_constant = true;
}
void ChromaMatteOperation::init_execution()
{
input_image_program_ = this->get_input_socket_reader(0);
input_key_program_ = this->get_input_socket_reader(1);
}
void ChromaMatteOperation::deinit_execution()
{
input_image_program_ = nullptr;
input_key_program_ = nullptr;
}
void ChromaMatteOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float in_key[4];
float in_image[4];
const float acceptance = settings_->t1; /* in radians */
const float cutoff = settings_->t2; /* in radians */
const float gain = settings_->fstrength;
float x_angle, z_angle, alpha;
float theta, beta;
float kfg;
input_key_program_->read_sampled(in_key, x, y, sampler);
input_image_program_->read_sampled(in_image, x, y, sampler);
/* Store matte(alpha) value in [0] to go with
* #COM_SetAlphaMultiplyOperation and the Value output. */
/* Algorithm from book "Video Demystified", does not include the spill reduction part. */
/* Find theta, the angle that the color space should be rotated based on key. */
/* rescale to -1.0..1.0 */
// in_image[0] = (in_image[0] * 2.0f) - 1.0f; // UNUSED
in_image[1] = (in_image[1] * 2.0f) - 1.0f;
in_image[2] = (in_image[2] * 2.0f) - 1.0f;
// in_key[0] = (in_key[0] * 2.0f) - 1.0f; // UNUSED
in_key[1] = (in_key[1] * 2.0f) - 1.0f;
in_key[2] = (in_key[2] * 2.0f) - 1.0f;
theta = atan2(in_key[2], in_key[1]);
/* Rotate the cb and cr into x/z space. */
x_angle = in_image[1] * cosf(theta) + in_image[2] * sinf(theta);
z_angle = in_image[2] * cosf(theta) - in_image[1] * sinf(theta);
/* If within the acceptance angle. */
/* If kfg is <0 then the pixel is outside of the key color. */
kfg = x_angle - (fabsf(z_angle) / tanf(acceptance / 2.0f));
if (kfg > 0.0f) { /* found a pixel that is within key color */
alpha = 1.0f - (kfg / gain);
beta = atan2(z_angle, x_angle);
/* if beta is within the cutoff angle */
if (fabsf(beta) < (cutoff / 2.0f)) {
alpha = 0.0f;
}
/* don't make something that was more transparent less transparent */
if (alpha < in_image[3]) {
output[0] = alpha;
}
else {
output[0] = in_image[3];
}
}
else { /* Pixel is outside key color. */
output[0] = in_image[3]; /* Make pixel just as transparent as it was before. */
}
}
void ChromaMatteOperation::update_memory_buffer_partial(MemoryBuffer *output,
const rcti &area,
Span<MemoryBuffer *> inputs)

13
source/blender/compositor/operations/COM_ChromaMatteOperation.h
View File

@ -15,23 +15,10 @@ namespace blender::compositor {
class ChromaMatteOperation : public MultiThreadedOperation {
private:
NodeChroma *settings_;
SocketReader *input_image_program_;
SocketReader *input_key_program_;
public:
/**
* Default constructor
*/
ChromaMatteOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
void init_execution() override;
void deinit_execution() override;
void set_settings(NodeChroma *node_chroma)
{
settings_ = node_chroma;

38
source/blender/compositor/operations/COM_ColorBalanceASCCDLOperation.cc
View File

@ -23,42 +23,10 @@ ColorBalanceASCCDLOperation::ColorBalanceASCCDLOperation()
this->add_input_socket(DataType::Value);
this->add_input_socket(DataType::Color);
this->add_output_socket(DataType::Color);
input_value_operation_ = nullptr;
input_color_operation_ = nullptr;
this->set_canvas_input_index(1);
flags_.can_be_constant = true;
}
void ColorBalanceASCCDLOperation::init_execution()
{
input_value_operation_ = this->get_input_socket_reader(0);
input_color_operation_ = this->get_input_socket_reader(1);
}
void ColorBalanceASCCDLOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float input_color[4];
float value[4];
input_value_operation_->read_sampled(value, x, y, sampler);
input_color_operation_->read_sampled(input_color, x, y, sampler);
float fac = value[0];
fac = std::min(1.0f, fac);
const float mfac = 1.0f - fac;
output[0] = mfac * input_color[0] +
fac * colorbalance_cdl(input_color[0], offset_[0], power_[0], slope_[0]);
output[1] = mfac * input_color[1] +
fac * colorbalance_cdl(input_color[1], offset_[1], power_[1], slope_[1]);
output[2] = mfac * input_color[2] +
fac * colorbalance_cdl(input_color[2], offset_[2], power_[2], slope_[2]);
output[3] = input_color[3];
}
void ColorBalanceASCCDLOperation::update_memory_buffer_row(PixelCursor &p)
{
for (; p.out < p.row_end; p.next()) {
@ -76,10 +44,4 @@ void ColorBalanceASCCDLOperation::update_memory_buffer_row(PixelCursor &p)
}
}
void ColorBalanceASCCDLOperation::deinit_execution()
{
input_value_operation_ = nullptr;
input_color_operation_ = nullptr;
}
} // namespace blender::compositor

24
source/blender/compositor/operations/COM_ColorBalanceASCCDLOperation.h
View File

@ -14,37 +14,13 @@ namespace blender::compositor {
*/
class ColorBalanceASCCDLOperation : public MultiThreadedRowOperation {
protected:
/**
* Prefetched reference to the input_program
*/
SocketReader *input_value_operation_;
SocketReader *input_color_operation_;
float offset_[3];
float power_[3];
float slope_[3];
public:
/**
* Default constructor
*/
ColorBalanceASCCDLOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
void set_offset(float offset[3])
{
copy_v3_v3(offset_, offset);

38
source/blender/compositor/operations/COM_ColorBalanceLGGOperation.cc
View File

@ -30,42 +30,10 @@ ColorBalanceLGGOperation::ColorBalanceLGGOperation()
this->add_input_socket(DataType::Value);
this->add_input_socket(DataType::Color);
this->add_output_socket(DataType::Color);
input_value_operation_ = nullptr;
input_color_operation_ = nullptr;
this->set_canvas_input_index(1);
flags_.can_be_constant = true;
}
void ColorBalanceLGGOperation::init_execution()
{
input_value_operation_ = this->get_input_socket_reader(0);
input_color_operation_ = this->get_input_socket_reader(1);
}
void ColorBalanceLGGOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float input_color[4];
float value[4];
input_value_operation_->read_sampled(value, x, y, sampler);
input_color_operation_->read_sampled(input_color, x, y, sampler);
float fac = value[0];
fac = std::min(1.0f, fac);
const float mfac = 1.0f - fac;
output[0] = mfac * input_color[0] +
fac * colorbalance_lgg(input_color[0], lift_[0], gamma_inv_[0], gain_[0]);
output[1] = mfac * input_color[1] +
fac * colorbalance_lgg(input_color[1], lift_[1], gamma_inv_[1], gain_[1]);
output[2] = mfac * input_color[2] +
fac * colorbalance_lgg(input_color[2], lift_[2], gamma_inv_[2], gain_[2]);
output[3] = input_color[3];
}
void ColorBalanceLGGOperation::update_memory_buffer_row(PixelCursor &p)
{
for (; p.out < p.row_end; p.next()) {
@ -83,10 +51,4 @@ void ColorBalanceLGGOperation::update_memory_buffer_row(PixelCursor &p)
}
}
void ColorBalanceLGGOperation::deinit_execution()
{
input_value_operation_ = nullptr;
input_color_operation_ = nullptr;
}
} // namespace blender::compositor

24
source/blender/compositor/operations/COM_ColorBalanceLGGOperation.h
View File

@ -14,37 +14,13 @@ namespace blender::compositor {
*/
class ColorBalanceLGGOperation : public MultiThreadedRowOperation {
protected:
/**
* Prefetched reference to the input_program
*/
SocketReader *input_value_operation_;
SocketReader *input_color_operation_;
float gain_[3];
float lift_[3];
float gamma_inv_[3];
public:
/**
* Default constructor
*/
ColorBalanceLGGOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
void set_gain(const float gain[3])
{
copy_v3_v3(gain_, gain);

118
source/blender/compositor/operations/COM_ColorCorrectionOperation.cc
View File

@ -13,18 +13,11 @@ ColorCorrectionOperation::ColorCorrectionOperation()
this->add_input_socket(DataType::Color);
this->add_input_socket(DataType::Value);
this->add_output_socket(DataType::Color);
input_image_ = nullptr;
input_mask_ = nullptr;
red_channel_enabled_ = true;
green_channel_enabled_ = true;
blue_channel_enabled_ = true;
flags_.can_be_constant = true;
}
void ColorCorrectionOperation::init_execution()
{
input_image_ = this->get_input_socket_reader(0);
input_mask_ = this->get_input_socket_reader(1);
}
/* Calculate x^y if the function is defined. Otherwise return the given fallback value. */
BLI_INLINE float color_correct_powf_safe(const float x, const float y, const float fallback_value)
@ -35,111 +28,6 @@ BLI_INLINE float color_correct_powf_safe(const float x, const float y, const flo
return powf(x, y);
}
void ColorCorrectionOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float input_image_color[4];
float input_mask[4];
input_image_->read_sampled(input_image_color, x, y, sampler);
input_mask_->read_sampled(input_mask, x, y, sampler);
float level = (input_image_color[0] + input_image_color[1] + input_image_color[2]) / 3.0f;
float contrast = data_->master.contrast;
float saturation = data_->master.saturation;
float gamma = data_->master.gamma;
float gain = data_->master.gain;
float lift = data_->master.lift;
float r, g, b;
float value = input_mask[0];
value = std::min(1.0f, value);
const float mvalue = 1.0f - value;
float level_shadows = 0.0;
float level_midtones = 0.0;
float level_highlights = 0.0;
#define MARGIN 0.10f
#define MARGIN_DIV (0.5f / MARGIN)
if (level < data_->startmidtones - MARGIN) {
level_shadows = 1.0f;
}
else if (level < data_->startmidtones + MARGIN) {
level_midtones = ((level - data_->startmidtones) * MARGIN_DIV) + 0.5f;
level_shadows = 1.0f - level_midtones;
}
else if (level < data_->endmidtones - MARGIN) {
level_midtones = 1.0f;
}
else if (level < data_->endmidtones + MARGIN) {
level_highlights = ((level - data_->endmidtones) * MARGIN_DIV) + 0.5f;
level_midtones = 1.0f - level_highlights;
}
else {
level_highlights = 1.0f;
}
#undef MARGIN
#undef MARGIN_DIV
contrast *= (level_shadows * data_->shadows.contrast) +
(level_midtones * data_->midtones.contrast) +
(level_highlights * data_->highlights.contrast);
saturation *= (level_shadows * data_->shadows.saturation) +
(level_midtones * data_->midtones.saturation) +
(level_highlights * data_->highlights.saturation);
gamma *= (level_shadows * data_->shadows.gamma) + (level_midtones * data_->midtones.gamma) +
(level_highlights * data_->highlights.gamma);
gain *= (level_shadows * data_->shadows.gain) + (level_midtones * data_->midtones.gain) +
(level_highlights * data_->highlights.gain);
lift += (level_shadows * data_->shadows.lift) + (level_midtones * data_->midtones.lift) +
(level_highlights * data_->highlights.lift);
float invgamma = 1.0f / gamma;
float luma = IMB_colormanagement_get_luminance(input_image_color);
r = input_image_color[0];
g = input_image_color[1];
b = input_image_color[2];
r = (luma + saturation * (r - luma));
g = (luma + saturation * (g - luma));
b = (luma + saturation * (b - luma));
r = 0.5f + ((r - 0.5f) * contrast);
g = 0.5f + ((g - 0.5f) * contrast);
b = 0.5f + ((b - 0.5f) * contrast);
/* Check for negative values to avoid nan. */
r = color_correct_powf_safe(r * gain + lift, invgamma, r);
g = color_correct_powf_safe(g * gain + lift, invgamma, g);
b = color_correct_powf_safe(b * gain + lift, invgamma, b);
/* Mix with mask. */
r = mvalue * input_image_color[0] + value * r;
g = mvalue * input_image_color[1] + value * g;
b = mvalue * input_image_color[2] + value * b;
if (red_channel_enabled_) {
output[0] = r;
}
else {
output[0] = input_image_color[0];
}
if (green_channel_enabled_) {
output[1] = g;
}
else {
output[1] = input_image_color[1];
}
if (blue_channel_enabled_) {
output[2] = b;
}
else {
output[2] = input_image_color[2];
}
output[3] = input_image_color[3];
}
void ColorCorrectionOperation::update_memory_buffer_row(PixelCursor &p)
{
for (; p.out < p.row_end; p.next()) {
@ -217,10 +105,4 @@ void ColorCorrectionOperation::update_memory_buffer_row(PixelCursor &p)
}
}
void ColorCorrectionOperation::deinit_execution()
{
input_image_ = nullptr;
input_mask_ = nullptr;
}
} // namespace blender::compositor

20
source/blender/compositor/operations/COM_ColorCorrectionOperation.h
View File

@ -10,11 +10,6 @@ namespace blender::compositor {
class ColorCorrectionOperation : public MultiThreadedRowOperation {
private:
/**
* Cached reference to the input_program
*/
SocketReader *input_image_;
SocketReader *input_mask_;
NodeColorCorrection *data_;
bool red_channel_enabled_;
@ -24,21 +19,6 @@ class ColorCorrectionOperation : public MultiThreadedRowOperation {
public:
ColorCorrectionOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
void set_data(NodeColorCorrection *data)
{
data_ = data;

94
source/blender/compositor/operations/COM_ColorCurveOperation.cc
View File

@ -16,72 +16,15 @@ ColorCurveOperation::ColorCurveOperation()
this->add_input_socket(DataType::Color);
this->add_output_socket(DataType::Color);
input_fac_program_ = nullptr;
input_image_program_ = nullptr;
input_black_program_ = nullptr;
input_white_program_ = nullptr;
this->set_canvas_input_index(1);
}
void ColorCurveOperation::init_execution()
{
CurveBaseOperation::init_execution();
input_fac_program_ = this->get_input_socket_reader(0);
input_image_program_ = this->get_input_socket_reader(1);
input_black_program_ = this->get_input_socket_reader(2);
input_white_program_ = this->get_input_socket_reader(3);
BKE_curvemapping_premultiply(curve_mapping_, false);
}
void ColorCurveOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
CurveMapping *cumap = curve_mapping_;
float fac[4];
float image[4];
/* local versions of cumap->black, cumap->white, cumap->bwmul */
float black[4];
float white[4];
float bwmul[3];
input_black_program_->read_sampled(black, x, y, sampler);
input_white_program_->read_sampled(white, x, y, sampler);
/* get our own local bwmul value,
* since we can't be threadsafe and use cumap->bwmul & friends */
BKE_curvemapping_set_black_white_ex(black, white, bwmul);
input_fac_program_->read_sampled(fac, x, y, sampler);
input_image_program_->read_sampled(image, x, y, sampler);
if (*fac >= 1.0f) {
BKE_curvemapping_evaluate_premulRGBF_ex(cumap, output, image, black, bwmul);
}
else if (*fac <= 0.0f) {
copy_v3_v3(output, image);
}
else {
float col[4];
BKE_curvemapping_evaluate_premulRGBF_ex(cumap, col, image, black, bwmul);
interp_v3_v3v3(output, image, col, *fac);
}
output[3] = image[3];
}
void ColorCurveOperation::deinit_execution()
{
CurveBaseOperation::deinit_execution();
input_fac_program_ = nullptr;
input_image_program_ = nullptr;
input_black_program_ = nullptr;
input_white_program_ = nullptr;
}
void ColorCurveOperation::update_memory_buffer_partial(MemoryBuffer *output,
const rcti &area,
Span<MemoryBuffer *> inputs)
@ -120,54 +63,17 @@ ConstantLevelColorCurveOperation::ConstantLevelColorCurveOperation()
this->add_input_socket(DataType::Color);
this->add_output_socket(DataType::Color);
input_fac_program_ = nullptr;
input_image_program_ = nullptr;
this->set_canvas_input_index(1);
}
void ConstantLevelColorCurveOperation::init_execution()
{
CurveBaseOperation::init_execution();
input_fac_program_ = this->get_input_socket_reader(0);
input_image_program_ = this->get_input_socket_reader(1);
BKE_curvemapping_premultiply(curve_mapping_, false);
BKE_curvemapping_set_black_white(curve_mapping_, black_, white_);
}
void ConstantLevelColorCurveOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float fac[4];
float image[4];
input_fac_program_->read_sampled(fac, x, y, sampler);
input_image_program_->read_sampled(image, x, y, sampler);
if (*fac >= 1.0f) {
BKE_curvemapping_evaluate_premulRGBF(curve_mapping_, output, image);
}
else if (*fac <= 0.0f) {
copy_v3_v3(output, image);
}
else {
float col[4];
BKE_curvemapping_evaluate_premulRGBF(curve_mapping_, col, image);
interp_v3_v3v3(output, image, col, *fac);
}
output[3] = image[3];
}
void ConstantLevelColorCurveOperation::deinit_execution()
{
CurveBaseOperation::deinit_execution();
input_fac_program_ = nullptr;
input_image_program_ = nullptr;
}
void ConstantLevelColorCurveOperation::update_memory_buffer_partial(MemoryBuffer *output,
const rcti &area,
Span<MemoryBuffer *> inputs)

40
source/blender/compositor/operations/COM_ColorCurveOperation.h
View File

@ -9,33 +9,11 @@
namespace blender::compositor {
class ColorCurveOperation : public CurveBaseOperation {
private:
/**
* Cached reference to the input_program
*/
SocketReader *input_fac_program_;
SocketReader *input_image_program_;
SocketReader *input_black_program_;
SocketReader *input_white_program_;
public:
ColorCurveOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
void update_memory_buffer_partial(MemoryBuffer *output,
const rcti &area,
Span<MemoryBuffer *> inputs) override;
@ -43,32 +21,14 @@ class ColorCurveOperation : public CurveBaseOperation {
class ConstantLevelColorCurveOperation : public CurveBaseOperation {
private:
/**
* Cached reference to the input_program
*/
SocketReader *input_fac_program_;
SocketReader *input_image_program_;
float black_[3];
float white_[3];
public:
ConstantLevelColorCurveOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
void set_black_level(float black[3])
{
copy_v3_v3(black_, black);

31
source/blender/compositor/operations/COM_ColorExposureOperation.cc
View File

@ -11,34 +11,9 @@ ExposureOperation::ExposureOperation()
this->add_input_socket(DataType::Color);
this->add_input_socket(DataType::Value);
this->add_output_socket(DataType::Color);
input_program_ = nullptr;
flags_.can_be_constant = true;
}
void ExposureOperation::init_execution()
{
input_program_ = this->get_input_socket_reader(0);
input_exposure_program_ = this->get_input_socket_reader(1);
}
void ExposureOperation::execute_pixel_sampled(float output[4],
float x,
float y,
PixelSampler sampler)
{
float input_value[4];
float input_exposure[4];
input_program_->read_sampled(input_value, x, y, sampler);
input_exposure_program_->read_sampled(input_exposure, x, y, sampler);
const float exposure = pow(2, input_exposure[0]);
output[0] = input_value[0] * exposure;
output[1] = input_value[1] * exposure;
output[2] = input_value[2] * exposure;
output[3] = input_value[3];
}
void ExposureOperation::update_memory_buffer_row(PixelCursor &p)
{
for (; p.out < p.row_end; p.next()) {
@ -52,10 +27,4 @@ void ExposureOperation::update_memory_buffer_row(PixelCursor &p)
}
}
void ExposureOperation::deinit_execution()
{
input_program_ = nullptr;
input_exposure_program_ = nullptr;
}
} // namespace blender::compositor

22
source/blender/compositor/operations/COM_ColorExposureOperation.h
View File

@ -9,31 +9,9 @@
namespace blender::compositor {
class ExposureOperation : public MultiThreadedRowOperation {
private:
/**
* Cached reference to the input_program
*/
SocketReader *input_program_;
SocketReader *input_exposure_program_;
public:
ExposureOperation();
/**
* The inner loop of this operation.
*/
void execute_pixel_sampled(float output[4], float x, float y, PixelSampler sampler) override;
/**
* Initialize the execution
*/
void init_execution() override;
/**
* Deinitialize the execution
*/
void deinit_execution() override;
void update_memory_buffer_row(PixelCursor &p) override;
};

Some files were not shown because too many files have changed in this diff Show More