повысить производительность opencl

Question

Я пытаюсь реализовать алгоритм обработки изображений с использованием opencl. Но, как я вижу, когда я использую opencl, для завершения одного процесса, т.е. одного кадра, требуется около 0,5 мс. Есть ли способ инициализации параметров opencl только один раз с помощью объявления объекта класса, а не только вызов функции, запускающей основное ядро? Я попытался так, создав класс, но когда я нашел контекст, устройство не может быть объявлено и использовано отдельно, и его нужно создавать каждый раз.

#include <CL/cl.hpp>
#include <chrono>
#include <iostream>
using namespace std::chrono;
using namespace std;
namespace Color {
enum Code {
  FG_RED = 31,
  FG_GREEN = 32,
  FG_BLUE = 34,
  FG_DEFAULT = 39,
  BG_RED = 41,
  BG_GREEN = 42,
  BG_BLUE = 44,
  BG_DEFAULT = 49
};
class Modifier {
  Code code;

 public:
  Modifier(Code pCode) : code(pCode) {}
  friend std::ostream& operator<<(std::ostream& os, const Modifier& mod) {
    return os << "\033[" << mod.code << "m";
  }
};
}  // namespace Color
class useOpenCL {
 public:
  int size = 294400;
  std::vector<cl::Platform> all_platforms;
  std::vector<cl::Device> all_devices;
  cl::Platform default_platform;
  cl::Device default_device;
  cl::Program::Sources sources;
  std::string kernel_code;
  cl::Kernel kernel_add;
  cl::Buffer buffer_A;

  useOpenCL();
  ~useOpenCL() {}
  void backgroundSub();
};

useOpenCL::useOpenCL() {
  Color::Modifier green(Color::FG_GREEN);
  Color::Modifier red(Color::FG_RED);
  Color::Modifier def(Color::FG_DEFAULT);
  // get all platforms (drivers)
  cl::Platform::get(&all_platforms);
  if (all_platforms.size() == 0) {
    std::cout << red << " No platforms found. Check OpenCL installation!" << def
              << endl;
    exit(1);
  }
  default_platform = all_platforms[0];
  std::cout << green << "Using platform: " << def
            << default_platform.getInfo<CL_PLATFORM_NAME>() << std::endl;

  // get default device of the default platform
  default_platform.getDevices(CL_DEVICE_TYPE_ALL, &all_devices);
  if (all_devices.size() == 0) {
    std::cout << red << " No devices found. Check OpenCL installation!" << def
              << endl;
    exit(1);
  }
  default_device = all_devices[0];
  std::cout << green << "Using device: " << def
            << default_device.getInfo<CL_DEVICE_NAME>() << std::endl;

  // kernel calculates for each element C=A+B
  kernel_code =
      "   void kernel simple_add(global const int* A, global const int* B, "
      "global int* C){       "
      "       C[get_global_id(0)]=A[get_global_id(0)]+B[get_global_id(0)];     "
      "            "
      "   }                                                                    "
      "           ";
  sources.push_back({kernel_code.c_str(), kernel_code.length()});
}

void useOpenCL::backgroundSub() {
  int A[size], B[size];
  for (int i = 0; i < size; i++) {
    A[i] = i;
    B[i] = i + 1;
  }
  auto start1 = high_resolution_clock::now();

  cl::Context context({default_device});

  cl::Program program(context, sources);
  if (program.build({default_device}) != CL_SUCCESS) {
    std::cout << " Error building: "
              << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(default_device)
              << "\n";
    exit(1);
  }
  // create buffers on the device
  cl::Buffer buffer_A(context, CL_MEM_READ_WRITE, sizeof(int) * size);
  cl::Buffer buffer_B(context, CL_MEM_READ_WRITE, sizeof(int) * size);
  cl::Buffer buffer_C(context, CL_MEM_READ_WRITE, sizeof(int) * size);

  // create queue to which we will push commands for the device.
  cl::CommandQueue queue(context, default_device);

  // write arrays A and B to the device
  queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, sizeof(int) * size, A);
  queue.enqueueWriteBuffer(buffer_B, CL_TRUE, 0, sizeof(int) * size, B);

  // run the kernel
  /*cl::KernelFunctor
  simple_add(cl::Kernel(program,"simple_add"),queue,cl::NullRange,cl::NDRange(10),cl::NullRange);
  simple_add(buffer_A,buffer_B,buffer_C);*/

  // alternative way to run the kernel
  kernel_add.setArg(0, buffer_A);
  kernel_add.setArg(1, buffer_B);
  kernel_add.setArg(2, buffer_C);
  queue.enqueueNDRangeKernel(kernel_add, cl::NullRange, cl::NDRange(size),
                             cl::NullRange);
  queue.finish();

  int C[size];
  // read result C from the device to array C
  queue.enqueueReadBuffer(buffer_C, CL_TRUE, 0, sizeof(int) * size, C);
  /*std::cout<<" result: \n";
  for(int i=0;i<size;i++){
      std::cout<<C[i]<<"\t";
  }*/
  auto stop1 = high_resolution_clock::now();
  auto duration1 = duration_cast<microseconds>(stop1 - start1);
  auto FPS = 1000000.0 / duration1.count();
  cout << "Segmentation FPS=" << FPS << "\t"
       << "Execution Time(sec)=" << duration1.count() / 1000000.0 << endl;
}

int main() {
  useOpenCL img;
  while (true) {
    img.backgroundSub();
  }
  return 0;
}

