[SYCL] Add convenience C++ wrappers for PI binary image hierarchy objects.

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Author: kbobrovs

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

@@ -541,15 +541,15 @@ using _pi_offload_entry = _pi_offload_entry_struct *;
 
 // A type of a binary image property.
 typedef enum {
-  PI_PROP_TYPE_UNKNOWN,
-  PI_PROP_TYPE_UINT32, // 32-bit integer
-  PI_PROP_TYPE_STRING  // null-terminated string
-} _pi_property_type;
+  PI_PROPERTY_TYPE_UNKNOWN,
+  PI_PROPERTY_TYPE_UINT32, // 32-bit integer
+  PI_PROPERTY_TYPE_STRING  // null-terminated string
+} pi_property_type;
 
 // Device binary image property.
 // If the type size of the property value is fixed and is no greater than
 // 64 bits, then ValAddr is 0 and the value is stored in the ValSize field.
-// Example - PI_PROP_TYPE_UINT32, which is 32-bit
+// Example - PI_PROPERTY_TYPE_UINT32, which is 32-bit
 struct _pi_device_binary_property_struct {
   char *Name;       // null-terminated property name
   void *ValAddr;    // address of property value

sycl/include/CL/sycl/detail/pi.hpp

@@ -165,6 +165,124 @@ void printArgs(Arg0 arg0, Args... args) {
   print(arg0);
   pi::printArgs(std::forward<Args>(args)...);
 }
+
+// C++ wrapper over the _pi_device_binary_property_struct structure.
+class DeviceBinaryProperty {
+public:
+  DeviceBinaryProperty(const _pi_device_binary_property_struct *Prop)
+      : Prop(Prop) {}
+
+  pi_uint32 asUint32() const;
+  const char *asCString() const;
+
+protected:
+  friend std::ostream &operator<<(std::ostream &Out,
+                                  const DeviceBinaryProperty &P);
+  const _pi_device_binary_property_struct *Prop;
+};
+
+std::ostream &operator<<(std::ostream &Out, const DeviceBinaryProperty &P);
+
+// C++ convenience wrapper over the pi_device_binary_struct structure.
+class DeviceBinaryImage {
+public:
+  // Represents a range of properties to enable iteration over them.
+  // Implements the standard C++ STL input iterator interface.
+  class PropertyRange {
+  public:
+    using ValTy = std::remove_pointer<pi_device_binary_property>::type;
+
+    class ConstIterator
+        : public std::iterator<std::input_iterator_tag, // iterator_category
+                               ValTy,                   // value_type
+                               ptrdiff_t,               // difference_type
+                               const pi_device_binary_property, // pointer
+                               pi_device_binary_property>       // reference
+    {
+      pi_device_binary_property Cur;
+
+    public:
+      ConstIterator(pi_device_binary_property Cur = nullptr) : Cur(Cur) {}
+      ConstIterator &operator++() {
+        Cur++;
+        return *this;
+      }
+      ConstIterator operator++(int) {
+        ConstIterator Ret = *this;
+        ++(*this);
+        return Ret;
+      }
+      bool operator==(ConstIterator Other) const { return Cur == Other.Cur; }
+      bool operator!=(ConstIterator Other) const { return !(*this == Other); }
+      reference operator*() const { return Cur; }
+    };
+    ConstIterator begin() const { return ConstIterator(Begin); }
+    ConstIterator end() const { return ConstIterator(End); }
+    friend class DeviceBinaryImage;
+
+  private:
+    PropertyRange() : Begin(nullptr), End(nullptr) {}
+    // Searches for a property set with given name and constructs a
+    // PropertyRange spanning all its elements. If property set is not found,
+    // the range will span zero elements.
+    PropertyRange(pi_device_binary Bin, const char *PropSetName)
+        : PropertyRange() {
+      init(Bin, PropSetName);
+    };
+    void init(pi_device_binary Bin, const char *PropSetName);
+    pi_device_binary_property Begin;
+    pi_device_binary_property End;
+  };
+
+public:
+  DeviceBinaryImage(pi_device_binary Bin) { init(Bin); }
+  DeviceBinaryImage() : Bin(nullptr){};
+
+  virtual void print() const;
+  virtual void dump(std::ostream &Out) const;
+
+  size_t getSize() const {
+    assert(Bin && "binary image data not set");
+    return static_cast<size_t>(Bin->BinaryEnd - Bin->BinaryStart);
+  }
+
+  const char *getCompileOptions() const {
+    assert(Bin && "binary image data not set");
+    return Bin->CompileOptions;
+  }
+
+  const char *getLinkOptions() const {
+    assert(Bin && "binary image data not set");
+    return Bin->LinkOptions;
+  }
+
+  /// Returns the format of the binary image
+  pi::PiDeviceBinaryType getFormat() const {
+    assert(Bin && "binary image data not set");
+    return Format;
+  }
+
+  /// Gets the iterator range over specialization constants in this this binary
+  /// image. For each property pointed to by an iterator within the range, the
+  /// name of the property is the specializaion constant symbolic ID and the
+  /// value is 32-bit unsigned integer ID.
+  const PropertyRange &getSpecConstants() const { return SpecConstIDMap; }
+  virtual ~DeviceBinaryImage() {}
+
+protected:
+  void init(pi_device_binary Bin);
+  pi_device_binary get() const { return Bin; }
+
+  pi_device_binary Bin;
+  pi::PiDeviceBinaryType Format = PI_DEVICE_BINARY_TYPE_NONE;
+  DeviceBinaryImage::PropertyRange SpecConstIDMap;
+};
+
+/// Tries to determine the device binary image foramat. Returns
+/// PI_DEVICE_BINARY_TYPE_NONE if unsuccessful.
+PiDeviceBinaryType getBinaryImageFormat(const unsigned char *ImgData,
+                                        size_t ImgSize);
+
 } // namespace pi
 
 namespace RT = cl::sycl::detail::pi;

