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 6a5427712..e84781bb5 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 = 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++) { @@ -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 = 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++) { @@ -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; } @@ -1322,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++) { @@ -1459,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++) { @@ -1510,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++) { @@ -1556,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++) { @@ -1603,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++) { @@ -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 = 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++) { @@ -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 = 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++) { @@ -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 = 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++) { @@ -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 = 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++) { @@ -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 = 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,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; } @@ -1878,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++) { @@ -1921,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++) { @@ -2006,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++) { @@ -2051,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++) { @@ -2096,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++) { @@ -2141,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. @@ -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 = 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++) { @@ -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 = 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++) { @@ -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 = 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++) { @@ -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 = 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++) { @@ -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 = 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++) { @@ -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(); 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/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..b78574ad0 --- /dev/null +++ b/tests/unit_tests/lapack/source/geqrf_orgqr_diagonal.cpp @@ -0,0 +1,171 @@ +/******************************************************************************* +* 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 "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(); + + 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); +#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 + } + 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()); + 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);