diff --git a/sycl/source/detail/property_set_io.hpp b/sycl/source/detail/property_set_io.hpp index a546a1d849022..bef16eb3796a0 100644 --- a/sycl/source/detail/property_set_io.hpp +++ b/sycl/source/detail/property_set_io.hpp @@ -14,7 +14,10 @@ #include "detail/base64.hpp" #include "sycl/exception.hpp" -#include +#include +#include +#include +#include namespace sycl { inline namespace _V1 { @@ -237,7 +240,54 @@ class PropertyValue { } Val; }; -using PropertySet = std::unordered_map; +// 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 class InsertionOrderedMap { +public: + using key_type = std::string; + using value_type = std::pair; + using StorageT = std::vector; + 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; /// A registry of property sets. Maps a property set name to its /// content. @@ -245,7 +295,7 @@ using PropertySet = std::unordered_map; /// The order of keys is preserved and corresponds to the order of insertion. class PropertySetRegistry { public: - using MapTy = std::unordered_map; + using MapTy = InsertionOrderedMap; // SYCLBIN specific property sets. static constexpr char SYCLBIN_GLOBAL_METADATA[] = "SYCLBIN/global metadata"; diff --git a/sycl/test-e2e/SYCLBIN/Inputs/spec_const_collision_kernel.cpp b/sycl/test-e2e/SYCLBIN/Inputs/spec_const_collision_kernel.cpp new file mode 100644 index 0000000000000..0895ad9786cc9 --- /dev/null +++ b/sycl/test-e2e/SYCLBIN/Inputs/spec_const_collision_kernel.cpp @@ -0,0 +1,16 @@ +#include + +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 SC_A{256}; +inline constexpr sycl::specialization_id 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(); + out[1] = kh.get_specialization_constant(); +} diff --git a/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp b/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp new file mode 100644 index 0000000000000..8d3d46991c734 --- /dev/null +++ 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(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 +#include +#include +#include + +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +inline constexpr sycl::specialization_id SC_A{256}; +inline constexpr sycl::specialization_id 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 resolveWithSCA(sycl::queue &Q, const char *Path, + int ValueA) { + const sycl::context Ctx = Q.get_context(); + auto KBInput = syclexp::get_kernel_bundle( + Ctx, std::string{Path}); + KBInput.set_specialization_constant(ValueA); + KBInput.set_specialization_constant(DefaultB); + auto KBExe = sycl::build(KBInput); + return {KBExe.get_specialization_constant(), + KBExe.get_specialization_constant()}; +} + +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; +} diff --git a/sycl/unittests/SYCL2020/CMakeLists.txt b/sycl/unittests/SYCL2020/CMakeLists.txt index 940450e8dec66..0cac763c423b2 100644 --- a/sycl/unittests/SYCL2020/CMakeLists.txt +++ b/sycl/unittests/SYCL2020/CMakeLists.txt @@ -18,5 +18,6 @@ add_sycl_unittest(SYCL2020Tests OBJECT SYCLBINSerializeMulti.cpp SYCLBINSerializeOverrides.cpp SYCLBINSerializeSpecConst.cpp + PropertySetIOOrder.cpp ) diff --git a/sycl/unittests/SYCL2020/PropertySetIOOrder.cpp b/sycl/unittests/SYCL2020/PropertySetIOOrder.cpp new file mode 100644 index 0000000000000..aff96ce5f6537 --- /dev/null +++ b/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 + +#include + +#include + +#include +#include +#include + +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: []\n=|\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 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 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 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]); +}