From bbdef2d2f61bac7ba0bf761f62575166934a352b Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Wed, 5 Aug 2026 16:33:29 -0700 Subject: [PATCH 1/3] [SYCL] Preserve property order in PropertySetRegistry (CMPLRLLVM-77316) std::unordered_map's hash order mismapped spec-const blob offsets, silently dropping set_specialization_constant on the SYCLBIN input path. Use an insertion-ordered container. Add unit + E2E tests. Co-Authored-By: Claude Opus 4.8 (1M context) --- sycl/source/detail/property_set_io.hpp | 56 +++++++++- .../Inputs/spec_const_collision_kernel.cpp | 16 +++ .../SYCLBIN/spec_const_collision_input.cpp | 104 ++++++++++++++++++ sycl/unittests/SYCL2020/CMakeLists.txt | 1 + .../unittests/SYCL2020/PropertySetIOOrder.cpp | 100 +++++++++++++++++ 5 files changed, 274 insertions(+), 3 deletions(-) create mode 100644 sycl/test-e2e/SYCLBIN/Inputs/spec_const_collision_kernel.cpp create mode 100644 sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp create mode 100644 sycl/unittests/SYCL2020/PropertySetIOOrder.cpp 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..a4ac70e30a722 --- /dev/null +++ b/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp @@ -0,0 +1,104 @@ +// REQUIRES: aspect-usm_shared_allocations + +// 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 yield 1024. + +// 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 +#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, launches, +// returns (out[0], out[1]). +static std::pair runWithSCA(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); + sycl::kernel Kern = KBExe.ext_oneapi_get_kernel("spec_const_collision"); + + int *Out = sycl::malloc_shared(2, Q); + Out[0] = Out[1] = -1; + Q.submit([&](sycl::handler &CGH) { + CGH.use_kernel_bundle(KBExe); + CGH.set_args(Out); + CGH.parallel_for(sycl::nd_range<1>{{1}, {1}}, Kern); + }).wait_and_throw(); + + std::pair Result{Out[0], Out[1]}; + sycl::free(Out, Q); + return Result; +} + +int main(int argc, char **argv) { + assert(argc == 2); + sycl::queue Q; + + int Failed = 0; + + // The regression case: SC_A set to SC_B's default. Must see A == 1024. + { + auto [A, B] = runWithSCA(Q, argv[1], DefaultB); + std::cout << "SC_A=1024 (== SC_B default): A=" << A << " B=" << B << "\n"; + if (A != DefaultB) { + std::cout << "FAIL: SC_A was dropped to " << A << "; expected " + << DefaultB << " (CMPLRLLVM-77316).\n"; + ++Failed; + } + if (B != DefaultB) { + std::cout << "FAIL: SC_B = " << B << "; expected " << DefaultB << "\n"; + ++Failed; + } + } + + // Control: a value that collides with nothing. + { + auto [A, B] = runWithSCA(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] = runWithSCA(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]); +} From 0a562c77dd55b7f78326190d81a8221c7d17e18f Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Wed, 12 Aug 2026 10:46:01 -0700 Subject: [PATCH 2/3] fix test --- .../SYCLBIN/spec_const_collision_input.cpp | 51 ++++++------------- 1 file changed, 15 insertions(+), 36 deletions(-) diff --git a/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp b/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp index a4ac70e30a722..f8e3f836d0a6a 100644 --- a/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp +++ b/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp @@ -1,5 +1,3 @@ -// REQUIRES: aspect-usm_shared_allocations - // UNSUPPORTED: cuda, hip // UNSUPPORTED-INTENDED: CUDA and HIP targets produce only native device // binaries and can therefore not produce input-state SYCLBIN files. @@ -7,7 +5,8 @@ // -- 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 yield 1024. +// -- 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 @@ -17,7 +16,6 @@ #include #include #include -#include #include #include @@ -30,57 +28,38 @@ 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, launches, -// returns (out[0], out[1]). -static std::pair runWithSCA(sycl::queue &Q, const char *Path, - int ValueA) { +// 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); - sycl::kernel Kern = KBExe.ext_oneapi_get_kernel("spec_const_collision"); - - int *Out = sycl::malloc_shared(2, Q); - Out[0] = Out[1] = -1; - Q.submit([&](sycl::handler &CGH) { - CGH.use_kernel_bundle(KBExe); - CGH.set_args(Out); - CGH.parallel_for(sycl::nd_range<1>{{1}, {1}}, Kern); - }).wait_and_throw(); - - std::pair Result{Out[0], Out[1]}; - sycl::free(Out, Q); - return Result; + return {KBExe.get_specialization_constant(), + KBExe.get_specialization_constant()}; } int main(int argc, char **argv) { assert(argc == 2); sycl::queue Q; - int Failed = 0; - // The regression case: SC_A set to SC_B's default. Must see A == 1024. + // Regression case: SC_A set to SC_B's default must stay 1024, not drop to 256. { - auto [A, B] = runWithSCA(Q, argv[1], DefaultB); + auto [A, B] = resolveWithSCA(Q, argv[1], DefaultB); std::cout << "SC_A=1024 (== SC_B default): A=" << A << " B=" << B << "\n"; - if (A != DefaultB) { - std::cout << "FAIL: SC_A was dropped to " << A << "; expected " - << DefaultB << " (CMPLRLLVM-77316).\n"; - ++Failed; - } - if (B != DefaultB) { - std::cout << "FAIL: SC_B = " << B << "; expected " << DefaultB << "\n"; + if (A != DefaultB || B != DefaultB) { + std::cout << "FAIL: expected A=1024 B=1024 (CMPLRLLVM-77316).\n"; ++Failed; } } - // Control: a value that collides with nothing. + // Control: value colliding with nothing. { - auto [A, B] = runWithSCA(Q, argv[1], 777); + 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"; @@ -90,7 +69,7 @@ int main(int argc, char **argv) { // Control: SC_A set to its own default. { - auto [A, B] = runWithSCA(Q, argv[1], DefaultA); + 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"; From 450c9cfe0fdf986d8239b17ad125833849dd93eb Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Wed, 12 Aug 2026 11:03:38 -0700 Subject: [PATCH 3/3] Format --- sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp b/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp index f8e3f836d0a6a..8d3d46991c734 100644 --- a/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp +++ b/sycl/test-e2e/SYCLBIN/spec_const_collision_input.cpp @@ -47,7 +47,8 @@ int main(int argc, char **argv) { sycl::queue Q; int Failed = 0; - // Regression case: SC_A set to SC_B's default must stay 1024, not drop to 256. + // 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";