cl_khr_command_buffer(3) Manual Page

  #define CL_CHECK(ERROR)                             \
    if (ERROR) {                                      \
      std::cerr << "OpenCL error: " << ERROR << "\n"; \
      return ERROR;                                   \
    }

  int main() {
    cl_platform_id platform;
    CL_CHECK(clGetPlatformIDs(1, &platform, nullptr));
    cl_device_id device;
    CL_CHECK(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, nullptr));

    cl_int error;
    cl_context context =
        clCreateContext(nullptr, 1, &device, nullptr, nullptr, &error);
    CL_CHECK(error);

    const char* code = R"OpenCLC(
  kernel void vector_addition(global int* tile1, global int* tile2,
                              global int* res) {
    size_t index = get_global_id(0);
    res[index] = tile1[index] + tile2[index];
  }
  )OpenCLC";
    const size_t length = std::strlen(code);

    cl_program program =
        clCreateProgramWithSource(context, 1, &code, &length, &error);
    CL_CHECK(error);

    CL_CHECK(clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr));

    cl_kernel kernel = clCreateKernel(program, "vector_addition", &error);
    CL_CHECK(error);

    constexpr size_t frame_count = 60;
    constexpr size_t frame_elements = 1024;
    constexpr size_t frame_size = frame_elements * sizeof(cl_int);

    constexpr size_t tile_count = 16;
    constexpr size_t tile_elements = frame_elements / tile_count;
    constexpr size_t tile_size = tile_elements * sizeof(cl_int);

    cl_mem buffer_tile1 =
        clCreateBuffer(context, CL_MEM_READ_ONLY, tile_size, nullptr, &error);
    CL_CHECK(error);
    cl_mem buffer_tile2 =
        clCreateBuffer(context, CL_MEM_READ_ONLY, tile_size, nullptr, &error);
    CL_CHECK(error);
    cl_mem buffer_res =
        clCreateBuffer(context, CL_MEM_WRITE_ONLY, tile_size, nullptr, &error);
    CL_CHECK(error);

    CL_CHECK(clSetKernelArg(kernel, 0, sizeof(buffer_tile1), &buffer_tile1));
    CL_CHECK(clSetKernelArg(kernel, 1, sizeof(buffer_tile2), &buffer_tile2));
    CL_CHECK(clSetKernelArg(kernel, 2, sizeof(buffer_res), &buffer_res));

    cl_command_queue command_queue =
      clCreateCommandQueue(context, device,
                           CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &error);
    CL_CHECK(error);

    cl_command_buffer_khr command_buffer =
        clCreateCommandBufferKHR(1, &command_queue, nullptr, &error);
    CL_CHECK(error);

    cl_mem buffer_src1 =
        clCreateBuffer(context, CL_MEM_READ_ONLY, frame_size, nullptr, &error);
    CL_CHECK(error);
    cl_mem buffer_src2 =
        clCreateBuffer(context, CL_MEM_READ_ONLY, frame_size, nullptr, &error);
    CL_CHECK(error);
    cl_mem buffer_dst =
        clCreateBuffer(context, CL_MEM_WRITE_ONLY, frame_size, nullptr, &error);
    CL_CHECK(error);

    cl_sync_point_khr tile_sync_point = 0;
    for (size_t tile_index = 0; tile_index < tile_count; tile_index++) {
      std::array<cl_sync_point_khr, 2> copy_sync_points;
      CL_CHECK(clCommandCopyBufferKHR(command_buffer,
          command_queue, buffer_src1, buffer_tile1, tile_index * tile_size, 0,
          tile_size, tile_sync_point ? 1 : 0,
          tile_sync_point ? &tile_sync_point : nullptr, &copy_sync_points[0]),
          nullptr);
      CL_CHECK(clCommandCopyBufferKHR(command_buffer,
          command_queue, buffer_src2, buffer_tile2, tile_index * tile_size, 0,
          tile_size, tile_sync_point ? 1 : 0,
          tile_sync_point ? &tile_sync_point : nullptr, &copy_sync_points[1]),
          nullptr);

      cl_sync_point_khr nd_sync_point;
      CL_CHECK(clCommandNDRangeKernelKHR(command_buffer,
          command_queue, nullptr, kernel, 1, nullptr, &tile_elements, nullptr,
          copy_sync_points.size(), copy_sync_points.data(), &nd_sync_point,
          nullptr));

      CL_CHECK(clCommandCopyBufferKHR(command_buffer,
          command_queue, buffer_res, buffer_dst, 0, tile_index * tile_size,
          tile_size, 1, &nd_sync_point, &tile_sync_point, nullptr));
    }

    CL_CHECK(clFinalizeCommandBufferKHR(command_buffer));

    std::random_device random_device;
    std::mt19937 random_engine{random_device()};
    std::uniform_int_distribution<cl_int> random_distribution{
        0, std::numeric_limits<cl_int>::max() / 2};
    auto random_generator = [&]() { return random_distribution(random_engine); };

    for (size_t frame_index = 0; frame_index < frame_count; frame_index++) {
      std::array<cl_event, 2> write_src_events;
      std::vector<cl_int> src1(frame_elements);
      std::generate(src1.begin(), src1.end(), random_generator);
      CL_CHECK(clEnqueueWriteBuffer(command_queue, buffer_src1, CL_FALSE, 0,
                                    frame_size, src1.data(), 0, nullptr,
                                    &write_src_events[0]));
      std::vector<cl_int> src2(frame_elements);
      std::generate(src2.begin(), src2.end(), random_generator);
      CL_CHECK(clEnqueueWriteBuffer(command_queue, buffer_src2, CL_FALSE, 0,
                                    frame_size, src2.data(), 0, nullptr,
                                    &write_src_events[1]));

      CL_CHECK(clEnqueueCommandBufferKHR(0, NULL, command_buffer, 2,
                                         write_src_events.data(), nullptr));

      CL_CHECK(clFinish(command_queue));

      CL_CHECK(clReleaseEvent(write_src_event[0]));
      CL_CHECK(clReleaseEvent(write_src_event[1]));
    }

    CL_CHECK(clReleaseCommandBufferKHR(command_buffer));
    CL_CHECK(clReleaseCommandQueue(command_queue));

    CL_CHECK(clReleaseMemObject(buffer_src1));
    CL_CHECK(clReleaseMemObject(buffer_src2));
    CL_CHECK(clReleaseMemObject(buffer_dst));

    CL_CHECK(clReleaseMemObject(buffer_tile1));
    CL_CHECK(clReleaseMemObject(buffer_tile2));
    CL_CHECK(clReleaseMemObject(buffer_res));

    CL_CHECK(clReleaseKernel(kernel));
    CL_CHECK(clReleaseProgram(program));
    CL_CHECK(clReleaseContext(context));

    return 0;
  }