[SYCL] Add initial support of AOT compilation for multiple devices

This patch adds handling of "target" device image field both in SYCL
runtime. This enables basic support of multiple device images.

The natural scenario goes as follows:
- a kernel needs to be run on a particular device;
- if it was AOT-compiled as part of a binary image for this device,
load the image;
- if the appropriate AOT-compiled image was not found, a SPIR-V image
is chosen and JIT-compiled instead.

This initial mechanism provides limited functionality, simply picking
the first image that has appropriate binary target.

Signed-off-by: Sergey Semenov <[email protected]>
Signed-off-by: Artem Gindinson <[email protected]>

Author: Artem Gindinson

sycl/include/CL/sycl/detail/pi.h

@@ -56,6 +56,7 @@ typedef enum {
   PI_INVALID_CONTEXT = CL_INVALID_CONTEXT,
   PI_INVALID_PLATFORM = CL_INVALID_PLATFORM,
   PI_INVALID_DEVICE = CL_INVALID_DEVICE,
+  PI_INVALID_BINARY = CL_INVALID_BINARY,
   PI_MISALIGNED_SUB_BUFFER_OFFSET = CL_MISALIGNED_SUB_BUFFER_OFFSET,
   PI_OUT_OF_HOST_MEMORY = CL_OUT_OF_HOST_MEMORY
 } _pi_result;
@@ -245,9 +246,20 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4;
 /// Target identification strings for
 /// pi_device_binary_struct.DeviceTargetSpec
 ///
+/// A device type represented by a particular target
+/// triple requires specific binary images. We need
+/// to map the image type onto the device target triple
+///
 #define PI_DEVICE_BINARY_TARGET_UNKNOWN "<unknown>"
+/// SPIR-V 32-bit image <-> 32-bit OpenCL device
 #define PI_DEVICE_BINARY_TARGET_SPIRV32 "spir"
-#define PI_DEVICE_BINARY_TARGET_SPIRV64 "spir64";
+/// SPIR-V 64-bit image <-> 64-bit OpenCL device
+#define PI_DEVICE_BINARY_TARGET_SPIRV64 "spir64"
+/// Device-specific binary images produced from SPIR-V 64-bit <->
+/// various triples for specific 64-bit OpenCL devices
+#define PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64 "spir64_x86_64"
+#define PI_DEVICE_BINARY_TARGET_SPIRV64_GEN "spir64_gen"
+#define PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA "spir64_fpga"
 
 /// This struct is a record of the device binary information. If the Kind field
 /// denotes a portable binary type (SPIRV or LLVMIR), the DeviceTargetSpec field
