attempts at folding opencl support in

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

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()
 {
 }

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:

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_;
+}

app/node/code.h

@@ -0,0 +1,23 @@
+#ifndef NODECODE_H
+#define NODECODE_H
+
+#include <QString>
+
+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

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();

app/node/node.cpp

@@ -278,6 +278,13 @@ QList<Node *> Node::GetImmediateDependencies()
   return node_list;
 }
 
+NodeCode Node::Code(NodeOutput *output)
+{
+  Q_UNUSED(output)
+
+  return NodeCode();
+}
+
 QList<NodeDependency> Node::RunDependencies(NodeOutput *output, const rational &time)
 {
   Q_UNUSED(output)

app/node/node.h

@@ -26,6 +26,7 @@
 #include <QObject>
 
 #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<Node*> GetImmediateDependencies();
 
+  /**
+   * @brief Generate OpenCL hardware accelerated code for this Node
+   */
+  virtual NodeCode Code(NodeOutput* output);
+
   /**
    * @brief Wrapper for Process()
    *

app/node/output/timeline/tracklist.h

@@ -111,6 +111,7 @@ private:
 
 private slots:
   void UpdateTotalLength();
+
 };
 
 #endif // TRACKLIST_H

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_;

app/node/param.h

@@ -87,6 +87,9 @@ public:
     /// Resolves to `Block*`
     kBlock,
 
+    /// Resolves to `QList<Block*>`
+    kBlockList,
+
     /// Resolves to `Footage*`
     kFootage,
 

app/render/CMakeLists.txt

@@ -15,6 +15,7 @@
 # along with this program.  If not, see <http://www.gnu.org/licenses/>.
 
 add_subdirectory(audio)
+add_subdirectory(backend)
 add_subdirectory(gl)
 add_subdirectory(video)
 

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 <http://www.gnu.org/licenses/>.
+
+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
+)

app/render/backend/cudabackend.cpp

@@ -0,0 +1,6 @@
+#include "cudabackend.h"
+
+CUDABackend::CUDABackend()
+{
+
+}

app/render/backend/cudabackend.h

@@ -0,0 +1,11 @@
+#ifndef CUDABACKEND_H
+#define CUDABACKEND_H
+
+
+class CUDABackend
+{
+public:
+  CUDABackend();
+};
+
+#endif // CUDABACKEND_H

app/render/backend/metalbackend.cpp

@@ -0,0 +1,6 @@
+#include "metalbackend.h"
+
+MetalBackend::MetalBackend()
+{
+
+}

app/render/backend/metalbackend.h

@@ -0,0 +1,11 @@
+#ifndef METALBACKEND_H
+#define METALBACKEND_H
+
+
+class MetalBackend
+{
+public:
+  MetalBackend();
+};
+
+#endif // METALBACKEND_H

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

app/render/backend/openclbackend.h

@@ -0,0 +1,40 @@
+#ifndef OPENCLBACKEND_H
+#define OPENCLBACKEND_H
+
+#ifdef __APPLE__
+#include <OpenCL/opencl.h>
+#else
+#include <CL/cl.h>
+#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

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_;
+}

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

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_;