[SYCL] Preserve property order in PropertySetRegistry

sycl/source/detail/property_set_io.hpp

@@ -14,7 +14,10 @@
 #include "detail/base64.hpp"
 #include "sycl/exception.hpp"
 
-#include <unordered_map>
+#include <string>
+#include <string_view>
+#include <utility>
+#include <vector>
 
 namespace sycl {
 inline namespace _V1 {
@@ -237,15 +240,62 @@ class PropertyValue {
   } Val;
 };
 
-using PropertySet = std::unordered_map<std::string, PropertyValue>;
+// Insertion-ordered map. The serialized format is order-sensitive (spec-const
+// default-value blob is laid out in descriptor iteration order), so a hash
+// container would mismap blob offsets (CMPLRLLVM-77316). Mirrors the MapVector
+// in upstream PropertySetIO.h; linear lookup, sets are small.
+template <typename ValueT> class InsertionOrderedMap {
+public:
+  using key_type = std::string;
+  using value_type = std::pair<key_type, ValueT>;
+  using StorageT = std::vector<value_type>;
+  using iterator = typename StorageT::iterator;
+  using const_iterator = typename StorageT::const_iterator;
+
+  iterator begin() { return Storage.begin(); }
+  iterator end() { return Storage.end(); }
+  const_iterator begin() const { return Storage.begin(); }
+  const_iterator end() const { return Storage.end(); }
+
+  bool empty() const { return Storage.empty(); }
+  size_t size() const { return Storage.size(); }
+  void clear() { Storage.clear(); }
+
+  // Returns the value for Key, appending a new entry if absent.
+  ValueT &operator[](std::string_view Key) {
+    iterator It = find(Key);
+    if (It != end())
+      return It->second;
+    Storage.emplace_back(key_type{Key}, ValueT{});
+    return Storage.back().second;
+  }
+
+  iterator find(std::string_view Key) {
+    for (iterator It = begin(); It != end(); ++It)
+      if (It->first == Key)
+        return It;
+    return end();
+  }
+  const_iterator find(std::string_view Key) const {
+    for (const_iterator It = begin(); It != end(); ++It)
+      if (It->first == Key)
+        return It;
+    return end();
+  }
+
+private:
+  StorageT Storage;
+};
+
+using PropertySet = InsertionOrderedMap<PropertyValue>;
 
 /// A registry of property sets. Maps a property set name to its
 /// content.
 ///
 /// The order of keys is preserved and corresponds to the order of insertion.
 class PropertySetRegistry {
 public:
-  using MapTy = std::unordered_map<std::string, PropertySet>;
+  using MapTy = InsertionOrderedMap<PropertySet>;
 
   // SYCLBIN specific property sets.
   static constexpr char SYCLBIN_GLOBAL_METADATA[] = "SYCLBIN/global metadata";

sycl/test-e2e/SYCLBIN/Inputs/spec_const_collision_kernel.cpp

@@ -0,0 +1,16 @@
+#include <sycl/sycl.hpp>
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+// Kernel reads both spec constants so both are live; the CMPLRLLVM-77316
+// collision only triggers when both are referenced.
+inline constexpr sycl::specialization_id<int> SC_A{256};
+inline constexpr sycl::specialization_id<int> SC_B{1024};
+
+extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((
+    syclexp::nd_range_kernel<1>)) void spec_const_collision(int *out,
+                                                            sycl::kernel_handler
+                                                                kh) {
+  out[0] = kh.get_specialization_constant<SC_A>();
+  out[1] = kh.get_specialization_constant<SC_B>();
+}

sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp

@@ -0,0 +1,84 @@
+// UNSUPPORTED: cuda, hip
+// UNSUPPORTED-INTENDED: CUDA and HIP targets produce only native device
+// binaries and can therefore not produce input-state SYCLBIN files.
+
+// -- Regression test for CMPLRLLVM-77316: on the -fsyclbin=input path,
+// -- set_specialization_constant<SC_A>(v) must take effect even when v equals
+// -- another referenced constant's default. Setting SC_A=1024 (== SC_B default)
+// -- once dropped SC_A to its own default 256; it must resolve to 1024. Checked
+// -- host-side after build(), which exercises the same blob-offset resolution.
+
+// RUN: %clangxx --offload-new-driver -fsyclbin=input %{sycl_target_opts} %S/Inputs/spec_const_collision_kernel.cpp -o %t.syclbin
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out %t.syclbin
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/syclbin_kernel_bundle.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/specialization_id.hpp>
+
+#include <cassert>
+#include <iostream>
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+inline constexpr sycl::specialization_id<int> SC_A{256};
+inline constexpr sycl::specialization_id<int> SC_B{1024};
+
+static constexpr int DefaultA = 256;
+static constexpr int DefaultB = 1024;
+
+// Sets SC_A=ValueA, SC_B=DefaultB on the input bundle, builds, returns the
+// executable bundle's resolved (SC_A, SC_B).
+static std::pair<int, int> resolveWithSCA(sycl::queue &Q, const char *Path,
+                                          int ValueA) {
+  const sycl::context Ctx = Q.get_context();
+  auto KBInput = syclexp::get_kernel_bundle<sycl::bundle_state::input>(
+      Ctx, std::string{Path});
+  KBInput.set_specialization_constant<SC_A>(ValueA);
+  KBInput.set_specialization_constant<SC_B>(DefaultB);
+  auto KBExe = sycl::build(KBInput);
+  return {KBExe.get_specialization_constant<SC_A>(),
+          KBExe.get_specialization_constant<SC_B>()};
+}
+
+int main(int argc, char **argv) {
+  assert(argc == 2);
+  sycl::queue Q;
+  int Failed = 0;
+
+  // Regression case: SC_A set to SC_B's default must stay 1024, not drop to
+  // 256.
+  {
+    auto [A, B] = resolveWithSCA(Q, argv[1], DefaultB);
+    std::cout << "SC_A=1024 (== SC_B default): A=" << A << " B=" << B << "\n";
+    if (A != DefaultB || B != DefaultB) {
+      std::cout << "FAIL: expected A=1024 B=1024 (CMPLRLLVM-77316).\n";
+      ++Failed;
+    }
+  }
+
+  // Control: value colliding with nothing.
+  {
+    auto [A, B] = resolveWithSCA(Q, argv[1], 777);
+    std::cout << "SC_A=777: A=" << A << " B=" << B << "\n";
+    if (A != 777 || B != DefaultB) {
+      std::cout << "FAIL: expected A=777 B=1024\n";
+      ++Failed;
+    }
+  }
+
+  // Control: SC_A set to its own default.
+  {
+    auto [A, B] = resolveWithSCA(Q, argv[1], DefaultA);
+    std::cout << "SC_A=256 (own default): A=" << A << " B=" << B << "\n";
+    if (A != DefaultA || B != DefaultB) {
+      std::cout << "FAIL: expected A=256 B=1024\n";
+      ++Failed;
+    }
+  }
+
+  if (!Failed)
+    std::cout << "OK\n";
+  return Failed;
+}

sycl/unittests/SYCL2020/CMakeLists.txt

@@ -18,5 +18,6 @@ add_sycl_unittest(SYCL2020Tests OBJECT
   SYCLBINSerializeMulti.cpp
   SYCLBINSerializeOverrides.cpp
   SYCLBINSerializeSpecConst.cpp
+  PropertySetIOOrder.cpp
 )
 

sycl/unittests/SYCL2020/PropertySetIOOrder.cpp

@@ -0,0 +1,100 @@
+//==-- PropertySetIOOrder.cpp - property iteration order unit test --------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// Regression test for CMPLRLLVM-77316: PropertySetRegistry::read must yield
+// properties in insertion order. A hash container mismaps spec-const blob
+// offsets on the SYCLBIN path, silently dropping set_specialization_constant.
+
+#include <sycl/sycl.hpp>
+
+#include <detail/property_set_io.hpp>
+
+#include <gtest/gtest.h>
+
+#include <sstream>
+#include <string>
+#include <vector>
+
+using namespace sycl::detail;
+
+// Neither alphabetical nor libstdc++ hash order matches this sequence.
+static constexpr const char *ExpectedOrder[] = {"beta", "alpha", "mid"};
+
+static std::string makeBlob() {
+  // Format: [<category>]\n<name>=<type>|<value>\n; type 1 == UINT32.
+  std::ostringstream OS;
+  OS << "[SYCL/specialization constants]\n";
+  OS << "beta=1|10\n";
+  OS << "alpha=1|20\n";
+  OS << "mid=1|30\n";
+  return OS.str();
+}
+
+// read() preserves the blob's property order.
+TEST(PropertySetIOOrder, ReadPreservesInsertionOrder) {
+  std::string Blob = makeBlob();
+  auto Reg = PropertySetRegistry::read(Blob);
+  ASSERT_NE(Reg, nullptr);
+
+  auto SetIt = Reg->getPropSets().find("SYCL/specialization constants");
+  ASSERT_NE(SetIt, Reg->getPropSets().end());
+
+  std::vector<std::string> Names;
+  for (const auto &Prop : SetIt->second)
+    Names.push_back(Prop.first);
+
+  ASSERT_EQ(Names.size(), 3u);
+  EXPECT_EQ(Names[0], ExpectedOrder[0]);
+  EXPECT_EQ(Names[1], ExpectedOrder[1]);
+  EXPECT_EQ(Names[2], ExpectedOrder[2]);
+}
+
+// write() then read() preserves order.
+TEST(PropertySetIOOrder, WriteReadRoundTripPreservesOrder) {
+  std::string Blob = makeBlob();
+  auto Reg = PropertySetRegistry::read(Blob);
+  ASSERT_NE(Reg, nullptr);
+
+  std::ostringstream OS;
+  Reg->write(OS);
+  std::string Serialized = OS.str();
+
+  auto Reg2 = PropertySetRegistry::read(Serialized);
+  ASSERT_NE(Reg2, nullptr);
+
+  auto SetIt = Reg2->getPropSets().find("SYCL/specialization constants");
+  ASSERT_NE(SetIt, Reg2->getPropSets().end());
+
+  std::vector<std::string> Names;
+  for (const auto &Prop : SetIt->second)
+    Names.push_back(Prop.first);
+
+  ASSERT_EQ(Names.size(), 3u);
+  EXPECT_EQ(Names[0], ExpectedOrder[0]);
+  EXPECT_EQ(Names[1], ExpectedOrder[1]);
+  EXPECT_EQ(Names[2], ExpectedOrder[2]);
+}
+
+// operator[] appends new keys at the end.
+TEST(PropertySetIOOrder, InsertionAppendsAtEnd) {
+  PropertySetRegistry Reg;
+  Reg.add("SYCL/specialization constants", "beta", uint32_t{10});
+  Reg.add("SYCL/specialization constants", "alpha", uint32_t{20});
+  Reg.add("SYCL/specialization constants", "mid", uint32_t{30});
+
+  auto SetIt = Reg.getPropSets().find("SYCL/specialization constants");
+  ASSERT_NE(SetIt, Reg.getPropSets().end());
+
+  std::vector<std::string> Names;
+  for (const auto &Prop : SetIt->second)
+    Names.push_back(Prop.first);
+
+  ASSERT_EQ(Names.size(), 3u);
+  EXPECT_EQ(Names[0], ExpectedOrder[0]);
+  EXPECT_EQ(Names[1], ExpectedOrder[1]);
+  EXPECT_EQ(Names[2], ExpectedOrder[2]);
+}