From 4c74b5264fd7a3cbe2c6b952e05edeb7a5f0c8f1 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Mon, 27 Jul 2026 21:53:05 -0500 Subject: [PATCH 1/7] [LAPACK][cuSOLVER] Fix incorrect QR results on repeated runs The Householder-family cuSOLVER routines (geqrf, orgqr, ormqr, gebrd, orgbr, orgtr, ormtr and the complex ung*/unm* variants) passed nullptr for cuSOLVER's devInfo argument and skipped the lapack_info_check that every other routine (getrf, potrf, ...) performs. Besides losing all error reporting, this omitted the implicit synchronization that lapack_info_check performs (it reads devInfo back via a blocking queue.wait()). Without it, the SYCL event returned from the native-command submission could signal before the cuSOLVER kernels had finished, so a subsequent memcpy read partially-computed data. This produced nondeterministic, size-dependent wrong results - e.g. QR of a diagonal matrix returning Q diagonals stuck at the input value on the second and later runs for n >= 256 (issue #626). Allocate a real devInfo and call lapack_info_check in all of these routines (buffer and USM paths), matching the established pattern. This both restores error checking and removes the race. Also align CusolverScopedContextHandler::get_stream with the cuBLAS backend by returning the interop handle's native queue (ih.get_native_queue()) instead of the queue's default stream, so cuSOLVER work is enqueued on the stream the SYCL runtime tracks for native-command completion. Fixes #626. --- .../backends/cusolver/cusolver_lapack.cpp | 144 +++++++++++++++--- .../cusolver/cusolver_scope_handle.cpp | 8 +- 2 files changed, 127 insertions(+), 25 deletions(-) diff --git a/src/lapack/backends/cusolver/cusolver_lapack.cpp b/src/lapack/backends/cusolver/cusolver_lapack.cpp index 6a5427712..2df2c2a22 100644 --- a/src/lapack/backends/cusolver/cusolver_lapack.cpp +++ b/src/lapack/backends/cusolver/cusolver_lapack.cpp @@ -41,6 +41,7 @@ inline void gebrd(const char* func_name, Func func, sycl::queue& queue, std::int if (m < n) throw unimplemented("lapack", "gebrd", "cusolver gebrd does not support m < n"); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto d_acc = d.template get_access(cgh); @@ -48,6 +49,7 @@ inline void gebrd(const char* func_name, Func func, sycl::queue& queue, std::int auto tauq_acc = tauq.template get_access(cgh); auto taup_acc = taup.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); @@ -56,11 +58,13 @@ inline void gebrd(const char* func_name, Func func, sycl::queue& queue, std::int auto tauq_ = sc.get_mem(tauq_acc); auto taup_ = sc.get_mem(taup_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, m, n, a_, lda, d_, e_, tauq_, - taup_, scratch_, scratchpad_size, nullptr); + taup_, scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define GEBRD_LAUNCHER(TYPE_A, TYPE_B, CUSOLVER_ROUTINE) \ @@ -107,20 +111,24 @@ inline void geqrf(const char* func_name, Func func, sycl::queue& queue, std::int sycl::buffer& scratchpad, std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, lda, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, m, n, a_, lda, tau_, scratch_, - scratchpad_size, nullptr); + scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define GEQRF_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -470,20 +478,24 @@ inline void orgbr(const char* func_name, Func func, sycl::queue& queue, oneapi:: std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_generate(vec), m, n, - k, a_, lda, tau_, scratch_, scratchpad_size, nullptr); + k, a_, lda, tau_, scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define ORGBR_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -505,20 +517,24 @@ inline void orgqr(const char* func_name, Func func, sycl::queue& queue, std::int sycl::buffer& tau, sycl::buffer& scratchpad, std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, m, n, k, a_, lda, tau_, - scratch_, scratchpad_size, nullptr); + scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define ORGQR_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -540,20 +556,24 @@ inline void orgtr(const char* func_name, Func func, sycl::queue& queue, oneapi:: sycl::buffer& scratchpad, std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_fill_mode(uplo), n, - a_, lda, tau_, scratch_, scratchpad_size, nullptr); + a_, lda, tau_, scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define ORGTR_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -577,24 +597,28 @@ inline void ormtr(const char* func_name, Func func, sycl::queue& queue, oneapi:: std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, lda, ldc, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto c_acc = c.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto c_ = sc.get_mem(c_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_side_mode(side), get_cublas_fill_mode(uplo), get_cublas_operation(trans), m, n, a_, lda, tau_, c_, ldc, scratch_, scratchpad_size, - nullptr); + devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define ORMTR_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -632,23 +656,27 @@ inline void ormqr(const char* func_name, Func func, sycl::queue& queue, oneapi:: std::int64_t ldc, sycl::buffer& scratchpad, std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, ldc, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto c_acc = c.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto c_ = sc.get_mem(c_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_side_mode(side), get_cublas_operation(trans), m, n, k, a_, lda, tau_, c_, ldc, - scratch_, scratchpad_size, nullptr); + scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define ORMQR_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -999,20 +1027,24 @@ inline void ungbr(const char* func_name, Func func, sycl::queue& queue, oneapi:: std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_generate(vec), m, n, - k, a_, lda, tau_, scratch_, scratchpad_size, nullptr); + k, a_, lda, tau_, scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define UNGBR_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -1034,20 +1066,24 @@ inline void ungqr(const char* func_name, Func func, sycl::queue& queue, std::int sycl::buffer& tau, sycl::buffer& scratchpad, std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, m, n, k, a_, lda, tau_, - scratch_, scratchpad_size, nullptr); + scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define UNGQR_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -1069,20 +1105,24 @@ inline void ungtr(const char* func_name, Func func, sycl::queue& queue, oneapi:: sycl::buffer& scratchpad, std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_fill_mode(uplo), n, - a_, lda, tau_, scratch_, scratchpad_size, nullptr); + a_, lda, tau_, scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define UNGTR_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -1120,23 +1160,27 @@ inline void unmqr(const char* func_name, Func func, sycl::queue& queue, oneapi:: std::int64_t ldc, sycl::buffer& scratchpad, std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto c_acc = c.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto c_ = sc.get_mem(c_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_side_mode(side), get_cublas_operation(trans), m, n, k, a_, lda, tau_, c_, ldc, - scratch_, scratchpad_size, nullptr); + scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define UNMQR_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -1161,24 +1205,28 @@ inline void unmtr(const char* func_name, Func func, sycl::queue& queue, oneapi:: std::int64_t scratchpad_size) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, lda, ldc, scratchpad_size); + sycl::buffer devInfo{ 1 }; queue.submit([&](sycl::handler& cgh) { auto a_acc = a.template get_access(cgh); auto tau_acc = tau.template get_access(cgh); auto c_acc = c.template get_access(cgh); auto scratch_acc = scratchpad.template get_access(cgh); + auto devInfo_acc = devInfo.get_access(cgh); onemath_cusolver_host_task(cgh, queue, [=](CusolverScopedContextHandler& sc) { auto handle = sc.get_handle(queue); auto a_ = sc.get_mem(a_acc); auto tau_ = sc.get_mem(tau_acc); auto c_ = sc.get_mem(c_acc); auto scratch_ = sc.get_mem(scratch_acc); + auto devInfo_ = sc.get_mem(devInfo_acc); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_side_mode(side), get_cublas_fill_mode(uplo), get_cublas_operation(trans), m, n, a_, lda, tau_, c_, ldc, scratch_, scratchpad_size, - nullptr); + devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); } #define UNMTR_LAUNCHER(TYPE, CUSOLVER_ROUTINE) \ @@ -1210,6 +1258,7 @@ inline sycl::event gebrd(const char* func_name, Func func, sycl::queue& queue, s if (m < n) throw unimplemented("lapack", "gebrd", "cusolver gebrd does not support m < n"); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1223,11 +1272,14 @@ inline sycl::event gebrd(const char* func_name, Func func, sycl::queue& queue, s auto tauq_ = reinterpret_cast(tauq); auto taup_ = reinterpret_cast(taup); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, m, n, a_, lda, d_, e_, tauq_, - taup_, scratch_, scratchpad_size, nullptr); + taup_, scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -1275,6 +1327,7 @@ inline sycl::event geqrf(const char* func_name, Func func, sycl::queue& queue, s const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, lda, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1285,11 +1338,14 @@ inline sycl::event geqrf(const char* func_name, Func func, sycl::queue& queue, s auto a_ = reinterpret_cast(a); auto tau_ = reinterpret_cast(tau); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, m, n, a_, lda, tau_, scratch_, - scratchpad_size, nullptr); + scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -1662,6 +1718,7 @@ inline sycl::event orgbr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1672,11 +1729,14 @@ inline sycl::event orgbr(const char* func_name, Func func, sycl::queue& queue, auto a_ = reinterpret_cast(a); auto tau_ = reinterpret_cast(tau); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_generate(vec), m, n, - k, a_, lda, tau_, scratch_, scratchpad_size, nullptr); + k, a_, lda, tau_, scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -1701,6 +1761,7 @@ inline sycl::event orgqr(const char* func_name, Func func, sycl::queue& queue, s const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1711,11 +1772,14 @@ inline sycl::event orgqr(const char* func_name, Func func, sycl::queue& queue, s auto a_ = reinterpret_cast(a); auto tau_ = reinterpret_cast(tau); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, m, n, k, a_, lda, tau_, - scratch_, scratchpad_size, nullptr); + scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -1739,6 +1803,7 @@ inline sycl::event orgtr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1749,11 +1814,14 @@ inline sycl::event orgtr(const char* func_name, Func func, sycl::queue& queue, auto a_ = reinterpret_cast(a); auto tau_ = reinterpret_cast(tau); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_fill_mode(uplo), n, - a_, lda, tau_, scratch_, scratchpad_size, nullptr); + a_, lda, tau_, scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -1779,6 +1847,7 @@ inline sycl::event ormtr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, lda, ldc, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1790,13 +1859,16 @@ inline sycl::event ormtr(const char* func_name, Func func, sycl::queue& queue, auto tau_ = reinterpret_cast(tau); auto c_ = reinterpret_cast(c); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_side_mode(side), get_cublas_fill_mode(uplo), get_cublas_operation(trans), m, n, a_, lda, tau_, c_, ldc, scratch_, scratchpad_size, - nullptr); + devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -1836,6 +1908,7 @@ inline sycl::event ormqr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, ldc, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1847,12 +1920,15 @@ inline sycl::event ormqr(const char* func_name, Func func, sycl::queue& queue, auto tau_ = reinterpret_cast(tau); auto c_ = reinterpret_cast(c); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_side_mode(side), get_cublas_operation(trans), m, n, k, a_, lda, tau_, c_, ldc, - scratch_, scratchpad_size, nullptr); + scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -2234,6 +2310,7 @@ inline sycl::event ungbr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2244,11 +2321,14 @@ inline sycl::event ungbr(const char* func_name, Func func, sycl::queue& queue, auto a_ = reinterpret_cast(a); auto tau_ = reinterpret_cast(tau); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_generate(vec), m, n, - k, a_, lda, tau_, scratch_, scratchpad_size, nullptr); + k, a_, lda, tau_, scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -2273,6 +2353,7 @@ inline sycl::event ungqr(const char* func_name, Func func, sycl::queue& queue, s const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2283,11 +2364,14 @@ inline sycl::event ungqr(const char* func_name, Func func, sycl::queue& queue, s auto a_ = reinterpret_cast(a); auto tau_ = reinterpret_cast(tau); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, m, n, k, a_, lda, tau_, - scratch_, scratchpad_size, nullptr); + scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -2311,6 +2395,7 @@ inline sycl::event ungtr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2321,11 +2406,14 @@ inline sycl::event ungtr(const char* func_name, Func func, sycl::queue& queue, auto a_ = reinterpret_cast(a); auto tau_ = reinterpret_cast(tau); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_fill_mode(uplo), n, - a_, lda, tau_, scratch_, scratchpad_size, nullptr); + a_, lda, tau_, scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -2365,6 +2453,7 @@ inline sycl::event unmqr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2376,12 +2465,15 @@ inline sycl::event unmqr(const char* func_name, Func func, sycl::queue& queue, auto tau_ = reinterpret_cast(tau); auto c_ = reinterpret_cast(c); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_side_mode(side), get_cublas_operation(trans), m, n, k, a_, lda, tau_, c_, ldc, - scratch_, scratchpad_size, nullptr); + scratch_, scratchpad_size, devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } @@ -2409,6 +2501,7 @@ inline sycl::event unmtr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, lda, ldc, scratchpad_size); + int* devInfo = (int*)malloc_device(sizeof(int), queue); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2420,13 +2513,16 @@ inline sycl::event unmtr(const char* func_name, Func func, sycl::queue& queue, auto tau_ = reinterpret_cast(tau); auto c_ = reinterpret_cast(c); auto scratch_ = reinterpret_cast(scratchpad); + auto devInfo_ = reinterpret_cast(devInfo); cusolverStatus_t err; cusolver_native_named_func(func_name, func, err, handle, get_cublas_side_mode(side), get_cublas_fill_mode(uplo), get_cublas_operation(trans), m, n, a_, lda, tau_, c_, ldc, scratch_, scratchpad_size, - nullptr); + devInfo_); }); }); + lapack_info_check(queue, devInfo, __func__, func_name); + free(devInfo, queue); return done; } diff --git a/src/lapack/backends/cusolver/cusolver_scope_handle.cpp b/src/lapack/backends/cusolver/cusolver_scope_handle.cpp index af0881c10..318cbce24 100644 --- a/src/lapack/backends/cusolver/cusolver_scope_handle.cpp +++ b/src/lapack/backends/cusolver/cusolver_scope_handle.cpp @@ -141,7 +141,13 @@ cusolverDnHandle_t CusolverScopedContextHandler::get_handle(const sycl::queue& q } CUstream CusolverScopedContextHandler::get_stream(const sycl::queue& queue) { - return sycl::get_native(queue); + // Return the CUDA stream backing the current host task / native command + // submission. When `ext_codeplay_enqueue_native_command` is used, the SYCL + // runtime tracks completion of the native command via this stream, so + // cuSOLVER work must be enqueued onto it (not the queue's default stream) + // to avoid the returned SYCL event signalling before the computation + // finishes. See uxlfoundation/oneMath#626. + return ih.get_native_queue(); } sycl::context CusolverScopedContextHandler::get_context(const sycl::queue& queue) { return queue.get_context(); From ea9c9b9229e79b55d7121c421a293b9fc4133b20 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Mon, 10 Aug 2026 23:45:02 +0000 Subject: [PATCH 2/7] [LAPACK] Add geqrf+orgqr diagonal regression test for #626 Adds a functional test that repeatedly runs geqrf + orgqr on a known diagonal matrix (whose Q is the identity) and verifies Q's diagonal is 1 on every run. This guards against a regression of the cuSOLVER synchronization bug from #626, where run 2+ returned corrupted results. Co-authored-by: Cursor --- tests/unit_tests/lapack/source/CMakeLists.txt | 4 +- .../lapack/source/geqrf_orgqr_diagonal.cpp | 151 ++++++++++++++++++ 2 files changed, 154 insertions(+), 1 deletion(-) create mode 100644 tests/unit_tests/lapack/source/geqrf_orgqr_diagonal.cpp diff --git a/tests/unit_tests/lapack/source/CMakeLists.txt b/tests/unit_tests/lapack/source/CMakeLists.txt index f660465ba..a3deae888 100644 --- a/tests/unit_tests/lapack/source/CMakeLists.txt +++ b/tests/unit_tests/lapack/source/CMakeLists.txt @@ -19,7 +19,9 @@ #Build object from all test sources # TODO: add list of tests without Netlib dependency -set(LAPACK_SOURCES) +set(LAPACK_SOURCES + "geqrf_orgqr_diagonal.cpp" +) set(LAPACK_SOURCES_W_LAPACKE "gebrd.cpp" diff --git a/tests/unit_tests/lapack/source/geqrf_orgqr_diagonal.cpp b/tests/unit_tests/lapack/source/geqrf_orgqr_diagonal.cpp new file mode 100644 index 000000000..da4d6f167 --- /dev/null +++ b/tests/unit_tests/lapack/source/geqrf_orgqr_diagonal.cpp @@ -0,0 +1,151 @@ +/******************************************************************************* +* Copyright 2025 Intel Corporation +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, +* software distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions +* and limitations under the License. +* +* +* SPDX-License-Identifier: Apache-2.0 +*******************************************************************************/ + +/* + * Regression test for https://github.com/uxlfoundation/oneMath/issues/626. + * + * QR decomposition (geqrf + orgqr) on the cuSOLVER backend used to return + * incorrect results on repeated runs sharing a queue: for a diagonal input + * matrix the diagonal of Q should be 1, but from the second run onwards some + * entries stayed at the input value because the native command completed + * before the cuSOLVER kernels had actually finished. This test performs the + * geqrf + orgqr sequence several times on a known diagonal matrix and verifies + * that every run yields a Q whose diagonal is 1 (and off-diagonal 0), guarding + * against a regression of that synchronization bug. + */ + +#include +#include +#include +#include +#include + +#if __has_include() +#include +#else +#include +#endif + +#include "oneapi/math.hpp" +#include "lapack_common.hpp" +#include "lapack_test_controller.hpp" +#include "lapack_accuracy_checks.hpp" +#include "lapack_reference_wrappers.hpp" +#include "test_helper.hpp" + +namespace { + +/* columns: n, lda, num_runs */ +const char* accuracy_input = R"( +256 256 6 +512 512 6 +1024 1024 6 +)"; + +template +bool accuracy(const sycl::device& dev, int64_t n, int64_t lda, int64_t num_runs) { + using fp = typename data_T_info::value_type; + using fp_real = typename complex_info::real_type; + + const int64_t m = n; + const int64_t k = n; + + /* Diagonal input matrix (column-major) with a non-unit diagonal value. + * Its QR factor Q is the identity, so Q's diagonal must be 1. */ + const fp diag_value = fp(2.0); + std::vector A_initial(lda * n, fp(0.0)); + for (int64_t j = 0; j < n; ++j) { + A_initial[j * lda + j] = diag_value; + } + + const fp_real tol = fp_real(10.0) * n * std::numeric_limits::epsilon(); + + bool result = true; + { + sycl::queue queue{ dev, async_error_handler }; + + auto A_dev = device_alloc(queue, A_initial.size()); + auto tau_dev = device_alloc(queue, k); +#ifdef CALL_RT_API + const auto geqrf_scratchpad_size = + oneapi::math::lapack::geqrf_scratchpad_size(queue, m, n, lda); + const auto orgqr_scratchpad_size = + oneapi::math::lapack::orgqr_scratchpad_size(queue, m, n, k, lda); +#else + int64_t geqrf_scratchpad_size; + int64_t orgqr_scratchpad_size; + TEST_RUN_LAPACK_CT_SELECT( + queue, geqrf_scratchpad_size = oneapi::math::lapack::geqrf_scratchpad_size, m, n, + lda); + TEST_RUN_LAPACK_CT_SELECT( + queue, orgqr_scratchpad_size = oneapi::math::lapack::orgqr_scratchpad_size, m, n, k, + lda); +#endif + const auto scratchpad_size = std::max(geqrf_scratchpad_size, orgqr_scratchpad_size); + auto scratchpad_dev = device_alloc(queue, scratchpad_size); + + std::vector Q(lda * n); + for (int64_t run = 0; run < num_runs && result; ++run) { + /* Reset the input for every run and factor it in place. */ + host_to_device_copy(queue, A_initial.data(), A_dev, A_initial.size()); + queue.wait_and_throw(); + +#ifdef CALL_RT_API + oneapi::math::lapack::geqrf(queue, m, n, A_dev, lda, tau_dev, scratchpad_dev, + geqrf_scratchpad_size); + oneapi::math::lapack::orgqr(queue, m, n, k, A_dev, lda, tau_dev, scratchpad_dev, + orgqr_scratchpad_size); +#else + TEST_RUN_LAPACK_CT_SELECT(queue, oneapi::math::lapack::geqrf, m, n, A_dev, lda, tau_dev, + scratchpad_dev, geqrf_scratchpad_size); + TEST_RUN_LAPACK_CT_SELECT(queue, oneapi::math::lapack::orgqr, m, n, k, A_dev, lda, + tau_dev, scratchpad_dev, orgqr_scratchpad_size); +#endif + queue.wait_and_throw(); + + device_to_host_copy(queue, A_dev, Q.data(), Q.size()); + queue.wait_and_throw(); + + for (int64_t j = 0; j < n && result; ++j) { + for (int64_t i = 0; i < m && result; ++i) { + const fp_real expected = (i == j) ? fp_real(1.0) : fp_real(0.0); + const fp_real got = std::abs(Q[j * lda + i]); + if (std::abs(got - expected) > tol) { + test_log::lout << "run " << run << ": Q(" << i << ", " << j << ") = " << got + << ", expected " << expected << std::endl; + result = false; + } + } + } + } + + device_free(queue, A_dev); + device_free(queue, tau_dev); + device_free(queue, scratchpad_dev); + } + + return result; +} + +InputTestController)> accuracy_controller{ accuracy_input }; + +} /* anonymous namespace */ + +#include "lapack_gtest_suite.hpp" +INSTANTIATE_GTEST_SUITE_ACCURACY_REAL(GeqrfOrgqrDiagonal); From 1dc8b30c3611153c8b2339681dfcd90ccb170a66 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Wed, 12 Aug 2026 11:32:11 -0700 Subject: [PATCH 3/7] [LAPACK][cuSOLVER] Check devInfo allocation before passing it to cuSOLVER A failed malloc_device silently degraded into passing a null devInfo to cuSOLVER, which disables the routine's error reporting - exactly the situation that hid the bug from #626. Route every USM devInfo allocation through a create_devinfo helper that throws device_bad_alloc instead. Co-authored-by: Cursor --- .../backends/cusolver/cusolver_helper.hpp | 10 ++++ .../backends/cusolver/cusolver_lapack.cpp | 46 +++++++++---------- 2 files changed, 33 insertions(+), 23 deletions(-) diff --git a/src/lapack/backends/cusolver/cusolver_helper.hpp b/src/lapack/backends/cusolver/cusolver_helper.hpp index 52da00bb6..c9ae0fbf6 100644 --- a/src/lapack/backends/cusolver/cusolver_helper.hpp +++ b/src/lapack/backends/cusolver/cusolver_helper.hpp @@ -290,6 +290,16 @@ struct CudaEquivalentType> { /* devinfo */ +/* Allocates the device memory cuSOLVER writes its `devInfo` output to. A failed + * allocation must not be passed on to cuSOLVER as a null pointer, since that + * silently disables the routine's error reporting. */ +inline int* create_devinfo(sycl::queue& queue, const char* func_name, int dev_info_size = 1) { + int* devInfo = sycl::malloc_device(dev_info_size, queue); + if (!devInfo) + throw oneapi::math::device_bad_alloc("lapack", func_name, queue.get_device()); + return devInfo; +} + inline void get_cusolver_devinfo(sycl::queue& queue, sycl::buffer& devInfo, std::vector& dev_info_) { sycl::host_accessor dev_info_acc{ devInfo }; diff --git a/src/lapack/backends/cusolver/cusolver_lapack.cpp b/src/lapack/backends/cusolver/cusolver_lapack.cpp index 2df2c2a22..e84781bb5 100644 --- a/src/lapack/backends/cusolver/cusolver_lapack.cpp +++ b/src/lapack/backends/cusolver/cusolver_lapack.cpp @@ -1258,7 +1258,7 @@ inline sycl::event gebrd(const char* func_name, Func func, sycl::queue& queue, s if (m < n) throw unimplemented("lapack", "gebrd", "cusolver gebrd does not support m < n"); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1327,7 +1327,7 @@ inline sycl::event geqrf(const char* func_name, Func func, sycl::queue& queue, s const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1378,7 +1378,7 @@ inline sycl::event getrf(const char* func_name, Func func, sycl::queue& queue, s std::uint64_t ipiv_size = std::min(n, m); int* ipiv32 = (int*)malloc_device(sizeof(int) * ipiv_size, queue); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1515,7 +1515,7 @@ inline sycl::event gesvd(const char* func_name, Func func, sycl::queue& queue, using cuDataType_A = typename CudaEquivalentType::Type; using cuDataType_B = typename CudaEquivalentType::Type; overflow_check(m, n, lda, ldu, ldvt, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1566,7 +1566,7 @@ inline sycl::event heevd(const char* func_name, Func func, sycl::queue& queue, using cuDataType_A = typename CudaEquivalentType::Type; using cuDataType_B = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1612,7 +1612,7 @@ inline sycl::event hegvd(const char* func_name, Func func, sycl::queue& queue, s using cuDataType_A = typename CudaEquivalentType::Type; using cuDataType_B = typename CudaEquivalentType::Type; overflow_check(n, lda, ldb, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1659,7 +1659,7 @@ inline sycl::event hetrd(const char* func_name, Func func, sycl::queue& queue, using cuDataType_A = typename CudaEquivalentType::Type; using cuDataType_B = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1718,7 +1718,7 @@ inline sycl::event orgbr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1761,7 +1761,7 @@ inline sycl::event orgqr(const char* func_name, Func func, sycl::queue& queue, s const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1803,7 +1803,7 @@ inline sycl::event orgtr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1847,7 +1847,7 @@ inline sycl::event ormtr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, lda, ldc, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1908,7 +1908,7 @@ inline sycl::event ormqr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, ldc, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1954,7 +1954,7 @@ inline sycl::event potrf(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1997,7 +1997,7 @@ inline sycl::event potri(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2082,7 +2082,7 @@ inline sycl::event syevd(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2127,7 +2127,7 @@ inline sycl::event sygvd(const char* func_name, Func func, sycl::queue& queue, s const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, ldb, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2172,7 +2172,7 @@ inline sycl::event sytrd(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2217,7 +2217,7 @@ inline sycl::event sytrf(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); // cuSolver legacy api does not accept 64-bit ints. // To get around the limitation. @@ -2310,7 +2310,7 @@ inline sycl::event ungbr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2353,7 +2353,7 @@ inline sycl::event ungqr(const char* func_name, Func func, sycl::queue& queue, s const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, k, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2395,7 +2395,7 @@ inline sycl::event ungtr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2453,7 +2453,7 @@ inline sycl::event unmqr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2501,7 +2501,7 @@ inline sycl::event unmtr(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using cuDataType = typename CudaEquivalentType::Type; overflow_check(m, n, lda, ldc, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { From bd54ae78a48820a62bedc27ef2dc294dd44b979f Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Wed, 12 Aug 2026 11:32:17 -0700 Subject: [PATCH 4/7] [LAPACK][tests] Make orgqr depend on the geqrf event in the QR regression test The test creates a default, out-of-order queue, so the USM orgqr call was not ordered after geqrf. Pass the geqrf event as a dependency; the buffer path keeps relying on accessor ordering. The test needs no reference implementation, so stop dropping the whole LAPACK domain when Netlib LAPACKE is absent - only the tests that compare against a reference are skipped now. Co-authored-by: Cursor --- tests/unit_tests/CMakeLists.txt | 8 +--- .../lapack/source/geqrf_orgqr_diagonal.cpp | 40 ++++++++++++++----- 2 files changed, 32 insertions(+), 16 deletions(-) diff --git a/tests/unit_tests/CMakeLists.txt b/tests/unit_tests/CMakeLists.txt index ef2c728f4..5bf597512 100644 --- a/tests/unit_tests/CMakeLists.txt +++ b/tests/unit_tests/CMakeLists.txt @@ -31,9 +31,7 @@ endif() if("lapack" IN_LIST TEST_TARGET_DOMAINS) find_package(LAPACKE) if(NOT LAPACKE_FOUND) - # TODO: add list of tests without Netlib dependency - message(WARNING "Netlib LAPACKE headers or libraries are not found, LAPACK unit tests will be skipped") - list(REMOVE_ITEM TEST_TARGET_DOMAINS "lapack") + message(WARNING "Netlib LAPACKE headers or libraries are not found, LAPACK unit tests comparing against a reference implementation will be skipped") endif() endif() @@ -66,9 +64,7 @@ set(blas_TEST_LINK "") set(lapack_TEST_LIST lapack_source) -if(LAPACKE_FOUND) - set(lapack_TEST_LINK "") -endif() +set(lapack_TEST_LINK "") # RNG config set(rng_TEST_LIST diff --git a/tests/unit_tests/lapack/source/geqrf_orgqr_diagonal.cpp b/tests/unit_tests/lapack/source/geqrf_orgqr_diagonal.cpp index da4d6f167..b78574ad0 100644 --- a/tests/unit_tests/lapack/source/geqrf_orgqr_diagonal.cpp +++ b/tests/unit_tests/lapack/source/geqrf_orgqr_diagonal.cpp @@ -45,8 +45,6 @@ #include "oneapi/math.hpp" #include "lapack_common.hpp" #include "lapack_test_controller.hpp" -#include "lapack_accuracy_checks.hpp" -#include "lapack_reference_wrappers.hpp" #include "test_helper.hpp" namespace { @@ -106,17 +104,39 @@ bool accuracy(const sycl::device& dev, int64_t n, int64_t lda, int64_t num_runs) host_to_device_copy(queue, A_initial.data(), A_dev, A_initial.size()); queue.wait_and_throw(); + if constexpr (is_buf::value) { + /* The accessors on the shared buffers order orgqr after geqrf. */ #ifdef CALL_RT_API - oneapi::math::lapack::geqrf(queue, m, n, A_dev, lda, tau_dev, scratchpad_dev, - geqrf_scratchpad_size); - oneapi::math::lapack::orgqr(queue, m, n, k, A_dev, lda, tau_dev, scratchpad_dev, - orgqr_scratchpad_size); + oneapi::math::lapack::geqrf(queue, m, n, A_dev, lda, tau_dev, scratchpad_dev, + geqrf_scratchpad_size); + oneapi::math::lapack::orgqr(queue, m, n, k, A_dev, lda, tau_dev, scratchpad_dev, + orgqr_scratchpad_size); #else - TEST_RUN_LAPACK_CT_SELECT(queue, oneapi::math::lapack::geqrf, m, n, A_dev, lda, tau_dev, - scratchpad_dev, geqrf_scratchpad_size); - TEST_RUN_LAPACK_CT_SELECT(queue, oneapi::math::lapack::orgqr, m, n, k, A_dev, lda, - tau_dev, scratchpad_dev, orgqr_scratchpad_size); + TEST_RUN_LAPACK_CT_SELECT(queue, oneapi::math::lapack::geqrf, m, n, A_dev, lda, + tau_dev, scratchpad_dev, geqrf_scratchpad_size); + TEST_RUN_LAPACK_CT_SELECT(queue, oneapi::math::lapack::orgqr, m, n, k, A_dev, lda, + tau_dev, scratchpad_dev, orgqr_scratchpad_size); #endif + } + else { + /* orgqr consumes the factorization and reuses the scratchpad, so on + * a possibly out-of-order queue it has to depend on the geqrf event. */ + sycl::event geqrf_event; +#ifdef CALL_RT_API + geqrf_event = oneapi::math::lapack::geqrf(queue, m, n, A_dev, lda, tau_dev, + scratchpad_dev, geqrf_scratchpad_size); + oneapi::math::lapack::orgqr(queue, m, n, k, A_dev, lda, tau_dev, scratchpad_dev, + orgqr_scratchpad_size, + std::vector{ geqrf_event }); +#else + TEST_RUN_LAPACK_CT_SELECT(queue, geqrf_event = oneapi::math::lapack::geqrf, m, n, + A_dev, lda, tau_dev, scratchpad_dev, + geqrf_scratchpad_size); + TEST_RUN_LAPACK_CT_SELECT(queue, oneapi::math::lapack::orgqr, m, n, k, A_dev, lda, + tau_dev, scratchpad_dev, orgqr_scratchpad_size, + std::vector{ geqrf_event }); +#endif + } queue.wait_and_throw(); device_to_host_copy(queue, A_dev, Q.data(), Q.size()); From 1cf8017e57601acdb07be8568f9c801e18ee220c Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Wed, 12 Aug 2026 13:20:55 -0700 Subject: [PATCH 5/7] [LAPACK][rocSOLVER] Fix incorrect QR results on repeated runs rocSOLVER calls only enqueue work on the HIP stream of the surrounding host task or native command. A submission that depends on that work may be scheduled on a different stream of the queue's stream pool without being ordered against this one, so it can read a factorization that has not finished yet. The Householder routines (geqrf, orgqr, ...) have no `info` output, so unlike getrf/potrf/... they do not get an implicit host-side wait from `lapack_info_check` that hides the problem. Wait for the stream before returning from the rocSOLVER call, which is what this backend already did on toolchains without `ext_codeplay_enqueue_native_command`. On an MI100 this turns the geqrf+orgqr diagonal regression test for #626 from failing on repeated runs into passing. Co-authored-by: Cursor --- .../backends/rocsolver/rocsolver_helper.hpp | 18 ++++++++---------- 1 file changed, 8 insertions(+), 10 deletions(-) diff --git a/src/lapack/backends/rocsolver/rocsolver_helper.hpp b/src/lapack/backends/rocsolver/rocsolver_helper.hpp index 5d4e6e821..31feffd7c 100644 --- a/src/lapack/backends/rocsolver/rocsolver_helper.hpp +++ b/src/lapack/backends/rocsolver/rocsolver_helper.hpp @@ -150,12 +150,6 @@ class hip_error : virtual public std::runtime_error { throw rocsolver_error(std::string(#name) + std::string(" : "), err); \ } -#define ROCSOLVER_ERROR_FUNC_T(name, func, err, ...) \ - err = func(__VA_ARGS__); \ - if (err != rocblas_status_success) { \ - throw rocsolver_error(std::string(name) + std::string(" : "), err); \ - } - #define ROCSOLVER_ERROR_FUNC_T_SYNC(name, func, err, handle, ...) \ err = func(handle, __VA_ARGS__); \ if (err != rocblas_status_success) { \ @@ -166,14 +160,18 @@ class hip_error : virtual public std::runtime_error { hipError_t hip_err; \ HIP_ERROR_FUNC(hipStreamSynchronize, hip_err, currentStreamId); +/* rocSOLVER calls only enqueue work on the HIP stream of the surrounding host + * task or native command. A submission depending on that work may be scheduled + * on a different stream of the queue's stream pool without being ordered against + * this one, and then reads a computation that has not finished yet. Unlike + * getrf/potrf/..., the Householder routines have no `info` output either, so they + * do not get an implicit host-side wait from `lapack_info_check`. Waiting for the + * stream keeps the enqueued work complete once the call returns. + * See uxlfoundation/oneMath#626. */ template inline void rocsolver_native_named_func(const char* func_name, Func func, rocsolver_status err, rocsolver_handle handle, Types... args) { -#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND - ROCSOLVER_ERROR_FUNC_T(func_name, func, err, handle, args...) -#else ROCSOLVER_ERROR_FUNC_T_SYNC(func_name, func, err, handle, args...) -#endif }; inline rocblas_eform get_rocsolver_itype(std::int64_t itype) { From 2f0edb2ded62ec7daf305760a339dc15037baeb8 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Wed, 12 Aug 2026 13:21:01 -0700 Subject: [PATCH 6/7] [LAPACK][rocSOLVER] Take the HIP stream from the interop handle `ext_codeplay_enqueue_native_command` requires work to be enqueued on the stream of the interop handle, which is the stream the SYCL runtime records the submission's completion event on. Query it the same way the cuBLAS and cuSOLVER backends do instead of asking the queue for its native stream. Both currently resolve to the same stream inside a native command, so this is an alignment with the other backends rather than a behavioural change. Co-authored-by: Cursor --- src/lapack/backends/rocsolver/rocsolver_scope_handle.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/src/lapack/backends/rocsolver/rocsolver_scope_handle.cpp b/src/lapack/backends/rocsolver/rocsolver_scope_handle.cpp index 264515d6e..7717d305e 100644 --- a/src/lapack/backends/rocsolver/rocsolver_scope_handle.cpp +++ b/src/lapack/backends/rocsolver/rocsolver_scope_handle.cpp @@ -143,7 +143,10 @@ rocblas_handle RocsolverScopedContextHandler::get_handle(const sycl::queue& queu } hipStream_t RocsolverScopedContextHandler::get_stream(const sycl::queue& queue) { - return sycl::get_native(queue); + // Return the HIP stream backing the current host task / native command + // submission, which is the stream the SYCL runtime records the submission's + // completion event on. Matches the cuBLAS and cuSOLVER backends. + return ih.get_native_queue(); } sycl::context RocsolverScopedContextHandler::get_context(const sycl::queue& queue) { return queue.get_context(); From f20a8074e2be8bc2d421dbe37a58bca63be2917d Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Wed, 12 Aug 2026 13:21:07 -0700 Subject: [PATCH 7/7] [LAPACK][rocSOLVER] Check devInfo allocation before passing it to rocSOLVER A failed `malloc_device` for the `devInfo` output silently disabled the routine's error reporting, since rocSOLVER treats the resulting null pointer as "do not report". Throw `device_bad_alloc` instead, matching the cuSOLVER backend. Co-authored-by: Cursor --- qr_diag_repro | Bin 0 -> 194312 bytes qr_diag_repro.cpp | 98 ++++++++++++++++++ .../backends/rocsolver/rocsolver_helper.hpp | 10 ++ .../backends/rocsolver/rocsolver_lapack.cpp | 18 ++-- 4 files changed, 117 insertions(+), 9 deletions(-) create mode 100755 qr_diag_repro create mode 100644 qr_diag_repro.cpp diff --git a/qr_diag_repro b/qr_diag_repro new file mode 100755 index 0000000000000000000000000000000000000000..0aa52f494e431f89e9440483ae2108e679b0820f GIT binary patch literal 194312 zcmeFa4PabF)i-|A4+tOH@F7M~SOugYl%y>sr6^n4g{`z0DTONFHcitukR~D7loTzd zrqM3zV*IF7J~WEHMn%Yr7Sn1Vg-U`bghwBZ;-e9HWMYsf0zyQj|KIPNnY(v4yZwOT zWxp10|l36kf3!d{v8 z2K@J(-aEX5ke?(lmar_G@f@5WS7|DlP7uPzcMS&vWa2jnAc~B8T>-> zU(d^~x75~KYVqCE;k&1k{o?Ra@pIT z9d_?xOE2Fz6xx2fr*h?aXVg?LpEdK0niZ$lRM$2&pWa+H>-1SOr#041Qw1Ymq`_gl z;G!kogch%tFjwD|GtYAVNXibMcLdV&@Q-*Li?<$m%uz+*9__qlE%JFY8{#}B9{E^aSlJdL{;&}!BeHi~L@UO4=;n_1k z`<{)15551PqZY=?zkbiMql>?I(VdT6KY41&lYfYHTy@9~-#+WeFK?{BYWNe+9{l%5 z@BF~6pLq1npEfVJ`ktqrcxC-T5&7w5)rKQ=eKA4;2x?uU=g4L>0dpD*OW zzb6mf)AGn?-dl3h?SeAr;&Xi-{Dbqz=cjMY-A*XaIDRb;|8LC0|IhQ#y)+NqhCJ;Y zokvf0=IQTEdHDQRo^~G0BhMQ#@pH-f+B|x7S04DSdD?$t9y#BZM?Poek>}BQ`a6~f z|3Mh7T>a|IqqmcQoC|*}kNkg{hwjOF_fI~V`gF+=C^!LR>bMrYVkNod|9Sh-K_V1WH{H%Fn z?)Izl%&$rq^moDw?fqz2NEl0KCfX;we4D~2)9~Z>Q;Was?y%(FEC2PpTY>)$@6Fz7 z>t|`H=S{_;%yvdUq7V#+{622sUF#I?=za(bIq|)3D}1A+`+h5*f%6pZ+22+>UPbdV z;jh^Cd$agF4)crccWqUC$Nvw|9^oZ-Df|)S`8x;WCH#ChEP46>{GAI2gzy2|uj3W$ zb=!H-p2g3thbVs0AHst9Kh}Xc8O{vb8M;p4P7c>QKG!SU$?c=o4wQryE{)+oQP8FQ zLstG#Tktsaf$fj>D!yyyTbBP(D+kxF4`ZDrez;ZfDHr}0U`Plr`9)arj{kEoACEwE z(Q@}(ffp6X!+Y5F%ezzYUHd0ne%iYg{$)k?)`2m$6SnQxX1udtA4s?3ql)kN@5g*6 zyz93L@3Zav3H%V=GEws!B^% zRmEc~s@7IlRz=okV$4!-=?tP&*RHIKRj;Y9S#Z(flGsdUZTZaFiZxY{NQ7{J;uY03 zrKNMvTX0cTQ@o*~<`Ptx6)T-x-Dt^ZQ!A<(D;Hc6iB&d&OE3l=Uffzi1jR$3aHTT@e4DU%U0%r&r_>I-8m z@zqMuW6Plm7hh7cCy~A;olEOg)8N|4uzoca^%a#@Tf;lObXCAq!SqQd>@~Evv4LS2fgD)Xc7`tyo@D6{}vkU}uHW7XRfWN zX~I%qE%CzgwIvW~jWzBvsV}}LYY~am)YMQ1(u<3o|7C9F*sDcj@0JKzr1DGP%({kE z*EG1rC_lgDD{C+n@&6jkD1WuM2!`m)s+j?N;6#!nDNE)(zYMj*L}sCJ{wD(PQ=D4cvGuk17c0J)$spf^>vM4aeCUH5w|B+ z4I@!O?}=k~@#2~P!qVIQ;sp6}{$?oKUKz}l{}hAy&+}=`QUq9SP64GLkvd#(c4>H*|#OD z6jrw$fm;)NoE75OEm(NbY($c^P12d^v4wS&SHtGjR9CKxOpipS&xB=cKs*Q^FjiC1 z2(ySpWnIk*4R`W)AfNi}!EF4St6#prnX6pxz!@braE@_b0)w#%v95%?%drO_u%_Bf z=!df1&sl4#)>PK7immtz#RST`HkU+-zLu=BVQAbv~c9vsRds;f9zM-xj3;Q~4c#r+}sm@%9 zy%b1IYp5`=#jpqQ`PK8tW?reZQ5q}Vt?X@tu&2$X+3>|z)HTGaS3re=oW1rPrX2AK z=jW>WraXh2r34gM!@4Y74IXA*3*(qo;>0zhQ8u95o*Sc&vXxD>l~j!d=U4mkml2O- z>WPlg^b$57gLcK@s~hUBjX?pL*X_f|X*2MfGl^xfvXGw#($J&KcIX>S%EXVns!~0z{Qix4Vtp{`h7y|$D?=B9=0 z4a>GM*+v_P>2xO;K#6HM7aFRTSJYJ0R#q*~v&{o_zId@JL|_-Fy0)^RY7HX(Jhgmb zl%Id`;=08WP-O913MZq!PPSIES;dlQHE%igV#K<0bGx#+BDNfBT`fi(`!w3C((1ZK zaa&->mp4>2tV8dr+=44KyUSwgfd+RFqTQAda#O8VwB<>teN4*Tz)W zMI~Wm_H___C6@;Fy}qmi>Q64l>bYRNF2mg%r^c1ogMy0c)U3T$r^1;lYnmEY=b8?= zny@ormJM3>wcm`Ep%IwOk&D-WzKj*lC}?R#7xr0_&Rf`b8N#S7f}xDtp%lx2*f;TiGcvO!sAm0IU4>o+X|}PVx)D1BRni7F)8Sgjtw69U`>LDY-sCl4(lL5v zHEKvzkpr`GMb#c>CFFp4rz>Dh-6{krnO1#IQcCYg?W>BEPPZG*tE*R3H9*nJUC^_n z_S$Ml=HjZxrZrU_{L(dbYpWnKOlZVVh@i0L3XZEPR#dI$xrfXZT@F> zU5%6ZiQamfkDu~c+%%ZrZ9<;E!MpU_8BBY)65;=b0`PNvC}d#I>(QupX#no-Q4oH9 z0PfDm33q4eEOzJbgkMg8`*kTlI7XK{jSgHGfEPF*o<#vT1n2*X18^Me`@c{C&O4ay zH#Gn!PwrO|fKx{9R~CSiSN97C;1IO`iw59O0RLAWfMa0&-_iga3gG{i1>m^Kg*BCiuUW0GvAQewzYtTy61x?E!f3UeuNV99N6{Uq=9rt0Dfc zGXTfcAOF`CfaA)D|LYFGaizxp^#pSwcv%2G5P;7Pzy|~Hvjgy<0DMjWZUXRg0`SoQ{5=8qSOEUs z0Nk_psc2Kf0eE2mt}{dnivsZT0{91#2PAMn0tX~;KmrFOa6keFByd0i2PAMn0tX~; zKmsa(=cXL@QndBy!f0~xV|4{yw5>NjVb?&kb!XuNg0Sn%9|OGWq@UowQ;rWK#r$fP z4G-v&Udi9hJh$G4yCnZD=DFoI+#&g|GS98H;daU2%sjW) zhFc{6S>_oA4%bWm2Ije?HoQ#oA7`FhX~X4`zm9osp$&&6-@rV#&W1}Qe--oGG8+y_ zemV2pDjP16{D+w57TK^T`HPw7*4XgqD*&AIKIXY4HasNx^O)yW*zkbl&t{%mV8gwV zpUynDzJ|Lbe=76b@*3`t{5zTFR@ZR5Rua>>6`i#)fshQpG7mU(Vz4VOs%Pt0>GYd9qNCz$6J)^L&Jf6Y9%u7*9y{|EEj zvKk)!2mAkH=0nU6N&bH3xkWWRAo=agb8Bk2SMql=&n>CpF3Eq3d2U4wcS!!L%ySEB zxLxu$GtaH3;TFk%mU(VD4cANl2IjfdG`vjmA7`FhOvB}pzm9osEe(ey-@rV#j)qGl ze--oGG8zs^emV2pDjF`5{D+w5R?x5~`HPw77SQnM%QF7VmoPsh`SY0PR?zT(JhyU&TO@xZ^W4H2u9y6w%ya8z zc$wrUFwZTU;d05pv<7)@)eMIv|19&|q8Toc{GXWT*3587@=q|&Et%mW$^V*pZp93H zlK&6px%EFh`ga+B=EKYnN&bH3xfL@!Ao=aga|>p;7x^X8#D7OOJUxI3_F!AhUw7>) zh^E>GA=qf5Z2*szeQl56!O};8&eDF6Md?x?B~YD3dIO5~UDMQxHg^r-T%K4h@Brq3u=GQ2ZK17v5 zD1*W@piJ~2n=y=5N6B1@lSzovQ-Xzfa5;#kQZGtDigq-bY?DVK(jQH371ZQrqC&Ps zc$nLjY_dn_rM8k#G__gEE^sWZR+fq~Sb82*CfiO^1;HFO(+-?K58pinK7*_ETUlOtf*SK$ql^h>W63f|}e+RCH;J zJW`La?FE+m9#Q+$R_Z%?MildQH=MM+km1<#@~|`G$EaQau6F&+&Wx9&)U(}3jxz%f ze`XK?Gebb4w;hBedJ5V7Ay_SRU#dg++I?0xe?*DS3_))*%0@48pztL;+{|bp0wjWm z6|g`OJ)@{%u2%|DJ-L}62$6j{dG4kMQiy{>v z2Q_>n9!}9mfe3y*0g0juB+=7>D&|Ej=jgsvPi`g%-Dh?450uCdmLUfUEg2M^A$ey> zl)?}m&XN$4=-EUHo0Wo8Pi`g%6j=JO96r>7OjpUJt^g@jU6=YeCL!8`8Ig82kZ3E(pzvM*o!*9s(AQRi zht(>9Bznq7;lJP*gMw5~ZYBt0&g$k5C=m)>Yybkxpl}?3PH#OT^tBb?VJQeC(KD44 zKB5$)dU7*Cpup;8gQGCO2C#-^PK52uC$M1Y$lAW=hsBzneB#dIUOLieS5ax*~~ zLRL5TqC}LW7!*N&FoVJ)Kygaar4)M86a<1*oD?oq3Q`@bxKdzsbA_W&rWE=!DBKAY zH-zmBCZ`(49fyMxZz~tXB#xX$k_tDozTY zRSHraW<@Emy4fNWbU3Iab~u_z4?zVU-2Q;AS_5n-3HK#ntK^Z03_wW)HMyCnP?AmZ zNY%3KmzDcuk7z(@D~BFRLKHI~vePh&bZwY5N|k}bj)zswtCVh;(v?RdG6uSWiXjK; zeOtau=RS2csovv8g*7qmVX4`}hxtY3^Z$6~yFw+-l zT2`_Nl9u%m+}10%e3(sf_Q@@&%ShXCyPBC~k35vyrGS`Ui0-EZHQ7dxa!WjwCymws zu~Kd70z^8Z$@LxPiuV)NLqEiHNqJ8f(_ho{0Mnd5z!_!wQ<@eZGA!vfdZy7tqy_D@ zn4dHXjed1wqvT|AuS&X{nPiVVl=m4Xl&(hLkuK2fGDC}mcWwiinTlJ^^i)mPGyOJA zcQ9SV^nLCNM$vm|5%`O4582x`<~t)`Wy4Nx?_BvnG_fNw5KW}EpmL(^^Z4)Zy@fy@ z556TgxH&g?NpA3|xxokL20wXlPEM}Q4UXmpKYUQmMmlqY+j4`K=LVmf8~nE1;M*tX zwK6GtOLvb!dy(F z;TGl)GuUe5Ll@>EB+Y&AK#gQdP_eniQdyXZ*YqfzCh^bF0wQ#raxI%i?^sre$%CA^n1S z!A%oJ7Qv?)#hMK#M);Y%1J-sF1mdOOg9-F~2rhyl9~0S5vW;oK0-X}#kq_y;_~P3_ z(nq`Ouh!o4Va{D6C3e{#WK)wnwD&{IBzxqcR0jYtDYlhN32L&9AniTzZ0}{^ZzEo+ zM;88dnwEusjizPcuh6tC{7aC2LBmoj2?nF+nD8c&=m!ar1QS*gy-1qB9Z#OxnD#4B z62$W*Auh#T#MF=*<;^MyKFpalQeva@E;co}!vieQ*3C?^M; z1o506i~;W>O4B+Oomkj9&A}ItnmQGiOSw7~@7J_A74tMLPQ_VBzo45lY3Pz z?aU;5oMgQ$}Ssr+9Z-}LKM4X@V)j8d`NG5rJ^!#{u37h1a&ls`V$nv@%6S}rKrrCQ(ddY*cI@FMXiUU?QL&VROZdc z94c=f_^3NkbO=$`DJt{k=MHss5VaRYhZ1#zqB3v3=}_B)s6Rvz2GQG|P*mp4D-Kn+ z$?U*(qv$ZAeoj%DH;WzWP_WhSq3Cd;-lC|?n@%UJLeX5?Y8Q$S_ITUBqNvQ9^|*K; zaxDp>-hrYw617uNnK%F8TBYOU3#$`FQ;6E7sLY#h*jCqL;#uln1H6cM+Z2y^v&%NQ zzSGC+0GyXowm+bF%$v(>-`iCGbuJ$XEX?IE0U_sH9@Dw}Wj1h+Yk=7PT<+M!No5;t z+r*;D1L97%wKIdojgQv8L9o|?q`3&~!2JumqTbXRaB<_H@kK%eIV2R1^3t^3$ zj-qil*I_z6IQhGP6kPehIC{h%xgpJiT@OZ4zc2C|;Ws7IrweAm{1y`roApfJ%P&_) zQ|tT9DgG=ozeVyszQTI%qnKXicfZCMCfYiIE=GWBP&C!lZ;ttug@})w+A53F4s8d= zw7|!DBaU;8HqoI?Bzu^+caV*L93|4K;Kk7*7UlI1a{-sFW8`9*3xOn_Xlpv1llobgTMSLBh!taZ^1sT;@WoSY)eB(e1K6 z8I3>K_6l0yc?1zbLU)&Y_%NSGHGm9hAQ&U)|0+oAajHj3Qd^lbeJJTHXqr5{7_5OX z%gnPxPezNd;WU8!lp+eE7bzW(UsGlmK}w*q1yJw^UkC|oa}&!vp^|QQWlY@rN7gFV zTAN%66R1Th<_@jZ;>wt~_di+di9w;&>`Iupw`g8i+ar0Q3w7#839KI56NPGqmThXg zRw~{g&l3wawcUq&OAr5xnH24wFkNRgy0e;?0R0(cMPn_9Bv>W_Fc&L7%ZSXrGohMP zhLREQ2BEOAcgp4+H^@ja*&FKLBpDUrv+UR*+R#$h5ED*_x2dSgT^SQ9#Gn0A#1VBR zOgJH~(pq6x#)Jy-Hmz0WN|;a~vc262@%DF7h-XQ~8|1lrAxcl5W8X-0+WRQED0io?a?dCB(9qK-u5 zR;WU!N#UaB@G$3}ZkkHT@S|A(wA0MFSNEUTV!eA@p__{84_`-Osc}d;d;bi(H`3_F zzxx%8_q_*0)1p%%qbh;J&`;=Gv6BE3DuKU%I@D-0yWEX-ywZnSUh<|yb|fOdan>OP zxafF%y5}VU->{RW7V9-FX$-(9PX*ouF=L6FZ{w5Jr^};>rqO6($yhY8-l!4cb3J0W zrXllvZ50UE)!t4I1KZnJiF?gYHAA%CXySg;159X5>&|kJMA+xVI*^F`#0)9&fcXac z*q%YZ2u(wL;32hs%={9bF09}f+9|@CGA^EIeL9Wb`e32dxjmy!#`=x>)*-ZyzM=gj z+P_I~G3H`DN6jahPwW6m6P;5a+62-eEcETNUa>CONkWk1Ca|VvmKR1_cN9kxg;zch z^dCS!4EjT$J%nd}GsH-?y0NNK6p(=RcnGWf5LxW=MFsEosfHwiQqyr;qP9jW36pcm;F%8%w zj9%t=FVGcq+L%7Br-+_GZq&RLaiBIm8ol}R;>ACE#NagE0WaK?q=ge1V-I z2pNK=!>|avg1y5i?be+o=qvrpjQ;ka45fWx=({`_(TEIuA5rj?9@=k0`xrOu?MaOl@giQX5nYqz=aQI-i0Fi6?*@aMN1 zZ?%D-Po>t5L7_1Iq$`80`lI-x;2*e9d@8!g31hHriT);&m$yc{?N}cgogFc2C!1Tr zk@@5!VC(Ajrr4Z5fRwhJF<(m{pP^{C-`PYtjCxl8LKJs;eF>pI&`()o&;!`uRvk+U zdep3CMy%t5B5-k)MJxe*%&dRA)~61PnlGdqGv3^o<5NTwX$ zf^^3G@b$s*3JC+oU-tYM*B?#-Vg21oA*26x{m5t^>xrM=J9(z(^EiBJ19UIT*bZ^h z7F>TcLR(8=EqP_p!sM033zJLO-qb1Ie_>%_ z=84h76(<%h#D9twCN4R#IGUJ$V#r)idC<#DWrgwhCxHNXW(&7q$UyXv0o&=Wod_+LDNT9v%`~ z$27v!Dkk2WkOZ3={}RPK1_u+sEv6XLB+I^q&?(3|yl)|EbLe;hLQTPJjJ!b2(l$ z96vw&@;vp!7>^+g))1ck+3{oNN$mY@;>W%VC_ygoh#zktwjV#Ea*WF$=F0$PJ;q04w4~=-H^eqe5VCOf6<-Jb@7cE305evGb2@O2G+mya%H4^` z_iV5;3Y1av9is3!eZ`N+qxBbL|7IWQPaIFbe9!bxcl4j!`5NiB;uTtGq(*YK|8w@- zf1*hLUn~8I9*8W4=vipxAfokn`{D~s{f35Tv8)5WqJQ?(d4kOk$dr>0mEB6$!NMrMa(dnqUD=aNy)ovO!&)ST6 z%UtQNY@>i4%#1GLCJ}kk+^Py`%wl0e<5?%~;1VvRTG+A*3YVkmuO+#rO#lp;6Yk^g z^1TD+g(=B)cp~C<{6uizIpCc**#*P)ObT1VGeM5(i|f^=G@}_fFDA@HwlrzNAUw;n z8Z=aTp6#;IVXQ=r{w$Ox$NKXFRmv=XvL5o=qE1(bZ+eq3zu4wRzSC^%p{t6-$ZE8Y zjYlvm&Katumi!>Y6pzic3I(t=#SvP0Zkhjt)svA93}JoRF83xPKOA|x>lAY^WUQ@L z`e0Z4MyZWnZiRIH{Sd|<8=jWkJ_(qb(1j)F&7yXpgxyj5Z0T$JD1;8R3WNDA0-bW# zXkXhE3I&U)%Q-6WMp4)$?BlXLL8;9|GC$d%eQQ%Zz6uC-j=2H7(o;z^jt=S@CUYVk z>#u@?8D;}h?gGV$n$|YWE;RNWPCDm(o6jTsdab<$Se999X^er~{!5wH+t_7K`pfqJ_KguNU%FdDIq=V+XkArc% z!TmBQ(B;Mp@!Sc->+L`?p9dJ~ycRow+c)umCDk*Aa}%ib5K6Ya1HGQx&6wT1mF?gi zT4`siw2@2^2_lrHW`h^B&btBBcy0d@*>!YZHYg!q$w&*(2BaO~XHXjHk#o#RQagj6 zZsn&JCEMO9{P3$J<_3~W_Gmj>rHy2YNZ?1B`euvp(=RQ89|7(pLejlQKgy=Lq0ZUO zE%D@{E#?~SN4M5XwlM*QyT;?TAoAW~HYk3V!)F5coy1S|NVU{f=FB$~v(sTR!M(>s zTZc5R51b1NTx8S1B#zZs-Mlnztd_AJUM2$8wy7eUw@@~Ni5p)>e%@+tk9&})LgHcfL2rFgOApa_Dyfhg1TLrR|BX1os_s(q z=3Yr_TgIH7$9Pp0IrF^soN4r7H9>R9{Ends$k!vEp<6|uNfQqEw$s7d-0ms(ubIXw z$*oFlvye%q*w{k7i?a}KL?H*-Jd!oBZC2$a_m%GlbFIRVjOA2TNy2L3SPNj~+M-$@I*RvRHMd0=ZY$QDC>P#XMNU5uGQ{bIH&pRJKQJD*2_*45b1e8+ zn7oas3-Ow#)bG(ow<=?sm8ldNLO;m~`YBu_hN6k3*fCm)O)l(YCd$xrk3DaD+V)j) zH&JY<9+ETL!f zjT1}z-lcu-wS8Zh+Dgd6#Ea&5@PfXx2Kuhe^zd687-MB_voe+m9^)AiDSk~I%At@gGTjb>;a-@t;_^#i_*U+A$~*Zome=^;88W2yNo zTEVo^cDArU=NoFbvwX5pku=2BAvK*GJBcu#ZL0EjA)moUy|U2___i5oHcEw!twIf7 zDUdc(Yz%^78xRbufD?XO#4NwZs(|Fa_Pt>46YFt(sjUWRXr3$SGM7xIm=qJxv8nCT zXm5Q{YY#CePq`H)fRG#HzhVL&nkdzmrmHh4)i-TwyK!7~LT(VaRKJGRahDBsV0=g! zX8>HGZ@V4q=VcFJJxy&y6|SqB@QfF6MGZojX9+QMP1`EXU;%&!0Cz?k$MrC45dXf8 zOLdWF?uHCls|B@32$B>b<=}&P`%JOj=cZh-!+gwE`KVuosVURi&VezW@`)QJxZ7Qy z!HR7PwxT<*ZGsH_VFLO?6KghOa8S0|l`(Pe&8+n>1_~w1TnQ62dEmjswANBr#>Bl} zX06RytK5|^fm#%WIYnzlT^SSieucGSS}W{Im~b68pZ%^dQsxSoxc6(UHdU*YxDqB< zjdC)VYqhDakO_PL10^F~Qq&F}7j9z%UsoY-_@V-6f;!m&*K7Fq5b1a@#!@_+Pe0FB z2J@i|P|!6*xg{S;Gf@^0r8y6ZTQ_wl@=GvS!;5mF4xgG63|3|FIr=)(JUAc95G7Yi zlqYk_0j-GfdLvPG@@Rxu>^Q6kJ?kzQF_;MIfDUeB`1xyZ4bOlt;s1`6e z{M3uoBxk#dFq+ot6?0&N{%IN7f&B{sz70F)u%5UN;F;?r0_)M5TUiUS7?Di--If`< zj*P*NvWrH~w~QTx)J+Ei_IfGs88$g&o3@m2RB=S_bEZ+W08o+s;on z-iu>h%v-Iwm9@|~k<2gtFF6moM!#dJH8IT?0;I)5t3mpPHkO`PfY_2jlN>xb+=T($ zFw=MenV#KzucU}k%_JomonzxJ!d=`YX)A=w(AcxE>QCLqr^H=SY}E|4(s38B zj3(OVvn9L=&WC&g3~|?75Fd^p?s`_@E;RUtPe_9w+<`_Vlre7rCB$7Qll-%Y$yg@w zjb-L$pg=mKP9?x#$87$ zKibY#X#);{G$l7kO@Z>$3y|<5?Jrb*2n0WdO=rhl39Z+q_26qLZs&Mhjk~_6_?-@) z3A8PtRLWgmX3m@;EK^7_N%5APbS+1mRZsMmKaZ2eQmHU5&N`@3#%2pCXUAF6)L%a? zW7DoY;(J1DKMYx8YzPEP?QGiW%*s}xGtN3!2xi1tVb*;f7?f(8BG@?V#!ZZ~-U|fA zStWSN?8BS+NM^@bM9Udx{c&8JC3*uVft)IfFudi+nct98*bwc*YJ$GWzU~;Bh(6UL zpP}1@poyx@fCuBO2d`l(dlP4c(13aI+g7C{H~U+`+_5qgdycbgm%?P6&YK@+iB9@) zmW&FnmkOH3@_QF&eYash$5~@oG|ZYI+gHh*0U8(&SL#h_g-+Cp>1^&1>p3&P}H$5{=mo*id#MP7uP zXNa?~My8q}2CkJYc*cvkf(9YXvxFFoi)?TqQ#Cuz;sBXS#7d~-$@`V8Rgc;u1WAgJ zDB5A(K2vP>S;kpsVbF-M+ONXYN(5S5$j1d*TOf0c{xZRa0x<#pXk*Q3-;}bYu8au{ zw3chFa#z9xZ625~FJMe*fm|6A8fcxSwZg812`dZJtF_8p850_4z3_jey%JZ#gzLHa zl~$YT3YpMA>pra(awSZ#8Ui?f`L{Bip=!$h>|zZ`Vfpb4VUb!*mz0;tqHDx z>S+s@9R8yQT2tu~V(Sf~Y0X|C2R0aJm4HJSz#_ndfz}aMagz6F&8@73aD_{_a!jM!9$GZ~E+YvYWlf`Qies-mUgJcM$Ca)T!}q4XZ7;T|?{ zc|$DYyj9x50cRg=JM4V+{q}7BHay)$JQob=>bRPU!U;2g5%Wc~#1X{>fiF>Qg$ceP!p})reko6=AP&Io-6Q04tKq}+j_oTM^ZojqZ)0-b zkN25n?m_%9nI-yOR02a=D2jwfatYoXegba|KgOHiXileOx;^%1)XG}t`OkdTbIFhV zgJ0ZW-iP;1Lc-o&y>G()jiPnC4Qc2sj=1DE#lEj$v&Y-_HJ-i=lkS(difJ&> z9ZyX(*L_I@h#m*^UrbWE|Ll*-6XqcOe#LHd@FgI8NOe}`y05{`avn=jTmEib_w+Y&ES394Be9VD14F0 zd~mV{$sn`-c#V61f@_g|e`AhPl;KT}XOZ-ew%~E$*3QDr_X9!8$s;|VxEyd!x%TaP zm(_k2K18MObI7~-936R+nE? zmS)R`lUlAJ$vcboK4G?fO3M?JbG)zM-fwhw`(MTT94Gn%9aMki$E^AyCgF7aC$D@V zFU;*7f1GIkesg+yW{;n|!=hsmcwZw(FlbMM_HdBnr}mI`covr)ylWgk2P2|o^9bV0 zZ`9>M6PLwl5fr*`gu^!ImM+6vTWe7|a*&!~j;Yyv3zsm*?}n#4D2@bft0kIRAaXo) zAx|!8iKf~<1|tPpTVTxEp$hFtcO%_}bQjY7NcV!_e&l=60y+YFo#4x*^ikle;N=Ry zXaI)*>i`QK=HtS#ULHy|wHWKKARHt$UTemmWas@Z@r9#7X_ z=a?qKs2?tdqZ(rEar^-vu*~|RyI4Os|MUgp7S^@ve>jKz#V|uvkEL%T`iIvcrb&kH z{7@k=dW(!f)B{!H0sw1$dK$amge9{G>I55s!t!+CfT@(s66g{A1GIL8u$Ejf6)wS4 zfMB!a^(Hf4ny)1)Y@z1zk%KIFc%S&({k?qSH4Dy022;endt|6_>#a=uHhDJ@D zGStnjWmW4TgNeksA@}_jVMdU(Xz|E2H~;u1Iw}vCBR2yPelL`1&{UzN5zpLs69UQq zqCljcTm~F9Q;S6`nO{p780lbZi#lx0VtrbMZc>7y!&jlR>C2V!lKrvh0`B@BiVApXcjq$b zBYQg&U@GPv2=LOgA#?r8J0G3?X3ZskJ_ql+9*e*cbTJRo;~l7v`Is@qtp57*$sD%F z`o-hsOSZr7d`QuM6*08nLHGXf5M(nd>)sGG8=Be|F2bNf;F2!()8Z$m#pV1l*PmoK zgBM>O>GF)e`sG3YL+=1%vOu@uvqdYmP=#sc!17NtJ!)?Efm8x4Mn_8UNF9rej`XS( z^w&p~r*AXLq(%w5m=(fg%O8e5WIjoS7sT=YjK&9;pCT_ZDg1z^C8Gu}8c~QwxykA; zP+-IVP87bj+!q)A<28;N43a@@qv@)@~yP;=Sd*hzSk{hSQk$>u{9FVzdz% zf4vFl7l$wFT|QXK2Mf(-cpR9eZsYYy{bP|o{I`S5oB;dI9}3zR>px>+jk7OV@e)gm8fcY^4XuLZM?cv;`Mxk{) z?H3q;G-^)bo056gKW86Wn~S{o8&Me|q)QMP8+nSkYuooA+>LgycQThzN-&AhPBX(k8dwr7LxfLq_P?FC+E z0Es=@i#qMXID29BH!wea{gU;DKC~)|vk!i`+t1kR`R?q=)lFupdFE5nS;#bJe2_2x z0KYN?bJ!o5PkW{B=lvxxpF;mMeg8ZC9hr2y*Z1dqytKNx@ptxSzxX%#^d0B>WOQ%z z{U*^ry4BG4n>DjH`Yr~TYJLxH=>+1be-r&KF>kzKKk0Y5ub(=9V4LAm?b?Lq;akBF zi$Cyag5=<+Hl1+aQM}3eB_N2g2TEXP{#yD`bGXK!nJ!bNBi3y(lZ(1xsk-E=PTY@> zjkf_a9az9{ujUV#4d~o>d+?D@N|!V2f&aZKu+h{F7qJTjkj9k}N6dEtt^1eaiOmvp z)470LRglKn1z+oLUEeun?P=%%UXE+(6yNkn%9;SIN(ZT6(0o&;S%-p=LQq5z-2r)_ zuEVZRc+KvcM8@Awk?yyk9Nk~iW$t!u^hG+*by#n-bvk+nycG9i2-~o;1rP2iq+i^Z z(iJS_%ZHa^OwuzrD6h=*4cCaQT}WdG=er*(P$04yNXl-1!P;Y4tKr?1dGR827!=CP zTSW(CAW$ zT>Ri{&YZ|-oSWoh`h6}=X>1e0W{!>gPRV7S_q+Py-)na>)`NX0myPYb{%6469q#Pi zp#g$6zV_E2E^JtTBAXtJd>@}wy9fWa7W75BLa5Uhxr2}BHip^!hqb3ewy5P>hS<$d z0))>S%C|iQVRodw37_0szp3wnLR9ejE;yVIeE(xtIqWZAW=E=c7;hs2i1mjX-UtQy zBABH3n(a9!h3}s0dGS-RwDl(=J3)~r=Rh68w<#in$wiW$SePipZeb$w^O31(%&|n| zdE~~Dk>?R%!Pt*Z*-?;)Je`a@4a{WZ=ZVP1WMpIOSV7Yvy^?H_!E*51SanH{<#8M%WU!tSYMqZs#L&)#$d$9IM z8?zkqn_q$*{tYNw$NHy~HeaBCll^;gBx(5)EI;vSPqKIOG1yzRv34lj7Oa^aaGcE; z8W&stgmk}*3<(hsP|DE$X zL4&}c3*(K(<`gvY$O_7#xwz+a13ChQJ@jLhW7pehufWgp`n%n}exMw6cd{a&*Lr??Zu+GnJZx#nby zKV>UxlU<*)zmMwf$MCp}Aro)3@BjuLFek~Fc4-p;;8*NzyO}|46EocotJ$;`o!Vqx zAtvteti|6hME297k!9vA;Z3%BphAP@MWEq*`w1iOA{vg$`Z@K?GHLRD(8O&N{MC^V zstwQPazHFzD|jy(;^`W;nud>;63h?%Te<^g!10OeHk|0~N@nu#5XV5;5axTe6Z##< zZ&Ux-4}Cv%PC*m?x*~Z0oXiXIEH{@;tK{{%GKQ9)ml`F1U*u%rhnJ4zK|A(u-siZ` zOB9BG>{B{B5AEXnOzb2?)N~yBEw9JI1aK;{&)w!UU=K&6O(=5;IzvjIAf+=1b&{A# zDozq_(vku5h!d8jeIxSnVS`_1N`PuRfS#?!Vq}J?_#y#jv2rxIVd+FI&pagMZdf`^ z0*wxTRT`x5^}Wo{(1E5PLhUMq z&|yvn_Q*UWy8y;NNMIxSJ5u6z|4fpw-5+@iF!|@)wrE3?FO%LL^j^0$bC)k2S*g<5a6BW_}`*i$wi&$ zF>q!-28{UWR9tl75wk&L5$QmcA#HE!-;2r2UMDD|%{fInDDT&5^l)|U0 z1{>3|{!zY<2{``247KhkoN~`t>r(~h{A-8>|1TMx3NXGNg(uNyM&3g4su{uZ5nR3J z2iCiweaO9EIV$~6mPn%vh;5@8E^em14^Pb`Xyxndk;GPV){QQjmH1N;pT*xH3086b zQxV2P)`y(ukD2WO4-SynI5sP(sO-SH{_6b)j4f=o#uS}q@e}M2J!KfAu#2<`2 z%CR2w;vYh1X=yMjGT0r!9W>YLD0aG0JPuNk@#9d%rwGp3;%-E?(6FP*VK$3)wvg7y zmr;2DGVL{dDM>26myDRl*f?EVl!=TaPh~N!$mghhD=TvX^q${kHts@5S#A|CIJ^ z(5n7erhkeF1|Q*QP9MD zoHUpL@Z;n=CfWY!bj9kI%B{U+)TGzLKfM_dj??6;A1cf3pDx3A55V5^V?F7GG?3=k zLx^n8)9pLpL>tzk7nX zuW5hxD9LB|yB_Ykj`w%3a$MNXo~m3-fCDG~E|=2${_Z(rXmgyu`vpoO!{41xD%Rl} z_Wj*woTq7N->GiqkE9GIjRk?_Ro6`ij=z=l<*XyZP>Or`O~C?+=equvzz$96cUiwe}VV9cp`i zRIlq!!2dmJkRrAI@4GZq*`xnktE1}szteK~zo)Wc9LZV#_wT>dEL1hy|6P{p|LW`R z0sprJ)N;764*_ibUwq}z`M*EQ@_#><+yDI(E9dclpA#k12{q3D6<1B&Kz@R* znL|)jc(mM1U$R&C@2~!a=yUG-_rFyRLRsf!x$oa^F~@6L9kOt~lNmj6f!1`pjR^2J z0Pqm(E*)|>C&61+3@UK({y~`sc`n{pSFmAxx0~COSKyCenx`tU*koS3KN0$ub@6_# zsAQknh8wTh@f76=#kTt;5`XjT0jC+;u8cnxdGnQVci(=y`D$|7b~C}=xBm^8F@4x3 z;@@vuYSt`AN={CK@29%?jJ|n&f;d_4R_~r>gC^Kg}E^LhPKIOv2H^=@5Cwkuf{3~$Fa(sMq4<(Ti-yEGY zzL8HHF_5#gANi}5!0zIkZ%V9?KfXywc{;w?hC7EjKwCY}dbJd~>v#+%7Dt9gM|K#M?Q!;+vO6oAwmn zto@+$6|MbS#5ZpJhOcGzSG=B}cHcn|S3f_?($PL0)UTax^t?AL5mtU$!^+S@Fub{}1iY-@YIB z=iPkx--uUk`H84^?s$dgN6!7;h_pQ))W3(;4}Yx}{#rMtUB6kZ?Z8Cgu<24}iy%!~6H@2L5Zl9UR8 z^H1UmBHsPL$smQnnr5UZ+Wo82xgvbYW9f|;Qfnbh{5$ZAJ)I!^W$>rrO9##uCaOumAbY)ZFpPc=^7H z^P8d{q-C7%{O0gSMabReXVBRL=Qo&7IG@@d_m%c0zIl!3H}$`RBb{+y>0j&oCXBPe zbbRCPEA8R@Ci!D3ZQl6iqrXL)K}lr9H~)+0H^0J-i=6R|yRWq8^P5lLFZzlQ za>qBcCD)H`=Pw2$XEU)HudFb{jpSDAS|&TqbX z0UNgGH^0H(EjA}Z*|N@WRuGi=;iL_sl6~fH-`S`0o4#w_#{awXn;ZX7#{awTCA}>Q z`UjIE#8*-C9c0JHR~h3=yNCB5(|^lbPjCKkzSPV(y$Sli5?kdtzX`sNqvv_1^o8xk z2$#&$Y&MAtYDdWi9S)`X!Yts1=BPOiK-px$EnTqBL&JAFylq{74-*klS+qF1Z77{9 zay^{y$HyVD=k9z8ce)(?8b@Ehx06jDJFMCCWz!XS+e&r?=`Ty?a?lT~Cp>=1x)0#w z&-Lcv9OW+l?bH7D>DSvBbNcZnZo!W6@^R+%T>t%u%>Gl351_3PyP3e(m0IyPSzLRL zA1g0&FlC(K-36ZQGX4n5!MZ zw?g&1fcE>iLxQwGMwl;O6(ZQCzJTB4lBZ2M3QEyK9o~%in)FBIe3NR!@4?F6MS>^! zINi-UU$*k}@7rbTC*?Cw`cbo|hvp=hg_cm4`$@K;@^#5%siV&QVk zKfX`1SM%>C*^A|Cley<#fM1Ti*p1?6JNZ z(!GlxF-C!Y?^r*`V>j9Q6Znb%i!6S>dHlTKE8Okl@3+|gP(O8?a@m9Qd_0WSZrpfk zc~IYv+Kcw~9-s3aa}($LJlmIRDtYfYf2%KZTlG-=~KKt=~1hpO)pmDtY5hXjX^7|mSel0kBf%^zBQGxZq zA0N@9WZ?TC6V0zc9tuO_GR+@6E}>zxomU_Bdi)^MoZkmw-(TzZLBtHhpTnNEZhpV# zy=iOKf_}O9H;jHm_IQJZuSM!TntpxL<`jq?Lz44$lT2ll8=Ws%Z);|~KLY!Vxk=aF z+5i0WB?Nxd9!dzq1&|_%!qsjA&w&Td%izZ&A8I~0tz7D6;Upms?L!&;X$6wy(7hQ%TiY^S9+qo|A zuPVq@j3J}<+SkM^YS@)?1R`=>ME$f={r8Z^Y^V)14GL z-aZ~O|0yq7%U8aEYmQ-%!1oDd+r*s+^(G=eWK-rvVl&2_meCu3{!y0vdke|Ag>-r3 z@7GM=1PZ(Pg82q<^ZFcLMPD+SSoHM#1cDy@TMVN~WZ~MIXM_-aMS9tp^|Qwdb8DBN zqeXd7UAs2r`(x&*IifRY!pYB{Zy;`P*oD8{ZYRzUaYgz->4h8^&ZbdwD8@?k4s0n>;LCxhL-zq3G8@n4 z^z`LlB>E%1(OG894+;Hv&M!BA_Y2eZFX*ox%h#XA%}36sAri(Db0~9tUCk=HTVHPi zg3MJOG`~*?a8_Z`o^xA{`J?C4I)7Dm-4qRUmGpbfn`epI_nK#zu-|sYStJwOboObJYe(?8c%F%DA3;W0`!Z?!OoA0BX z%cl9qw1VG9(fe<>v6SQcDA%7!v4BGM_fbkfU&lk&Q>4KUb$JeHP?y;;zCVyxcfX>T zL#F>M>h4Cm56*s2zp~qhzERKE zu0)Rb7<5LXLH}%mc?*UMOw&Jmg)56FKXtR?X5T+!|I^=lIi74NDcU$;3$0r3q#L0u zf5b|^(|E#mik7X@-1rwFtiQqb(T_d**CwefqLJYI4f*n;4d6S79F9*DS2npJxWZ_% zP(*_V?dgsuup;5u8XmXqc?8=lLL2EG(a)L4(G$_6nRoJ)^#N<0zu%+|yT9+0uF)Ro zFL>hnDmtj}6^Fd|$uebR46)ou)Q0Y$EAoZWkaTB1{XHxE0V}y*qGPq|Tzh~0Z8Nz4 ziErgYANala3FcdXsDylbi~eT$i)s1W^#|33afsUJ{*nk;`)}nHh#xZQ^LsrR^=&*A zSno3MGjL&AI8ML~jwkhN6!M15`Kp1NP}oNmee&0$A(#r_;SRiwKQffnjkkWkQpb-U z7vb1ln_eEdM@I$3Kg`A*L(nsQZ}JTC#sTM7u+$?WEr9q<6q4F2`RuX38MY*1V9Gt( z^e2T7gQa6>-6ab+SpxmuN#teZ@!;(A*UvEIHw5`%j=sA1=4c%jSo!p1@bxLUzmXZQ zOfVOgrqyr0Z!bBXm>^gaGyEO6N<02IL8iC$1nKJB2LtDs`%9#ui~;rgS0wFmo?HDX z%G95|E|;tPzr=}z+KTjeT3!>2o`hpa6~wT^FQ!@I#!s{`x#&sQB4|b;dNhe=tPds1 zANt1+*j$G9(j;*w60+Z3xUw(u0?hbFp?1U>GBXrs$eewK#aU|02F#naY`}aNg24wL zH?057Ct5tONm`lUjhqI+Z}GN~0w&^xG5{v{fJr!qA-&1s)TGlZlaUv!9%tM4wEP~& zF!0NfsrAph@7+Px*qjdZ*V+euZY^Lhb-aE1&b4t6a)b=w@PAFCKmcGV&6$2# z4>9oR`J2;zIDTL|GWiMU=Yx^|@om%I;KMeIdTVFNcqcgXF??m53xPT3&*#=>&^GMl z1?)`PFpT5J$_IazoYOZrr84ym!U+}z=39C^_Mq>g&j;hIBdE_oe>r!3%2zszr@xNo zm6z?$Uay}PAAQsQ2KZ5Wf&HgR=9>_APXDI&GroVL^4R_Y=GbTJp67EcHDl(r>8=8}R|Q9wqdNf3=Jo7U%zsczO3TPL^T>iXX<{1c%ec%QAnm_A`Qik2hfN z-1;gyQYYEHr#&1|C}oAz*1-1-^})CJMHlpBxUrx z$M@lkSrv@QPKEa<AFsb^Gw-vkHz zpFxc80N8+!gL+(F#gXJTGbw4nxIg3Lefs;&|8DVdp1=3}@(GBKpL!?9O5)>*=7tl+ zXvqSYwy$Esun61n3t_xuO$*vnG5!DQ@$uDoZ%fT4W7yXvK8BuYe0;n*#!zU9kL9lt zr`=%A3%9;x$}_mvVExf-f5eX!F&kmJJ=$;>5bL}hOKUD!z%kVU)=%aCtHr|;%^d1p z&i#!H{nGVT;^B$G%r5Tmt70YX? ze2~S?1g=;ak7Ys{YFA}J8tW?>8mp8K#a(lCO=W$PhsP>Bls7z-4I~;WYFBu*RdMj@ z)zw$kl6SA7VNHFF*HBj}IXndmPg`wORc%#6btS5=Tp1@GAN6$&aX!(wa6!0DHPtpY z)w8^6g*U%yZFObU*&**x@A8#3brtc@m7&`Ac%cR$y~jI$?t+EhmwG}nhE4<2zT4pL6sw(1Dp~m=%v(H{zRT-~q2-Q?Htg31V#aCCcTGcVMRmohSVL8PL*2B>`g#ux#wsfstGp0?$b`JJLm!yxo#vea?yH*X(OlJv zkXO65KHfZQrdQEONTbst&l|-3ZvLpv{EgUrfoAbyjF#ivO@2q(JoehTG>ndwz#4atJwL%6i5^;FV zwtlC@&xh{l3}w;rW-9&i^Vg=Sl+LVbuB@u(P{wMi8{;77=oDG}J)k@G{O9-aXSew~ z!`655mm6s;I8)2RwtSg=hAf;Y!QVtr1$mck*L}*zJ;S!%&>U^I-L~hRu6>s;umkDx zL0U{N{@}cYXQ=Q)4OQ1PRX0?vsj5wDmXu#z)lgei6T-x+SQDzNj|Unl|MQaV1zwTY zZU-Lp3d&I)wm9e5^aSq^??MaX z+Vn%G#W}a)p~f|_v(7wpTE&VL4ONYeGQ&c%W&&GUngD*nNcG|ZeehQZxn%GYEDWSS z7Fl@ygo1*(ZagP=_qwSeSmn0e0Xy)MF7yrxdu5o1Vc4Jo&zsh`8ukO`5-Imww?^~w zj5V&VUD51Kg9WLYR=aju{Oao3m38TSWs}XXou-KuRm+=J#a3c0s(`h2T4hB|&8miq z`qh$yL5bJJ*VSuTLsiYR3b9@PN}L4bo1HubfQeq#U;#E;;Lf_FYsCcaMkfvbh`@0d zKkLUo*xQbO132=%=+XuA7t9UKt81wD?IuRdTS@W7!RA4DsQ(s+E zAuB<=uCAu>G)$SADp>7M>5OSJP77Vi$F#F%c;_`$*Q^LtqEbvi4gBvJS|984jI!x7 zXO^B-8Y}Tme}8C3==6$)%GIjKq0`qudqW=iWUmYib|#+}dToGDWXJKNlfWllT9qe` z3FL_#c6kPth$Bw&CbL_w=|}mvH6++j*k*1B0+A+u@WF-4Ys?~ z5^AgtKk&<7!`1K9GXKABd2|5kv8|s$vDvm3 zvWRX5WY=Gy<@FZLe{8Y!9ln#X%R9sv5+FRT{;gT`v+FxKa$aE!$JOt(^@|)K`*ii( zlQWoob9OTe-EY_taP=sxS5wcemmO|>I(Fim!;kd5F!(K*v};#6()ESAcGZ-F08W-? zegMytcJ11TXD_bVoCf$1o-O=C$GEey1E|?tE%=tNmB?AU8@F(>5syeK0m-E!gW!noOEN9z!xMEmR$$V!#Sfa|IP8V1+1E14KZn60kr-)F{CV zT6DmQ;aW9bbJd?0$?v=NUi+LqXJ$_xEtNkwAIP`!U2Cti&wj4`IQ#5}ehe781bVD9 zLnH@|5)bq)wJZzoN01v40EU3wz&NlM*aaK}_5m%t#4-f*07rp^K<_d~p>I0PI7jsh*bQ{-I^J)j@x1BQSBU{^WP3HTzg^Ww3w!c~ytZLxA- z61WxVZsFC+JAeUT7xJkexChvamqs52KLXrO`F9EIkbVI90W1X0LcRom1<2=KU>SZ7 z0yj{;FGc>~cM`Z8zax01dI*>VzCpj2SdaLmqIPT<8@Il}b{Pt9ijqL;`-VYu)vJrMrK0LJu5A*>a1xA1`0y}{P zA>6MH&x79Pt8Tn`)pZUqJ!V29+uM@bHR5g2Slcwh`T3ujz(0?UB?!1cf+a4T>Gcn7fX zDuf4i0$&970{;Q*2hKu+odlKv16LzFFbJf5e+gh0uot)k&*5uAJ^?+=C}-e(z#(7^ zI120pdQX8}pdUC$a_CzvC~x@f1IB>?U>7g~>;ro6ov}Bd*BygCe)j`Mfk~kERD{0< z_JD=J5HJ9Y10%pL;2^LMI0763dfE^k7z27wgI%B>*bfW=t#*V527q(k5B)go;CCUg z510TB0SAGjK+k5_nGJoQ9~c3KfH7bk*bD3edbS`uFaR7PIdB#1_X0=pdk`3>`&|or z?}8lY2X+EOzqP6GQ#4jck{w!;qT z0lkHYFVGK+0Ykt}U>w*@@^z4Z7+_eqbGV>!YxX-yYyj{0;)W zXCR(HKX4Ej0$K^!2NnXmfB|41FhPiLBfufzKL&fiAkcdz?hEt-JAol!H!uzy1a<*O zfPFyg280I|0{6jQCvX(MyMcKhKzabZXCa<9!X7XP3;|=nIIt7g1?&a(0h7QXp!IRs z0Tu$ib8ugvALzRY;eiog9M}o$0(JxYfc?NBU=lbF_xF7QcJMm_^q!6T0{y^FUw!h==Es zu!G-0p!Xcu0s4Xcz!1>d2|K_DunQOi_5l;XAz(Lf6xa*gPS5oz*eSwqpdUC03;{j2 z!VWM7>;hW1!4A*|90CS`qrkvt5FU;;B^Nnjl4`5eLn1He9D1ULjt07rqHK=1jm5A*{^fFWSv?FbJH0K0%e z;4VDxAh3^q-vK+o7;qHW3w(~!;=YwgTOAJbr-?|3xPwx0B{tT=tlVYNH?G#I0y^@t-BE(=mT~E zgTOvu3^)W#07rqnKrh}EO9K5s&pika3;^T61h5O(4;;nwc76qR@OuO}1oZ6!4-5jm zK0F6-_IlU@`sw#qVGo!9#)197E@0ukumcPLhk%2?QJ{4n>=Z)|^aDL#M|hwQ7$^RI zga-zIeZUxS2-pc61@;2HC9ntd14n=%pzj+94-5jkfC*q9uopN4Oae!N);AHp6!!u8 zfj(de7y-tCoxm<&Kd=ut0vrPRzJ>6>7|`nn5A*|ry$BDC0pq|C3-*ax&8KzN`B zSdMh-1dh`02VrLc>;nD37_bTQ5nu?vt%qO_SP1L__5%BWN#GF4`(Ot+@-XZyguO># z2WUNt@W4V~92f+40b{^EU?*@0*bf{9CV}2X2=^U?2L`^2@W2Q#4(tSW0sDb{z!BgO zkj`r!1^R&AGK2&AfiYkR*bR&W2Z3EcYY)N$1Hd6*0@$@4@%tX^;CC<3yBPL?exT?3 zumkJ_#_9Lt&>;uMtL%>epDCq&c@5OxvU=J7shJXoR9M}o$0`>y?fP=sxVC;wR zM`6DkIEvr>K<_1pH_#9CJPSR@gTOggA{;Qj9CBb6(E1VlGw|KOKK$+n4iWzx>;VHn zZ#nJ{^aBThap(trjQHVq1Q-VLcH_-b* z+-GoXEP&q;;1c}q21fAP`WeQ-zyPot*a_?h_5+i^5umjK&-HWI1NH&~z+OQ=mYixBfuoE7ie9Icn%>vun-sk27wV^H!uOTUO;%D57-Y(0FxvKT9vpTun?F8 z27raXKzRBMOaMLqh491!`++fF64(v2s_^*0LZI(Oga-zJ5nv3M0CobqfxWI5#R_g0rdP5c7QQpKd>8^1P%geaykht1bTjj@W4Xg zPRb8p1i$-%3E&{G4}5GN?BRDOu%CEf5?J^${6ok)fz~R-1K0!{1mKZ(9xw@91@ycEyTBN5CvXtBo8-XfNd79q z0}Fu#YmlzMdB6m46>t#P1oZtD;eiq0ZeTC)IiP1h!UF@qg3Iwdz)p;B&wv@Bq;B8o~#WU%0#yteorCS~*)iIj0pInctZM9q_a9|H#Q>V{=KT z{74UR)IAr>8XKEV9IiR|A63kH@h@+V1@iEZ<2);ln|0ao$6TAAur4|E;`2)9kf$ze zH31{%jg6f~#M0wDH{~ur_K0=3A{*_Mn z11vY)zW~pjE&M#lW3X=uzY6lme6zfXf>U1MWs>CgWm z?fJ<*&ixNSp1d5ssV@J3BQKaH<;UsrDn~vKa$j(43|*L$KG!+&RgecE zH>H0QbBI4)T7;(S$nj%dHR!^vnk!AA$TVU9Odn z0;<37!(5Xt&y^WS;pyAL0mwI`$YlbOd==yg$dA_L^7gX+Lmq>Cxi0^R6aSr%cSCM^ zzTJ@bL;ij}d~VQwzULsf){c!erpW7T`2om-kk8cR?{n^7fJTkvC+PCkj(i@=&)4Pk zj(ip55rjWomtWz?n^HbBK??CvJXVcn-O2~IXK3BhgjdTA7$on8S)r}4ge~KPH~G z`Q5ra_f9+S2O!_O&YV8OPG|H zr|Ay6ZEAJWrvq~OX5RaBd9HLWME*k_g8Up^uK6wlkk@6A4@2IUCI0!x;Q1jpm5(CG z{a0prewLfs4?4529N|rUOUF3!9+sQZZvgTk*f-`sRD`F_|h((|uY%YP1UN}nEC?b+Q~q^8z6uTXSvCK z1LSjH-&B4&IJ_x+dRT5MKLc#v6#rq!yRtn$2DN)3H|1XuydF^4p${~RWHjX8e#NO+Rpr`ylH zN8EtI(|(^W$k(RGrSDJjO2~&GH>F<#Nh1`_CJsjSYz5^^brQa~* z{*7Z}=jr#4I_Z~xBI*a^ru!E`PT^^sW2dh+w^7M*Q~T4va#Q>}#*z21-01&6PVqlm zkDoU8F${SKa#Q}~d(rI0u_s{(=yZkgjKBv|^chtf1i}diDo&4#6 zyd2@T>+)RbU{m=UfP5F^7wYn>obbbtd+V@XpvyJ?Gyf#yFXW}Vyu}G$1bGN@Q~9lg zybkh!9{zeKd;{b=Ajh(YQ+{&g0FUC|0r_6YP3hkQd0!U!0OY$MM=^KoYjYRFEH{dnxjpp<#V)-k&{oH+a z{3;F<(94Jw4X_-|oFmufwuf16s^9rY__{Xp z{Adx&P4&AH@;Jhqkat0Ds^2{ESj1t^@Kp2ygNadLXZZ+|>RIu-w%C z3`0JI@TUCDN5R_e>DH0h~=jCrxNnMt>*M?V7aON>41F4HuLlKu-sI>1|aW3 zcvJokk0Z}7l;KVFtBB>M_N9{LruL-)^1kcL?MnyCP3=byY4nw{l zawLP3zuMk_d^8LsH$FehP3=o147}&dUN|S0J#@( zQ~NRuxgT;<`;vbq`cKF=>+#Q(n=hzN6hYnx`4uU0XZ%tL`6%S3^4Gv}Q~GwW+?2mP zkmqe58(Xa3KkoQv1CY;w+|+&zL*50ssr|~wz<3mLlmAr2a#QDP$3Kwo`k49o8dz>> zzd9iALwHm9=@~~pz;ZmMQ~$KR4a1P{g?&@`%ZKAda#Q}%`iJrI;bZff5KAg_bm)V>Wv9@=4Uzw#0BF364ZKalr9ZYrOZEWcil zf3Dn|L+xh+EUyw14qwS3Ay(sbN)3zJ_mAB`Rd^Cruy5%;Z5=Z$o;Tyk`F^(4!J3P z^UsxXQ~YpfiLh@9UpbDvf#tWSq~B-k`J)cV_h*S;56b~g{>iOPBL4;;7|jxXnB~U& zKM(oyi7fdKxgT<4{zKjcxvBm(K)wrdQ~q>7-j_w*L-+46haZ5v9CB0s471#H|9lku zP!{_|kk>(eiJt!2-oHx7cR>CDU7owv&ff;e{WoVR|70I>Q~Y`$k3+slx38@~4M5%n zx$*f~ZYn?d=y-M^{94^U@9h%)ALRS9#IF+aQOI!{r+jMbI}MQM-D397Iyk&3{ypQ! z2Uu>ZU&D~^!u?J5^DmHcQ~DG^-j^kOCFI_}zG?J78z9ev+?2l^9R4?Y`sM!8EHY*X7#!X(i0@78I#pQTiy=iUt+diOb#4GizB=ODL`N4>g@svD=@ns?hY-pU~C z&^w+TS^7#Z8i{{EzFoJYtu^~0e-wSef27EzvrX{~L0h+;smChcDOdf7Ge#amZ(V%Is@(LB1Yxll?x(8z48?AA_GRLbufNqim(c}z1LPh({7NVN`Z)Y*U7lNM+aH3w zBa8h}$nVG^_rl5Uf!yTF`yn5O{3QMU@--CEzJyqQf-cwgV#HZ)%AYQlo6^5;obW>| z$B^Htf7)K@QON&+_?y}XFETd&wk+iz@-oOx?N12uO2}~==lVQ~Tw`Br}CKJ)a+P3U6v3LXekrn)~`V%T3|C$UeeT zUCrcc^s)RBU9OEqhaewDcvD|K3VFw;Gxybnu<6AhqT$cxzL@&jGU)7w&Uk%o1bRzw zFH>G8Am0tSsjuyZyy-LMzP2Co1CX!N(<0aL`;(B*`>eUIwJyecOOTuT+Cs>;LvAVy z0myekZmLTW$nS)FeM%*)N2A1LUT@CIIyhNw?}Q96F+))e23CYnDzRgw@>Nq zH0$kyo*(b@Z>Q(I!>pH=i}IoLcA53&L2tj(>ox0zpts?FIQ{nMda3ER9eU+2zBzWw zbi{A3o__7qRl-pE?SkHH%$xlOJ;$JCkNQiQe4<>@yW=6~eV_F9X?nRkY(=s+3cUi% zb3I0SBf8$U6usHgkbW<7c^=W?r&Mep?AlJJhC6l zJD};&5|EaT-aH)RGeho!&ijX^M|+txy(Q3_bFy3y%gevhdG6e1Talil4tm?6w-5i$ zo2BdRO3|Zt$46($c0R92)1x(D?Y@1`D>xPB-_Y}xX?oLIXs@f#r2FoJp7->zv7gX= z8#KM#N0pi=H+d+3v(J!cCgknZ^scehgx);pg_Pc%X1x&f?mTmBtd{P(Thr6Z4aILe z^hVF(`|i>7T=(4ty(MRlja^Chp40S%KP?l6?9sd9JLZmw`Io$Xy56Vl2O+&t=*>O{ z=YLT8jp%x>rs&N+0?%6{*JJbYzU1UDy{@k1M>+JiLl51ImFLy;T+h)2y}jphy3En^ zTrNH3)6xypUPku>ik{hYiy zO;7luPCEOcH*W#X0VR8JU5^Hcj@}07?Ye~Z5@x*}(AyUn8~ZTX+o|ilFU8(&=iItC&4E4gxAGF2Jy&_o!ytM$^gcs+ow^=P4m;^G z4|*Z&Z*3yIZe5QSAsoFB^xlA;sK33sUTXSnhhBL*$FE=4OD$KsptlEl!rq{+mul}h z=x54%PG&UbZ&y4 z7yGkBI?vYiQq#E$dP}g6Nu={UT`x7A_dt*KJBf5zqU)um%YKR<_9fByHE)%sr=?2~ zWCaB%FW7%0o_D>br=<(j#CcQk%P{}1%I1^F6zWjoo(HG#HqW^w!VC1 z!DDAziTr|}pKZOGU-0_b*50EE5_7G4k1n`+;WcfqvHzhKF9jT@db~bW9{-3{PY~_LC>kLon!syj1&KLj`j5!r{av$ z$4@A@zsS1##DXV_tQ)-rFBMt8^A?O0S~!y`=)blo0-GqmrY-1y$2Im$lHVYcl!NduVnfq)*^bo zv7F=lN8Oxb)gEmSqT z-#^GJxa}nChUvHF;KxsoD0uKBtNX};7f!Oigd#iBy8ozx+h!qCOcjsEa&Ot?|w7y{aMOKN*`!lDP zX?b0M&$d`))8B8s2hJFDkGVN#`qy%-TXU$d5M^cN^s{XJX&2<2X+=_B!9RViwV2Dl{_iEI7fUf`CDJQ8y)EY|q*qtY z^e^UE_vK7~JSS(jmR^U-4#juY^t0{wjrW`Gn|rCX5R;A>+YonJxNVty@ zm1R}ox!1@#TVP=_cir^3^+7Bkpd1!hOEC>TeT}t3m0LaSNe*Z5f?4aVicImzR_5tB z^|&5owQ zJxA&#_Nw2E{mdWsh#79ECQ`m9rv~kbl^?r)|I+pQb=U8|x_;-(mN$GC-YzWiR;!^h$`(v0Rz;d;?hQB4_Cnie4RI*_&%dw0lu9F$> zm=_TLQhYxD#`2H2IvKG%A})3RH^=1j6rCRs{22dM@fhLX8_tt*p>!Dk^l(17Z@G~3 zcLn1*#(Ksrj5jfUj`3c`hZ&z{`~~A{jDKT1VyS%YnT&H87c#D3T*p|?xP|d1#?LX{ z%lI(k(~Q4he2wvMj7Kcv`!mjET*$bBaUEkl;}*u77(d5&FXO|EPc#04@ioT3F&?oz zV9_;`aW3OR#ubd~80#6gFy6%YImUY#A7*@-@fVD*G5(G5h;qI^<6Optj4K$|F^b4) ze=c3V{G!65_ixrmT2lX-#oyZGG9Ix<%2Ab>~{`+iIIGkY8?ZYTS51Lu-3ni?92iIUw?^ zzl^zgs@FMaE!Iy92vMQUDP17&>mI}p(wxa3R zrhuSnpCVl+<#Ksw35egSpL#Fzo@U9PC8Uz6Yh)=Wz1f$|#&3rfW*rqM6gyPxGZc!x^M%YGC#3ka9rHfsss7M)(h;1_%!~c*;M>4cdeW9X{rWiTC$0?$TG@Gp z`NFM|S9abj91X8U+xhhCec(m9T_g2XJU`8R|7OXnbnat2wAD|)o?)K0_UYFznWwFM z`t?WVY3p9fMGYlAFSdN?*NM#2);#^1%RFtx)33$M)7CrvqWwQ)pSIfRSA==mTBl#M z4~X<>E1iDPcO;0Xt#kTC`+W}7wr!s{XA@4 zO1Xri=(W5qp7sfm{u~!y#Ju0d(|#e+FL&`l=0h%?_6?DKor~YfeB8y;{vpz*t!Dbw z%{*-_(=XabMEbOqOuy(m6vXdw@wA_ac-ktaUn9&9xp>-FMEbN9OuuF^Pg}qAi}n|h zK5g~VuK@G3wM)NfpAqTPRxbUb?^F;UaPhR?i1?t3?_@sW;%VOz>Bn6Bqu{At`(go7 zur3yVE$i>hTU_6_Ghd2<4(WHZ{(9!GX1@ras_ybKW!1yuV&VdaeU5Z;WwD?=lEQ~e3JRXwt!%y--C7)O|xjrZQ7W3WA zt91Sq^B%UN^u5Q+`v#bg@pEluzL)uV%Ed5vp0>{E*B_aut#bM$w{4&tVoRBRoiu}f z!GCt~=Q2-Q-}LJ}%+ppk{dynsw6#sYqRi7)HvQTLp30A!FH-XkUjUzJyzl_)tNA8X zPM!jvV;yfL{nCCR-{ck6ABhG8qxApI`rj&%`l7pm#tbNsotKIwpTl;F4g4DBA92~K zX8uBskFvjk`G6~)+nBc&$os1Kpj((PWPT>c=S$2-*p83+J%aaG0oLypeiimkvVJG) zv%2+X!OyUg7t8yKVguokC(3kC`aupb7ks9CDF#pZ7_1K{(*jGGPq_FM%)93WH-M+{ zR40#jRQl6-|2fu7YnE%?RjI#d*^MO2RN42*rQSiu)@4;04&CGY7 zB=r}u{XXWc2c*8TGxua^r}t};SMBXp%t!tuc{Ts}dGIKrmMi6?=sebkk;FV2rWB#-F-BLgC zP^z6jGw(ZArt@X2Usx#hyMHP5mCYxax48T${^B#Fe&N@pKHNWX-NpPMj}ukC>|s9e zJ*kgqh^zEWX{U3ybZ`D14;3*xZK8BzmoaLH>Ksd%+ooR zq@UO;dE8E19w?}1XG*|%;D!F<0r6XI&d^uHr7AgYS3EU|8 z$2hL#vk<_9@_RM;!emHLH0l=io<{@n8S8?B(53vTUH^;&n#E@ zInKxNd@UgKUXkk7kC+cUFZr#^w-rnM&L<_W;!o#2QF_Mile{YDN0h+lqk8I-&zr|~ zP6AK*g-NNe()NqYTaTvN9|o^<*o1&%OWDqisrnZ%Kf>vcYACKY<~ws_d@f;rfcfr+ z1A!`gzIBnb)B99F{KijlJS@KrQh)Ggl2`rL9n6mu$n?igaXrg? zB46?fKi6 zlmAcZt9$*4`Ccw>s(fx=F75QP-zY;_))UMJIsMTMiL17p^?xcg)qOwBeDIf&SN4l4 zq<-)tlHV?@V4i1%GKoGzry;*qhk`~%2f_mGT+JNLB+p*rPS|zNZL{9bH*yRb3pQ{eF!ig zpaO|Y*|{0K3M&Ns=ONY~Oaug@p6mG4(td#LtFYC~d-h0u#nVee^jw|$B(LJLL-B7& z{y1?LoO1(Sg=I>>L*QrRo@N!g`q_TA)88xYEKq?im-qE>KF(%-7xRhVNc}6AUlx@5 zzHO4<#{72X120Kl<=3_EllqD8OJ2qQx6BWoA=6ov&;MiI$NUG_{))BIP9gWp%1#&a zy+4=sRe3(TTIvtJE_szNLGYQN9 zw?^vsa(z_xmogvV@~rZu5xml|wCk&^Zym8zu$Qpxapn^|9yy))Uoqd!^+Lt-U*I#@ zpSzLc$@!)7Wh?WEe+7hPmHrQcKOy%xOYP6&+ggtjuYPFw69fMTwjVhV5ZcQA1+_9h z7MBx0+gi?iH~ULUe-HQ*axq@6mv*>&w}x1Mkn2Sm+sUny_B+2T^;LPkp84QMCBL2Z z8&Hra-$(wMs{a?}lRuEWs@Lo5GpEBA@WRgPQeVY?FY}QHB(I)#{U)j3zg_aGy?Rpd z-2Pl@$*ZbC>PNVJPHexJp{+Zv^v-X|shZszwhKk}^P?dQFU zpQ}&us(u}Fb>{olfTwgQd_?M_`4ZP_%v)RzRr;LLB=r;QU%?F!*QLz+ei9JB)id6~ zeCL4VRXh1Q^CO<6LjQd6*RoD*mi7l9m-@4Eq+mUGWs_9}+{yaECj;WQvNOhf^7oQ| z58Jt}h2zijPU^XKFdzF_s-5??O8x#5Wjs|selzp2JyIXVTU>#d)Gy@rT&4dL;4{U4 z#x1T^D$Syb0718uS)%yB9?U4hP0FXo8(owinmKX z_zlUc=PHUzzW;}kzliO8nECFfC68ubTwi0pm*b=CTbrdFi_7_YS^qQ4d$|2X@eo%p z_!IcwI{eVUFWDmP$6g5tbrsK#GVkGj>Oa`nH<&McLF%jg`V06>_r361P9ILUd2Ht% z<^x=ARet@6`Q#(gzDkGEt(*?`OMVBJ&%ZP8`HJM%v;Dc-q<;UqWI0szsGRx0Gg7~l z_1nSIyzbUgsecOdoyyMhQXh{gu8!;Yx&B-7sy#UXK2thuyk6>${7LF7`*~=nsoeT_ z{z0{epJd+pm(*A7+pn4L|E1*5w&YdtVHtmqU&ddh+lAnT{l827Gg$v#<|E9j_@8t? z#UFkV_eXPCf5}Iro#a=fo#o8m&3xpil2`iqAC>wM_ODd@gUpXSCG}N5btChIuSdV)ePdphQbxYYFW4`x*AD8#_Js|C?c4G_kp1qPkTio5UUS&S{u~hvHH%U7QUY}Cs^A_e~ zJb$I~y9(lrPV-zV$~LPZj@`Pf9yJ?(bAR zJ-~eM-Al!DmW#hw7uzZIN4T9?!o25Gk`KHr?W^bNX1;f~M7GOrg!>m7&ko+)9*M3^sv3ezc9P=+R-~D^Zt9ZWuGg80*Vae0I z==vFWniq1f-@VNGh1`FgA=RxjKP&C{z7-J6Pnds}dF$trSLJPem()-4cu>XjlgtOd zF7@wbI~V<@)bHeSxR&`w<~>hH{ri}Ig!z8%ADftefq838>R-To`R6ztIDJ(4>}9_5 zD4D-?tp6+KeLP<(ty$Kww@W(_ZqHRZuVy}ZkJMD{z(<+)aeP$$eTn(Nkkp@H$t!Y) zv>)L1LA5h?F(3Sw)K}@bhk1+rne}Yv_uy4tGbLc|=h^u%X)=* z-wOf3oWlHVcgy=mUY5KnCx!RO_n{wVdgu>I&)rT&O3-#^QIa!l&mNIM>` zU#cA5%e=+)>qBg(=4(wqfVUIXtzeJ|=;$`&O2f%SX$N;~^lddWAW{UnzMW&g*_d!Cp2 zDjm-HCO_AxJz6PGgWl8pz3R(DLw&Qy>AedX2zx_c@hpm!d z!2FdD@$-I9@+zJ$F+cKM$*cVRJM&3bJA7K7w3Fa;R_SvY^Mkun?eAmW!|jYJZ#R8g z+Uee&s-OR`%4` zFki_2o+{7JfTwz``f*mbo@f2wUDE#9EI9i6^1g$v`gJDrzJEylGS+{b`532z;u{{9 zb^>oo{R>$CXUq@sys>ImkLs8DFNC=5kmI zp4$0-e%=qUov$)Kc(=5#(&s7WJHILUceDQNy_xObV&H!Up8Vbh_Isr@oWJmtj8Ed* zl2OmQn|W)8Uzx*SQ0dAABSreyjU_nfYG!f7S{q&Z}qM;`t?|f7TDBox&eUJIem;%#Uz7tMcr9 zR_X_jm;K{%wqFWfJtI>BE;aBQ*-ori-uF@t^uNp}zAAauZk+uiX}@#1HjI_yYG|suV?*fKb88v?UJu!zL5EdYkhg0(s#|f-p;(m%XA%|19;@KBD`6&hg=L z7-#(f<^x;~mHu_lOa0h@v~vRM-=%n7zmQVPdYt(Y-uHV6>z5Bn|D=cglUdA1z*9bU zzb-AD!~8zxlV6s+YA657eE;_(ui|s-3v7qeP3b3@@8o?`>)HPFUr7D_Z=~AkV!n{; zmnxqxGauU}_218SO8-mR3G9{pSDXAS}#fJ4e4H?RdT~c~$;5fLCeDssh>#e9kYW zouF%+JB#`5Z%g|}v$5UGTbvFmp1)E0*GYZVZv5<5nd38OpS*8`$4wj9{wn6XdA@Wl z^Vc)q`;4?N|KccU<~^@U<`hd_$Gj}<`%abZ?RMrbVm|RVsju?)8_Zj;NnWMl0r2YI zY(hZpuciHNE^n&;EoVN-^Q$U-Hi6G%C+9bu{`biHUd%UH#C+`2lBaL>(6xd2#JvIW zd$#y%S&x9HaYVms9PumG_dPH5RX%P=a{7-+{yMhv6XrX)9Z>fFzY|mY6lJ|p8Xk>4%1(i`n|lqa5md1W8OMZ=9h>0R>l8S-WO)Y^`YOg9nRnN%-;w; zQ#$Ww{p7;|q3LJ+`u&;p-vlrEDK3YqJ(>MG$@{(|?VQGT-ow1-w~|-w@XO#c#m9P0 z+8N}2L8bE<%zI9d>F_YyuVy~!YWH4YzMsomA?yDWy!tm&0*?BeUh(b{&&m=6O!M~{PTa1`u%T8UgcN80m(-OB(KWn`ON#ceVEC12AJ>W zdZ+Y{{iC##pE9Tat>CE~_I^|9FJb+XKTF;k zl)Or}-!dO_J@5LF%=UMKC;Nk3-p*nBr~O6piMyozWjQRUZHpJrZ?(eVnvIR&c+DpK zt9I+w@TR8LjWtc-Ej4Z7#+Jso6~1!K!q%2(O{~%1ToZ3twHRkCwbisWM%ybaoCOt* zZmo^R;xMqWDe6mOO4AsB8nZBsDfVTK%>0tH=oA-=JBM3pnlr_yxFp?pslj+z2ICuR zYOjvA)PswKtgT6n4}*Hp(BZi&{$ zTidFxL~b>*-aMDA4= zbw_flmSQn!DSHZvgmzakydITCr5Co6^ukuMm~*!zyiOF6@J193Q#rAV19CY-iP1Br z!j;qGsXMuy{qc(zhgXN&aDH4=G(@t*qELvUyd@N_UKp;9m-y;i+peu?tF!Z6Hf2;5 zipPJ_j7eOI*5}GK2VY)L-H?o>hpPDbDZga7(HrE2O*JtzQmYnhieA%JU$J(LpU%T8 z4X-O+aLwlE=4eGzQ?R~?5^G_&3GuP3(HcZ;L3B&BC0;?*sSI_6o;#!M*0xR8v}LqC z3FGc^Te`q_zbwyE_f|g3gpH>?%fboQE>W&GMdRU6%xPjpzE;P}!sXHScw6hXs$fbL zqvBo?POJH)=^cnaLyw!@J}gM@Ko+HIElO`i%F@l3rAMOJm#$WvshccLze90Jx`$9) zn!%nwyuTaP$T{IdBIUFDrDgUbGzV95 zkgOREHxhpR*04x0mBYAAS$aNY$R$0CGhfqW&gqz9ll`dVXFI$0fBzzD&86qJ{U32S}D2m>yjlPOhWy<^t z6gEUf7CzfP#th@zM?7@BlrBv6VDwzY{*-7GPaz)Hw$;R3U1TVXEG!Y-WQLoROzWj7LiC~n2sKE>2jMlX_hihtUqwVb}4rfJ~3R9IjXemX$29cUUJboGHtl=KM@$Ny{ zT+`m(xTz(K0eOn&PHtvJ#TML!ZYGQ4IMt=tXRP$>gM@1~Z^iSY^uiae%jlO>kYnOc zcTD^lI#~aL^p44&p@a1=Oz-^s89E{VqV&$szbHe8lfkoDP@LW&F38YHF2JbC-AOLU z;EXKD;E65B;GitIT4V8cI+;7x*c3I&n;KiLh7fU&i9eV)YQp@A9d7? z!X#OHd~<#M{93CniuA*D37WxpbGWt%b+8?GtZP+MW_8FF+-&n!ixRCV8jsc?Wio1n z>){DdylS?EDVE!;`Zgp*xNdWE^EL$0en~GxR}F_(tX;agB3!YiJdC;?zI4s|!xfcG zR+g`|!sS=2S-QGvxn0yQU$LSZ{Xl5xvQ-tLb;vNtD=8L(t}x11O}w@tR#PWB#;Uqd zOK=S)TgvEs_O_Oqrp3_~n(GTU)>o}>Z(GpV9k`|D{IxnP`X!BEW8>t@93o)Teh@b9Syft zZCM*!vBkM{{pn{kVhbQ5J`BuZ&5lCxx9z>xLMi?$cSx}nlTPrHo zhQcCq>?xdL|EAbxG@VUy#%>EHl-9sSTgVTRG3%*_w>Kbybz$URm6MtX6|ZS*D)CX{ z8*Phk`v5+^P>yzU8BHvSNuUKt2sOV8J9fztqlbevoL(mvs1$bHX(9#DuIZGX8K+ZV zMdkBk3_8!RLf7M7^>)IDbV>EMQH;pVR(V~D)q>2W!if==JW6m?j;N`jr&48?yje6Q z_Koonb{%t@ZxyrsDM8~h8fdyzaY?)TV*Tz5<=soeRV!(>a2eehaicmc%3m?96jVWJ zdALeN6+bJW8%DLjNCL|4dI$v+QV3Fge&voM6@oO8&qEHFTA&&U{5yuKhR=+kAdB+1 zN}rQ039o9!5Ixq~)L6T%qNJjt#E){&hM1x-hns5J%ORk45~b5nd+aB+69kVi zrKfk-7Stiw6rF^sF;uOzDH^IXDHTjr{zI526}3%Nn3hp8h}${-Pb2)FM)WCqIU@I$ z93OHtT=_Si0_xU5#qSIOj>FO(Fm z`KsWRHjLgEJ27Z!N7)HuFpu^EYuz<1wb6}1MA9xFjt7Izc5QVyaOtY5Wy{yE50}pO z%RbjueH%5zS}T##mN*@ntDEJ14Vkk#T-z3{iAPbxTiY=35tEi#s*ICM>}R2P+d1}D zcQ~ol$Cwjh97pv}dV(T`sz>MESM;v#8n`sQ(t@g>L~3hojhkm9VaVP6shhisxQd!9 zI_!oz`Q?(+;6fYxOz38I<2I!R4SxA*8qvd%l+SH$MI9&O5G6R>x;b?d!9k|4$nLa9 zLbmx^HaADx(7`n|Vji`t7Ivb2I_cYbogyx4y&oEp(R#2eaLuN5AY7<19+ zDU&007%K%Prf>EgIk4))E%eV`8AMZ8BMi$^9nG?PIjl9BtmX zx!xHFrt}^M*`QN5R#R6en{%3dh_*F1wxH|ZG6^R~^W=TDe&YtC!gLLvu^iGPpeW_BPtH)=qQHVzyS4 z#Z0re7<=Nslvbg*lqT|N;z&%6VyJf+CVML|>zS&Ds*$0kTvWTct&KJtQ5WS8Uv7;77Q^pW5Z9W3U%ss-QlIQe&n0Q#Rxdo$3OzF}r0-DKX04)5rz|U=~GL3&Q8A$$z z)Iy1jI}>WY=8XC!La8hMT7h8V%!2_y+J<8#gt?n_JtZ z<60e$UW54m`vzoOsNOzw|qj8xq+Y`PvlfAx*En89!gh`W4^Z;E2K>=S!-vHBcq(t9!K=BBLL za-B|_)1T0af{Ylg1y@~;IiI@57Md*=HF7GY)t>q}I6qt)*iOMQR>f2_jSEpR)ue@J zw8LU72N$4D?#xSbckPDL)!J#UMCu^QJzNrNH)6Fq>n1jr5kzjhC0AMkpaH$pQd03i z7G4fN>aDD~*svts2U8uCH)jv+5ROvRD#F6I+aiLo=Q-q)JeoToCvz+Gg`ujqG~Z>% ztJY=2$foKt(Q?D**pjKILo-did-hoUd;H%z+I&PFRkkh~wJ^0_g2>uXkGF{Kn4{ct z+K|bvsKE~A=2cc;PzSeJ51EYM%}F~mu=+xn=;?oi(9wFEf$=qNVF z%_T&xU}Kmz<&h*!FUjSVssNo72Ha!;L7GMvu8qaolq$|#1k)0oknvx;l)iAiA=L)YO#{aT54;g7V2!49RRe+ zVz;lTE?cW;nHPo0*#xbT8>m0};gds-Q7!hDFD zIG#kUr>vxS>Kjhbvt*j!8Nc4zqngP!xnj=*JS%-ZbUxm=#V!TB3nJNY$^4S};&n1x(U)rI!oIIb>=&Rdh-*$29u)+yr?`C1|ul?Czz zLVfRA#%-bWrXsqwGWVN`3-mV?DNT4>?M%v1Bdd1R8MZWFe;uj^O@dG=@fJkP=#TG# zi|LDTJZ)98o%zbNp&61XeMc3Z#?D~g07eg}Mrf$wE0kMR&l8zdrU#^U9y_QoDQo*; znKm_04Z{Yu7JXYL8dpPNX5Am_r0L{}o?Z6e@(!N1h4JPFF|S6BtmtSmb%+qyeylT< z#?C$M`mK#^=7n&*HKZ*w#ieSINQ~Ymv{F=zZhTF2Gj>rmg%Dt2xOie~N4(2M%;)JZ zwjZW7Bv*^$>a6O}Mqw@yy&}J7PMuZCt6Z5oa@+q)aY<~C73FKPEjenNHl={OzBO%Q z$|SK|lck!TseFNVYNp+r)4Ifj+95R)p!ptY@5G@wb0!KWFq3T8#DnXF+?%IY%B{Uf z^lY1?<<2xXEOOf?w5K}DX+OrnO+g{;c)SH)o7YTG^4_{Bj>T`YDUOt`(Qf0^Mtpk1 z%%c5y-2W;!Ryn9NnxfTwvGW{TrlBt9sC2Wzw3##Y8kyKBKZ#0eXZl;+i|7N4^!MuO zT!p`m-=B4-4mnZvjZHYC#6E$AUM;R4$FoRZ2|aAjja`#+b3M=BH(-=X@3CT*8e6a1 z;kjg~ujowI%G^f{a;JsiEo=4iVx|h`%&{qB7%)(iNMmYk$v1OYm5Sr@cM^mP{ z=|+>7ggD3^|1eKmV4K8hDirRG=d^WSt@#*Bi7Jy^Vh`sg+6m3=+wDIO;}*&n037!t z{Zs?dzO~5LhN-2L@6yX=ucazj*9_MXokA(SFNfinIFW&-GQ+$;gCX{&cmozy+gjUD zz4bx3Y$h|+?t?2b>TxFHmZr>U2xEuDvzffQ|c2QFCF)-c87HwR#~lOYfD@=Tq5kcfbed zZOX`G-i(9ZQ$Az5zwz&K@zb}PIQ=?z#k5X|r$P)f?0A3>gKxRl)L2oRB>p$Fitc{4 z3RQH?MEhCfn5p|)m)g&|M$A;Br?9tHI3~2*Kfx2F@FE;qUUdde2}aYhah4~m)nS)V zea+@3EQN&WNDjO|g}Ecue%mr?!kreJO6Wn%fYOAoc(KCXk%|`4*jpZApGx0dLrwpr zM$IDT?tR1V?$Z4>9&M^>+s5-pnX8DQ)lrY`a(u;vV#@h@u$nrpkXkS4Afs{W-(l~l zN@&nXb7yPl4Mjf0r**vbgTo%p2{}wi{ZSWA%iv1N zMtcke?}>&zbmR+7GvOEh34aQ;dpW(htDhKUyBmjaR!wcsSoAiVyFlad7HpKtTv4(3 zg1Us43Rhw?o9w6$SHqQkEe4}KM2Ak-nCx^ZT8>HXHK6uHq7G#E3_?Vjs6_GiG`dPdS+G9K)$FGzo`2 zEi)`pVAz;(i9+sG&sr2Rx#*K9#9FRV($e;L;clu%1UWN_9TloR9g0}j($CbY@R{Fce&jvGY&kf)$d8H8zda?oV#Owv!m zj4AvS8lq5HBa~iBMcD2%UGA!r#=PV}253_)az-q7I$n@kXLKLq$U93E+yUpJHhbFO`s!o@Qn=7X@*`z+G;NZ))9yci+uZu4s-Y)u_?9Sk@o% zk+pp#cCuxf^Q#IP-^Hs~EB5Wsp?M-nTASJ{>QfU6uOx1y18Qn;vIi|t+7pMR;=~TS z>E@a25*)7DLPyBps69H;6Q96nksC3v5RYmmC%VyWmDk7xd#ZztEmh584-s0V@rKb( zMLk5-fxLeErS-ZdI_ikN&Ok?Y;MK$0aI8_xiMk)u^}afFaL)Uehk2ABdRNu{ak9Jx z!_u<1->`Jz4vpPkOkn_r;+oxoL8VWf9ipcApk^LzbJWMGK04#hDI8?SP1C_n7`nC} z8GP(RUIE4MIQ!cKa_o(q%QOuk9aXA*t|Ot;oFBC+bm)?K0v(+^8q|d}>9ACc)-vq4 zEGZUyEhl@9Cv9$Po5De!;v`S=NarAj3e=NmXTV?v4O7zEn_S`mRXPk(_UNw30zNF> zP(bXOW>1kJ1Ey5#s21rCjbTMD&7oP5Htreckd2hr11EB{1~u@8$3NutL*FA~yr&wi z>762v+}>*EcAR0YauGxB1eC2B?4`}3+m&=OMxEzLH(2l5v{kP@ zdNA38C&dt+51y1gT)Gb?PjFT5;geHb=Hv+-z=!uZR07G(?U~bHYQ44VNih)TXPrbf zokZyrCvnQn(BfDUZNC`r2SW#m0|0p6n`{aUqapa>=nk=~9ix2z;d~tlr9c!+_r9Km z*wD!CX<|m)E7%>>zV9hsMDIz>VchcQ!B9U zJp#t6u9|=raW()5VueGVL0rme-vUG$n^)D=_b=I`|CRd8X3M z>rSG}51J>Mx+yLP{Qu*OqJ_Q--X`_)kwTuj5=Ff$}raakR zY|BAsEZ%SmRh#Cx$LWs$o#W;#Ck$fC1P$qPW=}Lpvy6taz!xhMng2$(Ll^*Z88i`cG!EElpO%*Ro3#c_>-7Htk3=> zpo5t~L^X0P*N%5G{rKy&@bEIeddM9^>Z^yL>T!>Fp8U~?b|#xUS!1t_Y2a%4ttO+# zXq!J=UW!(q4DH@QxaMMag0_i~9nZ<{&{rcDnk# z4^P-~ZOW24n2_-RRS01jZoB?PNY_4Ac>5UCP%2?yHcr`rw~Eq1lVfsgg8eZzdt*X+ zH|@0g>IGJIwCkFJTblLJyZWvyK2yU^miY8c)x-u8Q+Ws98Sl9YjrB#BjRqqAjLobHyHjP{5w;_xdbJ@hD*D~_RT%S&}KelJ58s@$yx$lKGc?4Iit!}Lj zSGd}tDfFjkiI$48(a?wwIpU0;U@cOhH0_k1slL#nN`WiqWo34)Si5=~RHZ3ZWNJIb z>p$j`HZrY5+it@YQZa3#mKs6Q+O>FA)wqbZpAP^LgDiDKKuS5(wA02|+~{SSjn!Yp zp6pC;YJt!q;yN%~ZEDEU0vyyJMb>M)w6PK|Z%`GZnaPRIwyAp`lI^eh?A(NwWirhj zrXMP2s2j8c8RrL4gMRov6tpCmZu*&o9kbnZXa>cqSGHN|*^v#vt-w((WklyCZC@DRp!pmywS zsKMv>$m>QXsjvIukSXy26S!hxCz2Zbyj^Mnw##COCqHLqpZF!(noO+~1kQFRbrm7**%kn|36}jHEG2HgB*P)a>uUDq-{75-p@GbjzwZ1VUTb{&)*-KJl$6`I6McM$3ACDnFRS+a{K_Z)>To6vsQXVOkmUt?ut}%Q(rJ zrq5j3C36Z%J-$zi?9AMAmW%zEOR0??XRhv$*Vu7(vl(iz(^@%Q4GrJ$cqz-27_g#k zN%?RojbT%q4Eq)n-cz7_)hXS3{PRr0q1AXRcTIVi+~~>A)~d9j^r+xn0$H5?L)~Y~ z^)CJvu3WLFqbg}!pZUynNbH{L@WnHw2)O-BM9FS7GL8MHd7_PBp_p1@8}~i`Ltb&F zx~?evIO7tc8r{waGdeaW(# zluuq|?-?+|L+|h48;D{Y)Y^vaceF1(?bvepXm&f0Is#JFN%v7E(bQmg?*4Wa9!>SL zX%o9N7ecQIYU6l&qCreQOeDmVW?_tXqz;>Bx19+entanjWO}hGYvzbm5~@ZS->=DZLSK}}>vmg;>HWi4D}D9=-;v)Ii=s%^w^nG2GIm$O zTVTR9o43k?i4c-^zhf+mh6zJaH3~(kNLBY$;ETVGI(a%%+D0SpqRtJ=RIZT=>K)H4 z$L=1KrwUiVHPK$KwsL2xDfC8V*46~GZ}LENRMkl|<7!qJGj~q+teR$S_72JW$f6U- z{>oL#F=^T6idCy38YPuc4bd8`2GB6GEs7U|MbiQW>_re`ujbY*IFU%~KoUpr$Q?am z0Sa>;a!b#II$YY~Z#P3K%_=I46_HLfbt)qJ%PB49!IdL9IVE<&V*Lk;o~i3U@(@Qn zIjy;Ai*9sYqgZ)ESBrI6`O0l$%ciC%y_7|1S?9cWTf+7CAR8+&okLC|w-vMsru7TB z3Az+4*7Byv_3Oi>^TjK_A|I=3=&0T6#Me%%RyCKFRH&29s>5}y_z5)M|9E zzD%8~YdZ^{gghZw3oVkHkkJlIxMRD+Z@5gBA`%T!D3v}F_?PdSmPRjr#B#}%nVv81QA%(U%? z8b|gWQs+C=a9B)t$;GgCe47j3JfF(`Ev;2Rgbe%ZGsa^gO~ZK8B$G(-!!|UbF@yQb z&-SRq7j{5;n6nYyr!-EDKa|lPL0iJ5=FhKN@YX62o;=2 zlVZ;a;RH#Uov!{#Y|W;=vJJ(N=AJN2#ET)y@yY@%lAsaPKF#J-`tcWz3^O0P&q<%A%&=B8zSEh$ zs;`!i$fb0xM(g`d#hFRkQ@Ce<(ZuUBz{081$0FlDc6N_j&U?$#SZG}AOvR4(&~^qP zm&NjfA8+zv>JbeePL>JB>HK44Exk3EH45W&SO?W{WAI3;P5Mb5RhWWXgrkzmC;DXs zln6AW&iuIe#BsP0{eKO7QgJLKPT6x7SGn0tnebEP3>)X>1FxyU`5Hoa1|En8fzQlP1_KF4hpo@X=s6sAMP!r zYjJ<1OPQF;L3!KMvYD1TiYMBE+YKt3vB|EyATq9vW9bVNPVa7)I`{S(I6{}B@1xR$ zX}b(B&)IRTNPT-QTg$6@nnTnT(qKyacB#o>$QGmY>TuZl0j@Q59N*V zq_l*RaZ=DuT|4Ol|eC7(Zq;aux#a}+(zD^omDQ1ooa>=MHr$g zBTti?F~b`fhlNo_onYZ6nfhP()P9M9Gyp0(;XyT+9)=DYFi zqn#yva{Pk-N8X4B2Lk2fVPUvHK9@<>yUhlEa3V=^e$c8k5p)4d3MpzNkZM?p@@+B8 zW~3}&V^M!TF1tI>;?3cu2TGyd2`&EI)zht6Og7CZZ-!CG{RpcZ^v3i z(oSWENH93BG7pEAt2@PbYcsS_u$ijPWcFmlZYc;!5d#8m_ zb&QjFnNfGP_hAgd-#%#@`Uq=Kwwy>I%5I>Cv=o5tpVN~MDhNNvxJa?@W~*^@48ALk zMnZLi08?7nMq^v!5vJ(e0MpT;@p*?n()(~KpI>Vo*5Nd3*H*TdHuHwt)L{p@`C74q zF;M~H{Cs}>Wff)=!Zi^|IMaC#bHuti=(?|kDiRCWYTYcDSf+4)&xy5vr#isQax9U3 zL{mVUhpWG3~xg8@4qFg3LMB5PG>C&u}C54w2)Ty(z+hxEQ>7x`{4taDjG6;8eqPuCW z<4G)vgW<~w?@o7)j?!kW)o#?`9Y(Um{Yn23%}86H!b3wl$pmB4r1I`ttAS zeI#s?KF*7&dE|9I?xuH9#VmBvtfS)OI!>^%1ijr?mSiIRqo#R!kh)7XncxNDlo(?J+ zTQD1hNZOdASPt}{eqZ)_!>s!T7c-;QqXRi|lif67loPSDi=bm7xbGa+kj&G@*cHfD zu&>KvA$TI4P+~pBxWgt?g?OqgHi|oEEHDlANMMWlNcrY;LTfMfDnO&ac$~w6Vs%h@ z@Sa&z`#HthUqCn1&CR!57NTOF$~VBA_9Ci6dW0!VZQI$)76m2gaczI((<9LE9W*DI z@{D}hzBwSjUcjg(53E6{iG3&zAV*3@AeM=7nltyH0kY60W;ITvAhY&WisLZ>%|HMM zqa@9(xW>Oe6kxb)Ai2^y$>c17W{4B80=)|EC#=4qfu17s$i({TfcUWpRCG9kp zDAHv#^zb{%ofPGPO+FR?)B4w=bA3kZ)-Dj-8V{rzIEXX_RG8Mi8o>0>=55hqU_KTw zf!x8f6BW-fD8s~Cjku$g#SA<5c+}slKiGpxyisc(9zb8daRW)IljeCXT@nF&>MC={ z?S?6#gm9n9bcQnF<^1*g(5)M-bPv7|H4Pc_=hXbd${H_O`(r@OrBl~I#jA$lnN)hp zdEY{C8?<~x7aX?gT()g###C=J!7jElETlUaGqmTx8cK(&hvrDt0jg<&iwYTeHD&Vc z;3`j^iRNv>4wpv}c9>teQo;^Zn;CRM(CBft*wV&Ds=4xeB-VZW!iP(v+%($nzOn$xHQ@F)U9Yn5$<%Rb5uu4}O9*Q5Ij9VFigrugEa#wzAYe!keC(-@+*iIGpYTja>z+)L{$nXLQ_*0hQZ zFNxGyuG6@V;kwda5StqwFLj+E!PV5!Td%%TZ`@of=tf0uCE2nVBoVZ3f zgwUtMAtdt|o(+}mU=U;9QN)-=2v$V?O(;RvQsqNiD|%9YGb5bB)X0b>O%By zAy_G8j@K9`V;j$#^BLHoOiXw)*2EN`*3=X!8kne9L6FmgBcUaZ1LuyUL8HzwlcrWEuFIdQ*-@`@=*TNK)H7B)-&D3_ zY}E92q9u-~hdu9JC=@I~9tf*Ck_!=h#!Ff2MXQG;8+l=zd7?^xE)dH(PQ@)ePfwS| z>zEKZeKL%L8@b~u609;0?zP-6_afhHk;w(;Jsvgyk+UlQxd-j^Rpkew5>uo+sdqTU zA-~IXx-k#>@Wxhg#vv7$LCl1W5U^>4-C2iMC!MTA0rE2sms5x-VE5tT70YmHtNTC8 zJ|*4W76Xrvo}~>g0B9(nXQrco2$CRsL;%@=bmfRs%|>*=({E$48&ME%WizjERv$KA zC&+1r1T)nTGHP$cHI;+A3yOBJ3%NAYeIm1F^n(>qFh~f_S1@o!6^sJOt=|~XK_b-( zOg?8UI^_C*ez|l33Bk@;E7tSOT84*Ax``M=l zP-xHrUgjLn(^GDUkOre9Ntsnhl>?-;4{nz?foNNmW(677$4hy3NN=$^QjM;u-MoL3 zd4C3SNdr~M5oFMDX<%%T(vy_h#|NFtCbpdxZ+CUi6trz7CDhlgm5>BpLhQ1M8jwRU z+h3v3ZgXz8(%wx<)#Ezu$Jdi|N9fYm2K(yny1B`u%3-Xl@38JDt4)bf~rq{k?Kf zOLwbiNPnRRTU!!ziI{XG&2xikrPZn&bA!51jZzQMzImF;!QK|&8|WAQl}u+%$4W?A zN=jRegM(TfC#y9i+i$j$v-o@$H~9j-FUZZtUo4xw@fZ2P_vE;4=vw#lef;<`z6yU$ z;I9e%gnKdeF9yHo@Y{vIC-C$({BLlK#zZbll5!cp@F|=B z3E`*vk=J4m{Th2N{Di=^ z_CCXpz0dIBcqyE2zm3c2N^O5a;3owBSK6>_6;2QT0KoK!3jbK(9}E1XAeIh;!fU^- z1nqzP!MwTk#~)Cyl0>gFKk!fCbvqrvKNa|=0(OsuClCJeC)}%h J0YLCG`5e0L$twT= literal 0 HcmV?d00001 diff --git a/qr_diag_repro.cpp b/qr_diag_repro.cpp new file mode 100644 index 000000000..c042d059b --- /dev/null +++ b/qr_diag_repro.cpp @@ -0,0 +1,98 @@ +// Standalone reproducer mirroring the geqrf+orgqr diagonal regression test +// added for https://github.com/uxlfoundation/oneMath/issues/626. +// Runs geqrf + orgqr repeatedly on a diagonal matrix (whose Q is identity) +// and verifies Q's diagonal is 1 on every run. + +#include +#include +#include +#include +#include +#include + +#if __has_include() +#include +#else +#include +#endif + +#include "oneapi/math.hpp" + +template +bool run_case(sycl::queue& queue, std::int64_t n, std::int64_t lda, std::int64_t num_runs) { + const std::int64_t m = n; + const std::int64_t k = n; + + const fp diag_value = fp(2.0); + std::vector A_initial(static_cast(lda) * n, fp(0.0)); + for (std::int64_t j = 0; j < n; ++j) + A_initial[j * lda + j] = diag_value; + + const fp tol = fp(10.0) * n * std::numeric_limits::epsilon(); + + auto* A_dev = sycl::malloc_device(A_initial.size(), queue); + auto* tau_dev = sycl::malloc_device(k, queue); + + const std::int64_t geqrf_ss = oneapi::math::lapack::geqrf_scratchpad_size(queue, m, n, lda); + const std::int64_t orgqr_ss = + oneapi::math::lapack::orgqr_scratchpad_size(queue, m, n, k, lda); + const std::int64_t ss = std::max(geqrf_ss, orgqr_ss); + auto* scratch_dev = sycl::malloc_device(ss, queue); + + std::vector Q(A_initial.size()); + bool ok = true; + for (std::int64_t run = 0; run < num_runs && ok; ++run) { + queue.memcpy(A_dev, A_initial.data(), A_initial.size() * sizeof(fp)).wait(); + + oneapi::math::lapack::geqrf(queue, m, n, A_dev, lda, tau_dev, scratch_dev, geqrf_ss); + oneapi::math::lapack::orgqr(queue, m, n, k, A_dev, lda, tau_dev, scratch_dev, orgqr_ss); + queue.wait_and_throw(); + + queue.memcpy(Q.data(), A_dev, Q.size() * sizeof(fp)).wait(); + + int bad = 0; + for (std::int64_t j = 0; j < n && bad < 5; ++j) { + for (std::int64_t i = 0; i < m && bad < 5; ++i) { + const fp expected = (i == j) ? fp(1.0) : fp(0.0); + const fp got = std::abs(Q[j * lda + i]); + if (std::abs(got - expected) > tol) { + std::cout << " run " << run << ": Q(" << i << "," << j << ")=" << got + << " expected " << expected << "\n"; + ++bad; + ok = false; + } + } + } + } + + sycl::free(A_dev, queue); + sycl::free(tau_dev, queue); + sycl::free(scratch_dev, queue); + return ok; +} + +int main() { + sycl::queue queue{ sycl::gpu_selector_v }; + std::cout << "Device: " << queue.get_device().get_info() << "\n"; + + struct Case { + std::int64_t n, lda, runs; + }; + const std::vector cases = { { 256, 256, 6 }, { 512, 512, 6 }, { 1024, 1024, 6 } }; + + bool all_ok = true; + for (auto c : cases) { + std::cout << "[float ] n=" << c.n << " runs=" << c.runs << ": " << std::flush; + bool ok = run_case(queue, c.n, c.lda, c.runs); + std::cout << (ok ? "PASS" : "FAIL") << "\n"; + all_ok &= ok; + + std::cout << "[double] n=" << c.n << " runs=" << c.runs << ": " << std::flush; + ok = run_case(queue, c.n, c.lda, c.runs); + std::cout << (ok ? "PASS" : "FAIL") << "\n"; + all_ok &= ok; + } + + std::cout << (all_ok ? "ALL PASS" : "SOME FAILED") << "\n"; + return all_ok ? 0 : 1; +} diff --git a/src/lapack/backends/rocsolver/rocsolver_helper.hpp b/src/lapack/backends/rocsolver/rocsolver_helper.hpp index 31feffd7c..37e42046b 100644 --- a/src/lapack/backends/rocsolver/rocsolver_helper.hpp +++ b/src/lapack/backends/rocsolver/rocsolver_helper.hpp @@ -255,6 +255,16 @@ struct RocmEquivalentType> { /* devinfo */ +/* Allocates the device memory rocSOLVER writes its `devInfo` output to. A failed + * allocation must not be passed on to rocSOLVER as a null pointer, since that + * silently disables the routine's error reporting. */ +inline int* create_devinfo(sycl::queue& queue, const char* func_name, int dev_info_size = 1) { + int* devInfo = sycl::malloc_device(dev_info_size, queue); + if (!devInfo) + throw oneapi::math::device_bad_alloc("lapack", func_name, queue.get_device()); + return devInfo; +} + inline int get_rocsolver_devinfo(sycl::queue& queue, sycl::buffer& devInfo) { sycl::host_accessor dev_info_{ devInfo }; return dev_info_[0]; diff --git a/src/lapack/backends/rocsolver/rocsolver_lapack.cpp b/src/lapack/backends/rocsolver/rocsolver_lapack.cpp index 5b0c265b2..a362e6d38 100644 --- a/src/lapack/backends/rocsolver/rocsolver_lapack.cpp +++ b/src/lapack/backends/rocsolver/rocsolver_lapack.cpp @@ -1273,7 +1273,7 @@ inline sycl::event getrf(const char* func_name, Func func, sycl::queue& queue, s std::uint64_t ipiv_size = std::min(n, m); int* ipiv32 = (int*)malloc_device(sizeof(int) * ipiv_size, queue); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1411,7 +1411,7 @@ inline sycl::event gesvd(const char* func_name, Func func, sycl::queue& queue, using rocmDataType_A = typename RocmEquivalentType::Type; using rocmDataType_B = typename RocmEquivalentType::Type; overflow_check(m, n, lda, ldu, ldvt, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1462,7 +1462,7 @@ inline sycl::event heevd(const char* func_name, Func func, sycl::queue& queue, using rocmDataType_A = typename RocmEquivalentType::Type; using rocmDataType_B = typename RocmEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1508,7 +1508,7 @@ inline sycl::event hegvd(const char* func_name, Func func, sycl::queue& queue, s using rocmDataType_A = typename RocmEquivalentType::Type; using rocmDataType_B = typename RocmEquivalentType::Type; overflow_check(n, lda, ldb, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1818,7 +1818,7 @@ inline sycl::event potrf(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using rocmDataType = typename RocmEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1860,7 +1860,7 @@ inline sycl::event potri(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using rocmDataType = typename RocmEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1944,7 +1944,7 @@ inline sycl::event syevd(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using rocmDataType = typename RocmEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -1989,7 +1989,7 @@ inline sycl::event sygvd(const char* func_name, Func func, sycl::queue& queue, s const std::vector& dependencies) { using rocmDataType = typename RocmEquivalentType::Type; overflow_check(n, lda, ldb, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); auto done = queue.submit([&](sycl::handler& cgh) { int64_t num_events = dependencies.size(); for (int64_t i = 0; i < num_events; i++) { @@ -2074,7 +2074,7 @@ inline sycl::event sytrf(const char* func_name, Func func, sycl::queue& queue, const std::vector& dependencies) { using rocmDataType = typename RocmEquivalentType::Type; overflow_check(n, lda, scratchpad_size); - int* devInfo = (int*)malloc_device(sizeof(int), queue); + int* devInfo = create_devinfo(queue, __func__); // rocsolver legacy api does not accept 64-bit ints. // To get around the limitation.