@@ -264,6 +276,13 @@ struct pi_device_binary_struct {
   /// format of the binary data - SPIRV, LLVMIR bitcode,...
   uint8_t Format;
   /// null-terminated string representation of the device's target architecture
+  /// which holds one of:
+  /// PI_DEVICE_BINARY_TARGET_UNKNOWN - unknown
+  /// PI_DEVICE_BINARY_TARGET_SPIRV32 - general value for 32-bit OpenCL devices
+  /// PI_DEVICE_BINARY_TARGET_SPIRV64 - general value for 64-bit OpenCL devices
+  /// PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64 - 64-bit OpenCL CPU device
+  /// PI_DEVICE_BINARY_TARGET_SPIRV64_GEN - GEN GPU device (64-bit OpenCL)
+  /// PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA - 64-bit OpenCL FPGA device
   const char *DeviceTargetSpec;
   /// a null-terminated string; target- and compiler-specific options
   /// which are suggested to use to "build" program at runtime
@@ -405,7 +424,7 @@ pi_result piDevicePartition(
 /// and the IR characteristics.
 ///
 pi_result piextDeviceSelectBinary(
-  pi_device           device, // TODO: does this need to be context?
+  pi_device           device,
   pi_device_binary *  binaries,
   pi_uint32           num_binaries,
   pi_device_binary *  selected_binary);

sycl/source/detail/pi_opencl.cpp

@@ -59,13 +59,16 @@ pi_result OCL(piDevicesGet)(pi_platform      platform,
   return cast<pi_result>(result);
 }
 
-pi_result OCL(piextDeviceSelectBinary)(
-  pi_device           device, // TODO: does this need to be context?
-  pi_device_binary *  images,
-  pi_uint32           num_images,
-  pi_device_binary *  selected_image) {
-
-  // TODO dummy implementation.
+pi_result OCL(piextDeviceSelectBinary)(pi_device device,
+                                       pi_device_binary *images,
+                                       pi_uint32 num_images,
+                                       pi_device_binary *selected_image) {
+
+  // TODO: this is a bare-bones implementation for choosing a device image
+  // that would be compatible with the targeted device. An AOT-compiled
+  // image is preferred over SPIRV for known devices (i.e. Intel devices)
+  // The implementation makes no effort to differentiate between multiple images
+  // for the given device, and simply picks the first one compatible
   // Real implementaion will use the same mechanism OpenCL ICD dispatcher
   // uses. Somthing like:
   //   PI_VALIDATE_HANDLE_RETURN_HANDLE(ctx, PI_INVALID_CONTEXT);
@@ -74,8 +77,56 @@ pi_result OCL(piextDeviceSelectBinary)(
   // where context->dispatch is set to the dispatch table provided by PI
   // plugin for platform/device the ctx was created for.
 
-  *selected_image = num_images > 0 ? images[0] : nullptr;
-  return PI_SUCCESS;
+  // Choose the binary target for the provided device
+  const char *image_target = nullptr;
+  // Get the type of the device
+  cl_device_type device_type;
+  cl_int ret_err = clGetDeviceInfo(cast<cl_device_id>(device), CL_DEVICE_TYPE,
+                                   sizeof(cl_device_type), &device_type, nullptr);
+  if (ret_err != CL_SUCCESS) {
+    *selected_image = nullptr;
+    return cast<pi_result>(ret_err);
+  }
+
+  switch (device_type) {
+  // TODO: Factor out vendor specifics into a separate source
+  // E.g. sycl/source/detail/vendor/intel/detail/pi_opencl.cpp?
+
+  // We'll attempt to find an image that was AOT-compiled
+  // from a SPIR-V image into an image specific for:
+
+  case CL_DEVICE_TYPE_CPU: // OpenCL 64-bit CPU
+    image_target = PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64;
+    break;
+  case CL_DEVICE_TYPE_GPU: // OpenCL 64-bit GEN GPU
+    image_target = PI_DEVICE_BINARY_TARGET_SPIRV64_GEN;
+    break;
+  case CL_DEVICE_TYPE_ACCELERATOR: // OpenCL 64-bit FPGA
+    image_target = PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA;
+    break;
+  default:
+    // Otherwise, we'll attempt to find and JIT-compile
+    // a device-independent SPIR-V image
+    image_target = PI_DEVICE_BINARY_TARGET_SPIRV64;
+    break;
+  }
+
+  // Find the appropriate device image, fallback to spirv if not found
+  pi_device_binary fallback = nullptr;
+  for (size_t i = 0; i < num_images; ++i) {
+    if (strcmp(images[i]->DeviceTargetSpec, image_target) == 0) {
+      *selected_image = images[i];
+      return PI_SUCCESS;
+    }
+    if (strcmp(images[i]->DeviceTargetSpec, PI_DEVICE_BINARY_TARGET_SPIRV64) ==
+        0)
+      fallback = images[i];
+  }
+  // Points to a spirv image, if such indeed was found
+  if ((*selected_image = fallback))
+    return PI_SUCCESS;
+  // No image can be loaded for the given device
+  return PI_INVALID_BINARY;
 }
 
 pi_result OCL(piQueueCreate)(pi_context context, pi_device device,
@@ -290,7 +341,7 @@ _PI_CL(piDeviceRetain,       clRetainDevice)
 _PI_CL(piDeviceRelease,      clReleaseDevice)
 _PI_CL(piextDeviceSelectBinary, OCL(piextDeviceSelectBinary))
 _PI_CL(piextGetDeviceFunctionPointer, OCL(piextGetDeviceFunctionPointer))
-  // Context
+// Context
 _PI_CL(piContextCreate,     clCreateContext)
 _PI_CL(piContextGetInfo,    clGetContextInfo)
 _PI_CL(piContextRetain,     clRetainContext)

sycl/source/detail/program_manager/program_manager.cpp

@@ -263,6 +263,7 @@ RT::PiProgram ProgramManager::loadProgram(OSModuleHandle M,
               << getRawSyclObjImpl(Context) << ")\n";
   }
 
+  const RT::PiContext &Ctx = getRawSyclObjImpl(Context)->getHandleRef();
   DeviceImage *Img = nullptr;
   bool UseKernelSpv = false;
   const std::string UseSpvEnv("SYCL_USE_KERNEL_SPV");
@@ -309,6 +310,10 @@ RT::PiProgram ProgramManager::loadProgram(OSModuleHandle M,
       std::cerr << "loaded device image from " << Fname << "\n";
     }
   } else {
+    // TODO: There may be cases with cl::sycl::program class usage in source code
+    // that will result in a multi-device context. This case needs to be handled
+    // here or at the program_impl class level
+
     // Take all device images in module M and ask the native runtime under the
     // given context to choose one it prefers.
     auto ImgIt = m_DeviceImages.find(M);
@@ -318,8 +323,8 @@ RT::PiProgram ProgramManager::loadProgram(OSModuleHandle M,
     }
     std::vector<DeviceImage *> *Imgs = (ImgIt->second).get();
 
-    PI_CALL(RT::piextDeviceSelectBinary(
-      0, Imgs->data(), (cl_uint)Imgs->size(), &Img));
+    PI_CALL(RT::piextDeviceSelectBinary(getFirstDevice(Ctx), Imgs->data(),
+                                        (cl_uint)Imgs->size(), &Img));
 
     if (DbgProgMgr > 0) {
       std::cerr << "available device images:\n";
@@ -400,7 +405,6 @@ RT::PiProgram ProgramManager::loadProgram(OSModuleHandle M,
   // Load the selected image
   if (!is_device_binary_type_supported(Context, Format))
     throw feature_not_supported("Online compilation is not supported in this context");
-  const RT::PiContext &Ctx = getRawSyclObjImpl(Context)->getHandleRef();
   RT::PiProgram Res = nullptr;
   Res = Format == PI_DEVICE_BINARY_TYPE_SPIRV
             ? createSpirvProgram(Ctx, Img->BinaryStart, ImgSize)

sycl/test/aot/accelerator.cpp

@@ -0,0 +1,82 @@
+// REQUIRES: aoc
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga-unknown-linux-sycldevice %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+//==----- accelerator.cpp - AOT compilation for fpga devices using aoc  ------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===------------------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+
+#include <array>
+#include <iostream>
+
+constexpr cl::sycl::access::mode sycl_read = cl::sycl::access::mode::read;
+constexpr cl::sycl::access::mode sycl_write = cl::sycl::access::mode::write;
+
+template <typename T>
+class SimpleVadd;
+
+template <typename T, size_t N>
+void simple_vadd(const std::array<T, N>& VA, const std::array<T, N>& VB,
+                 std::array<T, N>& VC) {
+  cl::sycl::queue deviceQueue([](cl::sycl::exception_list ExceptionList) {
+      for (cl::sycl::exception_ptr_class ExceptionPtr : ExceptionList) {
+        try {
+          std::rethrow_exception(ExceptionPtr);
+        } catch (cl::sycl::exception &E) {
+          std::cerr << E.what();
+        } catch (...) {
+          std::cerr << "Unknown async exception was caught." << std::endl;
+        }
+      }
+    });
+
+  cl::sycl::range<1> numOfItems{N};
+  cl::sycl::buffer<T, 1> bufferA(VA.data(), numOfItems);
+  cl::sycl::buffer<T, 1> bufferB(VB.data(), numOfItems);
+  cl::sycl::buffer<T, 1> bufferC(VC.data(), numOfItems);
+
+  deviceQueue.submit([&](cl::sycl::handler& cgh) {
+    auto accessorA = bufferA.template get_access<sycl_read>(cgh);
+    auto accessorB = bufferB.template get_access<sycl_read>(cgh);
+    auto accessorC = bufferC.template get_access<sycl_write>(cgh);
+
+    cgh.parallel_for<class SimpleVadd<T>>(numOfItems,
+    [=](cl::sycl::id<1> wiID) {
+        accessorC[wiID] = accessorA[wiID] + accessorB[wiID];
+    });
+  });
+
+  deviceQueue.wait_and_throw();
+}
+
+int main() {
+  const size_t array_size = 4;
+  std::array<cl::sycl::cl_int, array_size> A = {{1, 2, 3, 4}},
+                                           B = {{1, 2, 3, 4}}, C;
+  std::array<cl::sycl::cl_float, array_size> D = {{1.f, 2.f, 3.f, 4.f}},
+                                             E = {{1.f, 2.f, 3.f, 4.f}}, F;
+  simple_vadd(A, B, C);
+  simple_vadd(D, E, F);
+  for (unsigned int i = 0; i < array_size; i++) {
+    if (C[i] != A[i] + B[i]) {
+      std::cout << "The results are incorrect (element " << i << " is " << C[i]
+                << "!\n";
+      return 1;
+    }
+    if (F[i] != D[i] + E[i]) {
+      std::cout << "The results are incorrect (element " << i << " is " << F[i]
+                << "!\n";
+      return 1;
+    }
+  }
+  std::cout << "The results are correct!\n";
+  return 0;
+}

sycl/test/aot/cpu.cpp

@@ -0,0 +1,82 @@
+// REQUIRES: ioc64
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64-unknown-linux-sycldevice %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+
+//==----- cpu.cpp - AOT compilation for cpu devices using ioc64  -------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===------------------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+
+#include <array>
+#include <iostream>
+
+constexpr cl::sycl::access::mode sycl_read = cl::sycl::access::mode::read;
+constexpr cl::sycl::access::mode sycl_write = cl::sycl::access::mode::write;
+
+template <typename T>
+class SimpleVadd;
+
+template <typename T, size_t N>
+void simple_vadd(const std::array<T, N>& VA, const std::array<T, N>& VB,
+                 std::array<T, N>& VC) {
+  cl::sycl::queue deviceQueue([](cl::sycl::exception_list ExceptionList) {
+      for (cl::sycl::exception_ptr_class ExceptionPtr : ExceptionList) {
+        try {
+          std::rethrow_exception(ExceptionPtr);
+        } catch (cl::sycl::exception &E) {
+          std::cerr << E.what();
+        } catch (...) {
+          std::cerr << "Unknown async exception was caught." << std::endl;
+        }
+      }
+    });
+
+  cl::sycl::range<1> numOfItems{N};
+  cl::sycl::buffer<T, 1> bufferA(VA.data(), numOfItems);
+  cl::sycl::buffer<T, 1> bufferB(VB.data(), numOfItems);
+  cl::sycl::buffer<T, 1> bufferC(VC.data(), numOfItems);
+
+  deviceQueue.submit([&](cl::sycl::handler& cgh) {
+    auto accessorA = bufferA.template get_access<sycl_read>(cgh);
+    auto accessorB = bufferB.template get_access<sycl_read>(cgh);
+    auto accessorC = bufferC.template get_access<sycl_write>(cgh);
+
+    cgh.parallel_for<class SimpleVadd<T>>(numOfItems,
+    [=](cl::sycl::id<1> wiID) {
+        accessorC[wiID] = accessorA[wiID] + accessorB[wiID];
+    });
+  });
+
+  deviceQueue.wait_and_throw();
+}
+
+int main() {
+  const size_t array_size = 4;
+  std::array<cl::sycl::cl_int, array_size> A = {{1, 2, 3, 4}},
+                                           B = {{1, 2, 3, 4}}, C;
+  std::array<cl::sycl::cl_float, array_size> D = {{1.f, 2.f, 3.f, 4.f}},
+                                             E = {{1.f, 2.f, 3.f, 4.f}}, F;
+  simple_vadd(A, B, C);
+  simple_vadd(D, E, F);
+  for (unsigned int i = 0; i < array_size; i++) {
+    if (C[i] != A[i] + B[i]) {
+      std::cout << "The results are incorrect (element " << i << " is " << C[i]
+                << "!\n";
+      return 1;
+    }
+    if (F[i] != D[i] + E[i]) {
+      std::cout << "The results are incorrect (element " << i << " is " << F[i]
+                << "!\n";
+      return 1;
+    }
+  }
+  std::cout << "The results are correct!\n";
+  return 0;
+}