How to create program objects from SPIRV with clCreateProgramWithBinary runtime API .

[malixian]

I can emit SPIRV with command clang++ -fsycl -fsycl-device-only -fno-sycl-use-bitcode test.cpp.
And I would like to know how to create program objects from SPIRV with clCreateProgramWithBinary runtime API .
If this way works, can you give me some examples?

[MrSidims]

For OpenCL it is: clCreateProgramWithIL (https://www.khronos.org/registry/OpenCL/sdk/2.1/docs/man/xhtml/clCreateProgramWithIL.html). I think this is an OK example.

[malixian]

thx your reply @MrSidims , I will try it.

[bader]

@malixian, could you clarify how are you going to use the program object?

[malixian]

could you clarify how are you going to use the program object

I want to build a program object by myopencl with spir that is generated by dpcpp. And codes the main program to emti kernel in spir style to myopencl. I think the examples from the link shows me how to program.

[bader]

I'd like to note that ABI for SPIR kernels generated by DPC++ compiler is implementation defined. Some implementation details are documented here: https://github.com/intel/llvm/blob/sycl/sycl/doc/CompilerAndRuntimeDesign.md, but it's subject to change. I would not recommend using this approach "in production".

[malixian]

Whether the SPIR generated by dpcpp can be used to program with amd-opencl clCreateProgramWithBinary ? Whether The SPIR generated by dpcpp has some extensions?

[bader]

Whether the SPIR generated by dpcpp can be used to program with amd-opencl clCreateProgramWithBinary ?

No.
NOTE: DPC++ generates SPIR-V by default and AFAIK, amd-opencl doesn't support SPIR-V.
DPC++ is able to produce executable binary for clCreateProgramWithBinary API in ahead of time compilation mode (AOT), but it's device specific and AMD devices are not support at the moment.

Whether The SPIR generated by dpcpp has some extensions?

It depends on the DPC++ features used by the application, but in most cases the answer is yes.

[malixian]

thx very much @bader .

[malixian]

BTW, if my amd-opencl supports SPIR-V and the SPIR-V generated by DPC++ don't have extensioncan, it's work?

[bader]

BTW, if my amd-opencl supports SPIR-V and the SPIR-V generated by DPC++ don't have extensioncan, it's work?

This setup should work.

[malixian]

Can you give me some code that can be compiled into spirv without extensions? thank you very much.

[bader]

For example:

#include <CL/sycl.hpp>

int main() {
  // Creating buffer of 4 ints to be used inside the kernel code
  cl::sycl::buffer<cl::sycl::cl_int, 1> Buffer(4);

  // Creating SYCL queue
  cl::sycl::queue Queue;

  // Size of index space for kernel
  cl::sycl::range<1> NumOfWorkItems{Buffer.get_count()};

  // Submitting command group(work) to queue
  Queue.submit([&](cl::sycl::handler &cgh) {
    // Getting write only access to the buffer on a device
    auto Accessor = Buffer.get_access<cl::sycl::access::mode::write>(cgh);
    // Executing kernel
    cgh.parallel_for<class FillBuffer>(
        NumOfWorkItems, [=](cl::sycl::id<1> WIid) {
          // Fill buffer with indexes
          Accessor[WIid] = (cl::sycl::cl_int)WIid.get(0);
        });
  });

  // Getting read only access to the buffer on the host.
  // Implicit barrier waiting for queue to complete the work.
  const auto HostAccessor = Buffer.get_access<cl::sycl::access::mode::read>();

  // Check the results
  bool MismatchFound = false;
  for (size_t I = 0; I < Buffer.get_count(); ++I) {
    if (HostAccessor[I] != I) {
      std::cout << "The result is incorrect for element: " << I
                << " , expected: " << I << " , got: " << HostAccessor[I]
                << std::endl;
      MismatchFound = true;
    }
  }

  if (!MismatchFound) {
    std::cout << "The results are correct!" << std::endl;
  }

  return MismatchFound;
}

[malixian]

Thank you so much for your generous help @bader !