Support running OpenCL kernels

[jerryct]

The PR https://github.com/KomputeProject/kompute/pull/250 should start a discussion about a potential change to run OpenCL kernels. Demo / testing function to compile shaders uses glslangValidator to compile GLSL into SPIR-V. Similar to this https://github.com/google/clspv exists to compile OpenCL kernels into SPIR-V. So the examples from the documentation can be rewritten as

#define VULKAN_HPP_NO_SMART_HANDLE
#include "kompute/Algorithm.hpp"
#include "kompute/Manager.hpp"
#include "kompute/Sequence.hpp"
#include "kompute/Tensor.hpp"
#include "kompute/operations/OpTensorSyncDevice.hpp"
#include "kompute/operations/OpTensorSyncLocal.hpp"
#include <fmt/format.h>
#include <fmt/ostream.h>
#include <fstream>
#include <string>
#include <vector>

std::vector<uint32_t> compileSource(const std::string &source) {
  if (system(std::string("clspv/_build/bin/clspv - -o tmp_kp_opencl.cl.spv << END\n" + source + "\nEND").c_str()))
    throw std::runtime_error("Error running glslangValidator command");
  std::ifstream fileStream("tmp_kp_opencl.cl.spv", std::ios::binary);
  std::vector<char> buffer;
  buffer.insert(buffer.begin(), std::istreambuf_iterator<char>(fileStream), {});
  return {(uint32_t *)buffer.data(), (uint32_t *)(buffer.data() + buffer.size())};
}

int main() {
  std::string shader(R"(
__kernel void foo(__global float* x, float y) {
    x[get_global_id(0)]+=y;
})");

  std::vector<uint32_t> spirv = compileSource(shader);

  kp::Manager mgr(0, {}, {});

  std::vector<std::shared_ptr<kp::Tensor>> tensors{mgr.tensor({23, 0, 0})};

  std::shared_ptr<kp::Algorithm> algo = mgr.algorithm(tensors, spirv, kp::Workgroup({3}), {}, {42.0});

  std::shared_ptr<kp::Sequence> sq = mgr.sequence()
                                         ->record<kp::OpTensorSyncDevice>(tensors)
                                         ->record<kp::OpAlgoDispatch>(algo) //, std::vector<float>{72.0})
                                         ->record<kp::OpAlgoDispatch>(algo)
                                         ->record<kp::OpTensorSyncLocal>(tensors)
                                         ->eval();

  for (float f : tensors[0]->vector<float>()) {
    fmt::print("{}, ", f);
  }
  fmt::print("\n");

  tensors[0]->data<float>()[0] = 23,

  sq->eval();

  for (float f : tensors[0]->vector<float>()) {
    fmt::print("{}, ", f);
  }
  fmt::print("\n");
}

The code above only changed glslangValidator in compileSource to clspv and shader is written in OpenCL C. It prints as expected.

107, 84, 84, 
107, 168, 168,

Currently kcompute hard-codes the entry point to main. The change I had to made to get the above example working was changing main to foo. The reason for this is that OpenCL does not allow a kernel to be named main:

$ cat test.cl 
__kernel void main(__global float* x, float y) {
    x[get_global_id(0)]=y;
}
$ clspv/_build/bin/clspv test.cl
test.cl:1:15: error: kernel cannot be called 'main'
__kernel void main(__global float* x, float y) {
              ^

The question now could we change the default? One idea could be to choose a different hard-coded name, e.g. entry_point. Or more flexible, allow to specify an entry_point here either by adding another parameter or create a small class wrapping the entry point name and the binary.