sycl/source/detail/pi.cpp

@@ -298,6 +298,146 @@ void assertion(bool Condition, const char *Message) {
     die(Message);
 }
 
+std::ostream &operator<<(std::ostream &Out, const DeviceBinaryProperty &P) {
+  switch (P.Prop->Type) {
+  case PI_PROPERTY_TYPE_UINT32:
+    Out << "[UINT32] ";
+    break;
+  case PI_PROPERTY_TYPE_STRING:
+    Out << "[String] ";
+    break;
+  default:
+    assert("unsupported property");
+    return Out;
+  }
+  Out << P.Prop->Name << "=";
+
+  switch (P.Prop->Type) {
+  case PI_PROPERTY_TYPE_UINT32:
+    Out << P.asUint32();
+    break;
+  case PI_PROPERTY_TYPE_STRING:
+    Out << P.asCString();
+    break;
+  default:
+    assert("unsupported property");
+    return Out;
+  }
+  return Out;
+}
+
+void DeviceBinaryImage::print() const {
+  std::cerr << "  --- Image " << Bin << "\n";
+  if (!Bin)
+    return;
+  std::cerr << "    Version  : " << (int)Bin->Version << "\n";
+  std::cerr << "    Kind     : " << (int)Bin->Kind << "\n";
+  std::cerr << "    Format   : " << (int)Bin->Format << "\n";
+  std::cerr << "    Target   : " << Bin->DeviceTargetSpec << "\n";
+  std::cerr << "    Bin size : "
+            << ((intptr_t)Bin->BinaryEnd - (intptr_t)Bin->BinaryStart) << "\n";
+  std::cerr << "    Compile options : "
+            << (Bin->CompileOptions ? Bin->CompileOptions : "NULL") << "\n";
+  std::cerr << "    Link options    : "
+            << (Bin->LinkOptions ? Bin->LinkOptions : "NULL") << "\n";
+  std::cerr << "    Entries  : ";
+  for (_pi_offload_entry EntriesIt = Bin->EntriesBegin;
+       EntriesIt != Bin->EntriesEnd; ++EntriesIt)
+    std::cerr << EntriesIt->name << " ";
+  std::cerr << "\n";
+  std::cerr << "    Properties [" << Bin->PropertySetsBegin << "-"
+            << Bin->PropertySetsEnd << "]:\n";
+
+  for (pi_device_binary_property_set PS = Bin->PropertySetsBegin;
+       PS != Bin->PropertySetsEnd; ++PS) {
+    std::cerr << "      Category " << PS->Name << " [" << PS->PropertiesBegin
+              << "-" << PS->PropertiesEnd << "]:\n";
+
+    for (pi_device_binary_property P = PS->PropertiesBegin;
+         P != PS->PropertiesEnd; ++P) {
+      std::cerr << "        " << DeviceBinaryProperty(P) << "\n";
+    }
+  }
+}
+
+void DeviceBinaryImage::dump(std::ostream &Out) const {
+  size_t ImgSize = getSize();
+  Out.write(reinterpret_cast<const char *>(Bin->BinaryStart), ImgSize);
+}
+
+static pi_uint32 asUint32(const void *Addr) {
+  assert(Addr && "Addr is NULL");
+  const auto *P = reinterpret_cast<const unsigned char *>(Addr);
+  return (*P) | (*(P + 1) << 8) | (*(P + 2) << 16) | (*(P + 3) << 24);
+}
+
+pi_uint32 DeviceBinaryProperty::asUint32() const {
+  assert(Prop->Type == PI_PROPERTY_TYPE_UINT32 && "property type mismatch");
+  // if type fits into the ValSize - it is used to store the property value
+  assert(Prop->ValAddr == nullptr && "primitive types must be stored inline");
+  return sycl::detail::pi::asUint32(&Prop->ValSize);
+}
+
+const char *DeviceBinaryProperty::asCString() const {
+  assert(Prop->Type == PI_PROPERTY_TYPE_STRING && "property type mismatch");
+  assert(Prop->ValSize > 0 && "property size mismatch");
+  return pi::cast<const char *>(Prop->ValAddr);
+}
+
+void DeviceBinaryImage::PropertyRange::init(pi_device_binary Bin,
+                                            const char *PropSetName) {
+  assert(!this->Begin && !this->End && "already initialized");
+  pi_device_binary_property_set PS = nullptr;
+
+  for (PS = Bin->PropertySetsBegin; PS != Bin->PropertySetsEnd; ++PS) {
+    assert(PS->Name && "nameless property set - bug in the offload wrapper?");
+    if (!strcmp(PropSetName, PS->Name))
+      break;
+  }
+  if (PS == Bin->PropertySetsEnd) {
+    Begin = End = nullptr;
+    return;
+  }
+  Begin = PS->PropertiesBegin;
+  End = Begin ? PS->PropertiesEnd : nullptr;
+}
+
+RT::PiDeviceBinaryType getBinaryImageFormat(const unsigned char *ImgData,
+                                            size_t ImgSize) {
+  struct {
+    RT::PiDeviceBinaryType Fmt;
+    const uint32_t Magic;
+  } Fmts[] = {{PI_DEVICE_BINARY_TYPE_SPIRV, 0x07230203},
+              {PI_DEVICE_BINARY_TYPE_LLVMIR_BITCODE, 0xDEC04342}};
+
+  if (ImgSize >= sizeof(Fmts[0].Magic)) {
+    std::remove_const<decltype(Fmts[0].Magic)>::type Hdr = 0;
+    std::copy(ImgData, ImgData + sizeof(Hdr), reinterpret_cast<char *>(&Hdr));
+
+    for (const auto &Fmt : Fmts) {
+      if (Hdr == Fmt.Magic)
+        return Fmt.Fmt;
+    }
+  }
+  return PI_DEVICE_BINARY_TYPE_NONE;
+}
+
+void DeviceBinaryImage::init(pi_device_binary Bin) {
+  this->Bin = Bin;
+  // If device binary image format wasn't set by its producer, then can't change
+  // now, because 'Bin' data is part of the executable image loaded into memory
+  // which can't be modified (easily).
+  // TODO clang driver + ClangOffloadWrapper can figure out the format and set
+  // it when invoking the offload wrapper job
+  Format = static_cast<pi::PiDeviceBinaryType>(Bin->Format);
+
+  if (Format == PI_DEVICE_BINARY_TYPE_NONE)
+    // try to determine the format; may remain "NONE"
+    Format = getBinaryImageFormat(Bin->BinaryStart, getSize());
+
+  SpecConstIDMap.init(Bin, PI_PROPERTY_SET_SPEC_CONST_MAP);
+}
+
 } // namespace pi
 } // namespace detail
 } // namespace sycl