Это дает мне следующие результаты:

Segmentation FPS=13.2557    Execution Time(sec)=0.075439
Segmentation FPS=15.7602    Execution Time(sec)=0.063451
Segmentation FPS=14.3872    Execution Time(sec)=0.069506
Segmentation FPS=12.7525    Execution Time(sec)=0.078416

Что не хорошо, так как fps составляет всего 12, 13 кадров в секунду. Так как я могу сделать эту программу быстрее?

Answer 1

Поместите в конструктор часть инициализации, которую вам нужно вызвать только один раз. Эта инициализация должна содержать ВСЕ выделение памяти, компиляцию кода OpenCL C и любые начальные передачи памяти с хоста на устройство:

useOpenCL::useOpenCL() {
  Color::Modifier green(Color::FG_GREEN);
  Color::Modifier red(Color::FG_RED);
  Color::Modifier def(Color::FG_DEFAULT);
  // get all platforms (drivers)
  cl::Platform::get(&all_platforms);
  if (all_platforms.size() == 0) {
    std::cout << red << " No platforms found. Check OpenCL installation!" << def
              << endl;
    exit(1);
  }
  default_platform = all_platforms[0];
  std::cout << green << "Using platform: " << def
            << default_platform.getInfo<CL_PLATFORM_NAME>() << std::endl;

  // get default device of the default platform
  default_platform.getDevices(CL_DEVICE_TYPE_ALL, &all_devices);
  if (all_devices.size() == 0) {
    std::cout << red << " No devices found. Check OpenCL installation!" << def
              << endl;
    exit(1);
  }
  default_device = all_devices[0];
  std::cout << green << "Using device: " << def
            << default_device.getInfo<CL_DEVICE_NAME>() << std::endl;

  // kernel calculates for each element C=A+B
  kernel_code =
      "   void kernel simple_add(global const int* A, global const int* B, "
      "global int* C){       "
      "       C[get_global_id(0)]=A[get_global_id(0)]+B[get_global_id(0)];     "
      "            "
      "   }                                                                    "
      "           ";
  sources.push_back({kernel_code.c_str(), kernel_code.length()});

  context = cl::Context({default_device});

  program = cl::Program(context, sources);
  if (program.build({default_device}) != CL_SUCCESS) {
    std::cout << " Error building: "
              << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(default_device)
              << "\n";
    exit(1);
  }

  // create queue to which we will push commands for the device.
  queue = cl::CommandQueue(context, default_device);

  // create buffers on host
  int A[size], B[size];
  int C[size];
  for (int i = 0; i < size; i++) {
    A[i] = i;
    B[i] = i + 1;
  }

  // create buffers on the device
  buffer_A = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(int) * size);
  buffer_B = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(int) * size);
  buffer_C = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(int) * size);

  // write arrays A and B to the device
  queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, sizeof(int) * size, A);
  queue.enqueueWriteBuffer(buffer_B, CL_TRUE, 0, sizeof(int) * size, B);

  // alternative way to run the kernel
  kernel_add.setArg(0, buffer_A);
  kernel_add.setArg(1, buffer_B);
  kernel_add.setArg(2, buffer_C);
}

Поэтому сделайте context, program, queue, buffer_A , buffer_B, buffer_C переменные-члены вашего класса useOpenCL. В частности, выделение памяти и компиляция занимают много времени, поэтому делайте это только один раз и повторно используйте буферы.

class useOpenCL {
 public:
  int size = 294400;
  std::vector<cl::Platform> all_platforms;
  std::vector<cl::Device> all_devices;
  cl::Platform default_platform;
  cl::Device default_device;
  cl::Program::Sources sources;
  std::string kernel_code;
  cl::Kernel kernel_add;

  cl::Buffer buffer_A;
  cl::Buffer buffer_B;
  cl::Buffer buffer_C;

  cl::Context context;
  cl::Program program;
  cl::CommandQueue queue;

  useOpenCL();
  ~useOpenCL() {}
  void backgroundSub();
};

Тогда для каждого кадра остаются только вызов ядра и, в конечном счете, передача памяти хосту <-> устройства:

void useOpenCL::backgroundSub() {
  auto start1 = high_resolution_clock::now();

  // write arrays A and B to the device (ONLY IF NECESSARY FOR EVERY FRAME)
  //queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, sizeof(int) * size, A);
  //queue.enqueueWriteBuffer(buffer_B, CL_TRUE, 0, sizeof(int) * size, B);

  // run the kernel
  queue.enqueueNDRangeKernel(kernel_add, cl::NullRange, cl::NDRange(size),
                             cl::NullRange);

  // read result C from the device to array C
  queue.enqueueReadBuffer(buffer_C, CL_TRUE, 0, sizeof(int) * size, C);

  queue.finish();

  auto stop1 = high_resolution_clock::now();
  auto duration1 = duration_cast<microseconds>(stop1 - start1);
  auto FPS = 1000000.0 / duration1.count();
  cout << "Segmentation FPS=" << FPS << "\t"
       << "Execution Time(sec)=" << duration1.count() / 1000000.0 << endl;
}

Последний код можно вызывать снова и снова, и он должен быть намного быстрее, чем если бы вы заново инициализировали все снова и снова. Также убедитесь, что size достаточно велико, иначе графический процессор может быть использован не полностью, а задержки при передаче памяти устройства <-> устройства будут непропорционально замедлять каждый кадр.