diff --git a/app/node/CMakeLists.txt b/app/node/CMakeLists.txt index dcaa35d5a..64c0ca0aa 100644 --- a/app/node/CMakeLists.txt +++ b/app/node/CMakeLists.txt @@ -24,6 +24,8 @@ add_subdirectory(output) set(OLIVE_SOURCES ${OLIVE_SOURCES} + node/code.h + node/code.cpp node/dependency.h node/dependency.cpp node/edge.h diff --git a/app/node/blend/alphaover/alphaover.cpp b/app/node/blend/alphaover/alphaover.cpp index e1cf15295..2a53dc36a 100644 --- a/app/node/blend/alphaover/alphaover.cpp +++ b/app/node/blend/alphaover/alphaover.cpp @@ -45,6 +45,22 @@ QString AlphaOverBlend::Description() return tr("A blending node that composites one texture over another using its alpha channel."); } +NodeCode AlphaOverBlend::Code(NodeOutput *output) +{ + if (output == texture_output()) { + return NodeCode("AlphaOver", + "void AlphaOver(const pixel *base_in, const pixel *blend_in, pixel *tex_out) {" + " int i = get_global_id(0);" + " tex_out[i].r = base_in.r - blend_in.a + blend_in.r;" + " tex_out[i].g = base_in.g - blend_in.a + blend_in.g;" + " tex_out[i].b = base_in.b - blend_in.a + blend_in.b;" + " tex_out[i].a = base_in.a - blend_in.a + blend_in.a;" + "}"); + } + + return Node::Code(output); +} + void AlphaOverBlend::Release() { } diff --git a/app/node/blend/alphaover/alphaover.h b/app/node/blend/alphaover/alphaover.h index 4ab025289..83fe15ffd 100644 --- a/app/node/blend/alphaover/alphaover.h +++ b/app/node/blend/alphaover/alphaover.h @@ -33,6 +33,8 @@ public: virtual QString id() override; virtual QString Description() override; + virtual NodeCode Code(NodeOutput* output) override; + virtual void Release() override; protected: diff --git a/app/node/code.cpp b/app/node/code.cpp new file mode 100644 index 000000000..6e3110e73 --- /dev/null +++ b/app/node/code.cpp @@ -0,0 +1,27 @@ +#include "nodecode.h" + +NodeCode::NodeCode() +{ + +} + +NodeCode::NodeCode(const QString &function_name, const QString &code) : + function_name_(function_name), + code_(code) +{ +} + +bool NodeCode::IsValid() +{ + return (!function_name_.isEmpty() && !code_.isEmpty()); +} + +const QString &NodeCode::function_name() +{ + return function_name_; +} + +const QString &NodeCode::code() +{ + return code_; +} diff --git a/app/node/code.h b/app/node/code.h new file mode 100644 index 000000000..dcbefc466 --- /dev/null +++ b/app/node/code.h @@ -0,0 +1,23 @@ +#ifndef NODECODE_H +#define NODECODE_H + +#include + +class NodeCode +{ +public: + NodeCode(); + NodeCode(const QString& function_name, const QString& code); + + bool IsValid(); + + const QString& function_name(); + const QString& code(); + +private: + QString function_name_; + + QString code_; +}; + +#endif // NODECODE_H diff --git a/app/node/input/media/video/video.h b/app/node/input/media/video/video.h index b1b90e887..4c6a84ada 100644 --- a/app/node/input/media/video/video.h +++ b/app/node/input/media/video/video.h @@ -18,6 +18,8 @@ public: virtual QString Category() override; virtual QString Description() override; + virtual QString Code() override; + virtual void Release() override; NodeInput* matrix_input(); diff --git a/app/node/node.cpp b/app/node/node.cpp index 9ffbda954..cbb6d9b6d 100644 --- a/app/node/node.cpp +++ b/app/node/node.cpp @@ -278,6 +278,13 @@ QList Node::GetImmediateDependencies() return node_list; } +NodeCode Node::Code(NodeOutput *output) +{ + Q_UNUSED(output) + + return NodeCode(); +} + QList Node::RunDependencies(NodeOutput *output, const rational &time) { Q_UNUSED(output) diff --git a/app/node/node.h b/app/node/node.h index c728891d3..8aa35a5c5 100644 --- a/app/node/node.h +++ b/app/node/node.h @@ -26,6 +26,7 @@ #include #include "common/rational.h" +#include "node/code.h" #include "node/dependency.h" #include "node/input.h" #include "node/output.h" @@ -124,6 +125,11 @@ public: */ QList GetImmediateDependencies(); + /** + * @brief Generate OpenCL hardware accelerated code for this Node + */ + virtual NodeCode Code(NodeOutput* output); + /** * @brief Wrapper for Process() * diff --git a/app/node/output/timeline/tracklist.h b/app/node/output/timeline/tracklist.h index ee0d20de6..bc1dd6182 100644 --- a/app/node/output/timeline/tracklist.h +++ b/app/node/output/timeline/tracklist.h @@ -111,6 +111,7 @@ private: private slots: void UpdateTotalLength(); + }; #endif // TRACKLIST_H diff --git a/app/node/output/track/track.h b/app/node/output/track/track.h index f72913eb5..bed2746ca 100644 --- a/app/node/output/track/track.h +++ b/app/node/output/track/track.h @@ -174,6 +174,8 @@ private: NodeOutput* track_output_; + NodeOutput* block_list_output_; + TrackType track_type_; int block_invalidate_cache_stack_; diff --git a/app/node/param.h b/app/node/param.h index 40a832442..0e4f3f26b 100644 --- a/app/node/param.h +++ b/app/node/param.h @@ -87,6 +87,9 @@ public: /// Resolves to `Block*` kBlock, + /// Resolves to `QList` + kBlockList, + /// Resolves to `Footage*` kFootage, diff --git a/app/render/CMakeLists.txt b/app/render/CMakeLists.txt index f360554ff..2cf3cbfb3 100644 --- a/app/render/CMakeLists.txt +++ b/app/render/CMakeLists.txt @@ -15,6 +15,7 @@ # along with this program. If not, see . add_subdirectory(audio) +add_subdirectory(backend) add_subdirectory(gl) add_subdirectory(video) diff --git a/app/render/backend/CMakeLists.txt b/app/render/backend/CMakeLists.txt new file mode 100644 index 000000000..02026e673 --- /dev/null +++ b/app/render/backend/CMakeLists.txt @@ -0,0 +1,28 @@ +# Olive - Non-Linear Video Editor +# Copyright (C) 2019 Olive Team +# +# This program is free software: you can redistribute it and/or modify +# it under the terms of the GNU General Public License as published by +# the Free Software Foundation, either version 3 of the License, or +# (at your option) any later version. +# +# This program is distributed in the hope that it will be useful, +# but WITHOUT ANY WARRANTY; without even the implied warranty of +# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the +# GNU General Public License for more details. +# +# You should have received a copy of the GNU General Public License +# along with this program. If not, see . + +set(OLIVE_SOURCES + ${OLIVE_SOURCES} + render/backend/cudabackend.h + render/backend/cudabackend.cpp + render/backend/metalbackend.h + render/backend/metalbackend.cpp + render/backend/openclbackend.h + render/backend/openclbackend.cpp + render/backend/renderbackend.h + render/backend/renderbackend.cpp + PARENT_SCOPE +) diff --git a/app/render/backend/cudabackend.cpp b/app/render/backend/cudabackend.cpp new file mode 100644 index 000000000..04329de24 --- /dev/null +++ b/app/render/backend/cudabackend.cpp @@ -0,0 +1,6 @@ +#include "cudabackend.h" + +CUDABackend::CUDABackend() +{ + +} diff --git a/app/render/backend/cudabackend.h b/app/render/backend/cudabackend.h new file mode 100644 index 000000000..d47cf7807 --- /dev/null +++ b/app/render/backend/cudabackend.h @@ -0,0 +1,11 @@ +#ifndef CUDABACKEND_H +#define CUDABACKEND_H + + +class CUDABackend +{ +public: + CUDABackend(); +}; + +#endif // CUDABACKEND_H diff --git a/app/render/backend/metalbackend.cpp b/app/render/backend/metalbackend.cpp new file mode 100644 index 000000000..c33c2983e --- /dev/null +++ b/app/render/backend/metalbackend.cpp @@ -0,0 +1,6 @@ +#include "metalbackend.h" + +MetalBackend::MetalBackend() +{ + +} diff --git a/app/render/backend/metalbackend.h b/app/render/backend/metalbackend.h new file mode 100644 index 000000000..4123d6497 --- /dev/null +++ b/app/render/backend/metalbackend.h @@ -0,0 +1,11 @@ +#ifndef METALBACKEND_H +#define METALBACKEND_H + + +class MetalBackend +{ +public: + MetalBackend(); +}; + +#endif // METALBACKEND_H diff --git a/app/render/backend/openclbackend.cpp b/app/render/backend/openclbackend.cpp new file mode 100644 index 000000000..be38e3c6c --- /dev/null +++ b/app/render/backend/openclbackend.cpp @@ -0,0 +1,128 @@ +#include "openclbackend.h" + +OpenCLBackend::OpenCLBackend() +{ + +} + +bool OpenCLBackend::Init() +{ + // Get platform and device information + cl_platform_id platform_id = nullptr; + cl_uint ret_num_devices; + cl_uint ret_num_platforms; + cl_int ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms); + ret = clGetDeviceIDs( platform_id, CL_DEVICE_TYPE_GPU, 1, + &device_id_, &ret_num_devices); + + if (ret != CL_SUCCESS) { + qWarning() << "Failed to find compatible OpenCL device"; + return false; + } + + // Create an OpenCL context + context_ = clCreateContext(nullptr, 1, &device_id_, nullptr, nullptr, &ret); + + // Create a command queue + //command_queue_ = clCreateCommandQueue(context_, device_id_, 0, &ret); + + return true; +} + +void OpenCLBackend::GenerateFrame(const rational &time) +{ + Q_UNUSED(time) + /* + // Copy the lists A and B to their respective memory buffers + cl_int ret = clEnqueueWriteBuffer(command_queue_, a_mem_obj, CL_TRUE, 0, + BITMAP_SZ, source_bmp, 0, NULL, NULL); + + // Set the arguments of the kernel + ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_mem_obj); + ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&c_mem_obj); + + // Execute the OpenCL kernel on the list + size_t global_item_size = BITMAP_SZ; // Process the entire lists + size_t local_item_size = 30; // Divide work items into groups of 64 + ret = clEnqueueNDRangeKernel(command_queue_, kernel, 1, NULL, + &global_item_size, &local_item_size, 0, NULL, NULL); + + if (ret != CL_SUCCESS) { + fprintf(stderr, "Failed to run\n"); + exit(1); + } + + // Read the memory buffer C on the device to the local variable C + ret = clEnqueueReadBuffer(command_queue_, c_mem_obj, CL_TRUE, 0, + BITMAP_SZ, dest_bmp, 0, NULL, NULL); + + // Clean up + ret = clFlush(command_queue_); + ret = clFinish(command_queue_); + */ +} + +void OpenCLBackend::Close() +{ + Decompile(); + + cl_int ret; + + /*ret = clReleaseKernel(kernel); + ret = clReleaseMemObject(a_mem_obj); + ret = clReleaseMemObject(c_mem_obj);*/ + ret = clReleaseCommandQueue(command_queue_); + ret = clReleaseContext(context_); +} + +void OpenCLBackend::Decompile() +{ + clReleaseProgram(program_); +} + +void OpenCLBackend::Compile() +{ + /* + cl_int ret; + + // Create memory buffers on the device for each vector + cl_mem a_mem_obj = clCreateBuffer(context_, CL_MEM_READ_ONLY, BITMAP_SZ, nullptr, &ret); + cl_mem c_mem_obj = clCreateBuffer(context_, CL_MEM_WRITE_ONLY, BITMAP_SZ, nullptr, &ret); + + // Create a program from the kernel source + program_ = clCreateProgramWithSource(context_, + 1, + (const char **)&source_str, + (const size_t *)&source_size, + &ret); + + // Build the program + ret = clBuildProgram(program_, 1, &device_id_, nullptr, nullptr, nullptr); + + if (ret != CL_SUCCESS) { + // Decompile failed, the user will probably want to know why + size_t error_len = 0; + + clGetProgramBuildInfo(program_, device_id_, CL_PROGRAM_BUILD_LOG, 0, nullptr, &error_len); + char* err = new char[error_len]; + clGetProgramBuildInfo(program_, device_id_, CL_PROGRAM_BUILD_LOG, error_len, err, nullptr); + + SetError(err); + //fprintf(stderr, "%s\n", err); + + delete [] err; + } + + // Create the OpenCL kernel + cl_kernel kernel = clCreateKernel(program_, "vector_add", &ret); + */ +} + +void OpenCLBackend::GenerateCode() +{ + if (viewer_node() == nullptr) { + return; + } + + +} diff --git a/app/render/backend/openclbackend.h b/app/render/backend/openclbackend.h new file mode 100644 index 000000000..0ca332071 --- /dev/null +++ b/app/render/backend/openclbackend.h @@ -0,0 +1,40 @@ +#ifndef OPENCLBACKEND_H +#define OPENCLBACKEND_H + +#ifdef __APPLE__ +#include +#else +#include +#endif + +#include "renderbackend.h" + +class OpenCLBackend : public RenderBackend +{ +public: + OpenCLBackend(); + + virtual bool Init() override; + + virtual void GenerateFrame(const rational& time) override; + + virtual void Close() override; + + void Compile(); + +protected: + virtual void Decompile() override; + +private: + void GenerateCode(); + + cl_context context_; + + cl_program program_; + + cl_command_queue command_queue_; + + cl_device_id device_id_; +}; + +#endif // OPENCLBACKEND_H diff --git a/app/render/backend/renderbackend.cpp b/app/render/backend/renderbackend.cpp new file mode 100644 index 000000000..61b99341b --- /dev/null +++ b/app/render/backend/renderbackend.cpp @@ -0,0 +1,34 @@ +#include "renderbackend.h" + +RenderBackend::RenderBackend() : + viewer_node_(nullptr) +{ + +} + +RenderBackend::~RenderBackend() +{ + +} + +const QString &RenderBackend::GetError() +{ + return error_; +} + +void RenderBackend::set_viewer_node(ViewerOutput *viewer_node) +{ + viewer_node_ = viewer_node; + + Decompile(); +} + +void RenderBackend::SetError(const QString &error) +{ + error_ = error; +} + +ViewerOutput *RenderBackend::viewer_node() const +{ + return viewer_node_; +} diff --git a/app/render/backend/renderbackend.h b/app/render/backend/renderbackend.h new file mode 100644 index 000000000..934d3a908 --- /dev/null +++ b/app/render/backend/renderbackend.h @@ -0,0 +1,38 @@ +#ifndef RENDERBACKEND_H +#define RENDERBACKEND_H + +#include "node/output/viewer/viewer.h" + +class RenderBackend : QObject +{ + Q_OBJECT +public: + RenderBackend(); + virtual ~RenderBackend(); + + virtual bool Init() = 0; + + virtual void GenerateFrame(const rational& time) = 0; + + virtual void GenerateSamples(const rational& time, const rational& length) = 0; + + virtual void Close() = 0; + + const QString& GetError(); + + void set_viewer_node(ViewerOutput* viewer_node); + +protected: + virtual void Decompile() = 0; + + void SetError(const QString& error); + + ViewerOutput* viewer_node() const; + +private: + ViewerOutput* viewer_node_; + + QString error_; +}; + +#endif // RENDERBACKEND_H diff --git a/app/widget/viewer/viewer.h b/app/widget/viewer/viewer.h index 182fb7969..62ba637d2 100644 --- a/app/widget/viewer/viewer.h +++ b/app/widget/viewer/viewer.h @@ -32,6 +32,7 @@ #include "node/output/viewer/viewer.h" #include "render/audio/audiorenderer.h" #include "render/video/videorenderer.h" +#include "render/backend/openclbackend.h" #include "viewerglwidget.h" #include "viewersizer.h" #include "widget/playbackcontrols/playbackcontrols.h" @@ -115,6 +116,10 @@ private: AudioRendererProcessor* audio_renderer_; + // FIXME: Test code only + OpenCLBackend opencl_backend_; + // End test code + ViewerSizer* sizer_; ViewerGLWidget* gl_widget_;