From de0a532ba1a24962e962a07b9316d6dc757baff4 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 30 Jan 2023 11:23:34 -0800 Subject: [PATCH 1/3] [SYCL] Add test for spec const info after link This commit adds a test for checking that information about associated specialization constants are available in the kernel_bundle produced by calls to sycl::link. Signed-off-by: Larsen, Steffen --- .../spec_constants_after_link.cpp | 106 ++++++++++++++++++ 1 file changed, 106 insertions(+) create mode 100644 SYCL/KernelAndProgram/spec_constants_after_link.cpp diff --git a/SYCL/KernelAndProgram/spec_constants_after_link.cpp b/SYCL/KernelAndProgram/spec_constants_after_link.cpp new file mode 100644 index 0000000000..a0ffbda71a --- /dev/null +++ b/SYCL/KernelAndProgram/spec_constants_after_link.cpp @@ -0,0 +1,106 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// FIXME: ACC devices use emulation path, which is not yet supported + +// This test checks that specialization constant information is available on +// kernel bundles produced by sycl::link. + +#include + +#include +#include + +using namespace sycl; + +class Kernel1; +class Kernel2; + +struct TestStruct { + float f; + long long ll; +}; + +bool operator==(const TestStruct &LHS, const TestStruct &RHS) { + return LHS.f == RHS.f && LHS.ll == RHS.ll; +} +bool operator!=(const TestStruct &LHS, const TestStruct &RHS) { + return !(LHS == RHS); +} + +constexpr specialization_id SpecName1; +constexpr specialization_id SpecName2; + +void UnusuedFunc(queue &Q) { + Q.single_task([=](kernel_handler KHR) { + volatile TestStruct A = KHR.get_specialization_constant(); + }); + + Q.single_task([=](kernel_handler KHR) { + volatile int A = KHR.get_specialization_constant(); + }); +} + +std::optional FindDeviceWithOnlineLinker() { + std::vector Devices = device::get_devices(); + auto Device = + std::find_if(Devices.begin(), Devices.end(), [](const device &D) { + return D.has(aspect::online_linker); + }); + if (Device == Devices.end()) + return std::nullopt; + return *Device; +} + +int main() { + std::optional Device = FindDeviceWithOnlineLinker(); + + if (!Device) { + std::cout << "No device with aspect::online_linker. Skipping test." + << std::endl; + return 0; + } + + context Ctx{*Device}; + + kernel_bundle Bundle1 = + sycl::get_kernel_bundle(Ctx, {*Device}); + kernel_bundle Bundle2 = + sycl::get_kernel_bundle(Ctx, {*Device}); + + assert(Bundle1.has_specialization_constant()); + assert(Bundle1.has_specialization_constant()); + + TestStruct Value1{3.14f, 1234ll}; + int Value2 = 42; + Bundle1.set_specialization_constant(Value1); + Bundle2.set_specialization_constant(Value2); + + kernel_bundle ObjBundle1 = sycl::compile(Bundle1); + kernel_bundle ObjBundle2 = sycl::compile(Bundle2); + + assert(ObjBundle1.has_specialization_constant()); + assert(ObjBundle2.has_specialization_constant()); + + kernel_bundle LinkedBundle = + sycl::link({ObjBundle1, ObjBundle2}); + + assert(LinkedBundle.has_specialization_constant()); + assert(LinkedBundle.has_specialization_constant()); + + int Failures = 0; + + if (LinkedBundle.get_specialization_constant() != Value1) { + std::cout << "Read value of SpecName1 is not the same as was written." + << std::endl; + ++Failures; + } + + if (LinkedBundle.get_specialization_constant() != Value2) { + std::cout << "Read value of SpecName2 is not the same as was written." + << std::endl; + ++Failures; + } + + return Failures; +} From 46db54f42ee3a8adabd0cc2e2568f39aace3a928 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 31 Jan 2023 00:59:46 -0800 Subject: [PATCH 2/3] Fix assertion Signed-off-by: Larsen, Steffen --- SYCL/KernelAndProgram/spec_constants_after_link.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/SYCL/KernelAndProgram/spec_constants_after_link.cpp b/SYCL/KernelAndProgram/spec_constants_after_link.cpp index a0ffbda71a..1fbc5ffbf2 100644 --- a/SYCL/KernelAndProgram/spec_constants_after_link.cpp +++ b/SYCL/KernelAndProgram/spec_constants_after_link.cpp @@ -69,7 +69,9 @@ int main() { sycl::get_kernel_bundle(Ctx, {*Device}); assert(Bundle1.has_specialization_constant()); - assert(Bundle1.has_specialization_constant()); + assert(Bundle2.has_specialization_constant()); + assert(!Bundle1.has_specialization_constant()); + assert(!Bundle2.has_specialization_constant()); TestStruct Value1{3.14f, 1234ll}; int Value2 = 42; @@ -81,6 +83,8 @@ int main() { assert(ObjBundle1.has_specialization_constant()); assert(ObjBundle2.has_specialization_constant()); + assert(!ObjBundle1.has_specialization_constant()); + assert(!ObjBundle2.has_specialization_constant()); kernel_bundle LinkedBundle = sycl::link({ObjBundle1, ObjBundle2}); From 7b101277d56f9ed9948b140054b7682e7acf6d22 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 31 Jan 2023 10:07:44 -0800 Subject: [PATCH 3/3] Expand test to also check the device values Signed-off-by: Larsen, Steffen --- .../spec_constants_after_link.cpp | 59 ++++++++++++++----- 1 file changed, 45 insertions(+), 14 deletions(-) diff --git a/SYCL/KernelAndProgram/spec_constants_after_link.cpp b/SYCL/KernelAndProgram/spec_constants_after_link.cpp index 1fbc5ffbf2..97655386df 100644 --- a/SYCL/KernelAndProgram/spec_constants_after_link.cpp +++ b/SYCL/KernelAndProgram/spec_constants_after_link.cpp @@ -31,16 +31,6 @@ bool operator!=(const TestStruct &LHS, const TestStruct &RHS) { constexpr specialization_id SpecName1; constexpr specialization_id SpecName2; -void UnusuedFunc(queue &Q) { - Q.single_task([=](kernel_handler KHR) { - volatile TestStruct A = KHR.get_specialization_constant(); - }); - - Q.single_task([=](kernel_handler KHR) { - volatile int A = KHR.get_specialization_constant(); - }); -} - std::optional FindDeviceWithOnlineLinker() { std::vector Devices = device::get_devices(); auto Device = @@ -95,14 +85,55 @@ int main() { int Failures = 0; if (LinkedBundle.get_specialization_constant() != Value1) { - std::cout << "Read value of SpecName1 is not the same as was written." - << std::endl; + std::cout + << "Read value of SpecName1 on host is not the same as was written." + << std::endl; ++Failures; } if (LinkedBundle.get_specialization_constant() != Value2) { - std::cout << "Read value of SpecName2 is not the same as was written." - << std::endl; + std::cout + << "Read value of SpecName2 on host is not the same as was written." + << std::endl; + ++Failures; + } + + TestStruct ReadValue1; + int ReadValue2; + + { + queue Q{Ctx, *Device}; + + buffer ReadValue1Buffer{&ReadValue1, 1}; + Q.submit([&](handler &CGH) { + CGH.use_kernel_bundle(LinkedBundle); + accessor ReadValue1Acc{ReadValue1Buffer, CGH, sycl::write_only}; + CGH.single_task([=](kernel_handler KHR) { + ReadValue1Acc[0] = KHR.get_specialization_constant(); + }); + }); + + buffer ReadValue2Buffer{&ReadValue2, 1}; + Q.submit([&](handler &CGH) { + CGH.use_kernel_bundle(LinkedBundle); + accessor ReadValue2Acc{ReadValue2Buffer, CGH, sycl::write_only}; + CGH.single_task([=](kernel_handler KHR) { + ReadValue2Acc[0] = KHR.get_specialization_constant(); + }); + }); + } + + if (ReadValue1 != Value1) { + std::cout + << "Read value of SpecName1 on device is not the same as was written." + << std::endl; + ++Failures; + } + + if (ReadValue2 != Value2) { + std::cout + << "Read value of SpecName2 on device is not the same as was written." + << std::endl; ++Failures; }