From 7055a717617cbe93a019361d1476b3409201aaf8 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sat, 15 Aug 2026 09:37:44 -0700 Subject: [PATCH] [blas][rocblas] Restore the pointer mode when a rocBLAS call throws The reduction routines switch the rocBLAS handle to device pointer mode and switch it back once the call returns. rocblas_native_func throws on any status other than success, which skips the reset. Handles are cached per thread and reused, so the handle stays in device pointer mode for every later call on that thread. Routines that pass a host address for their scalar arguments then have that address dereferenced on the device. rot is the most exposed: it forwards &c and &s and relies on the handle already being in host pointer mode. On a discrete GPU this is an illegal address, which matches the intermittent failure of RotTests.RealSinglePrecision reported in #486. Set the mode through an RAII guard so it is restored while unwinding. The guard restores the previous mode rather than forcing host mode, so it nests safely. Verified on an AMD Instinct MI210: after provoking a failure inside a guarded region, a following rot aborts with hipErrorIllegalAddress in 10 of 10 runs without the guard and in 0 of 10 with it. Co-authored-by: Cursor --- src/blas/backends/rocblas/rocblas_helper.hpp | 24 +++++++ src/blas/backends/rocblas/rocblas_level1.cpp | 66 ++++---------------- 2 files changed, 37 insertions(+), 53 deletions(-) diff --git a/src/blas/backends/rocblas/rocblas_helper.hpp b/src/blas/backends/rocblas/rocblas_helper.hpp index 77101f1bc..721d8bc95 100644 --- a/src/blas/backends/rocblas/rocblas_helper.hpp +++ b/src/blas/backends/rocblas/rocblas_helper.hpp @@ -173,6 +173,30 @@ class hip_error : virtual public std::runtime_error { hipError_t hip_err; \ HIP_ERROR_FUNC(hipStreamSynchronize, hip_err, currentStreamId); +// Sets the pointer mode of a rocBLAS handle and restores the previous mode on scope +// exit. rocBLAS handles are cached and reused across calls, so an unwound scope must +// not leave the handle in device mode: later routines pass host addresses for their +// scalar arguments and rocBLAS would dereference them on the device. +class rocblas_pointer_mode_guard { + rocblas_handle handle_; + rocblas_pointer_mode previous_; + +public: + rocblas_pointer_mode_guard(rocblas_handle handle, rocblas_pointer_mode mode) + : handle_(handle), + previous_(rocblas_pointer_mode_host) { + rocblas_get_pointer_mode(handle_, &previous_); + rocblas_set_pointer_mode(handle_, mode); + } + + ~rocblas_pointer_mode_guard() { + rocblas_set_pointer_mode(handle_, previous_); + } + + rocblas_pointer_mode_guard(const rocblas_pointer_mode_guard&) = delete; + rocblas_pointer_mode_guard& operator=(const rocblas_pointer_mode_guard&) = delete; +}; + template inline void rocblas_native_func(Func func, rocblas_status err, rocblas_handle handle, Types... args) { diff --git a/src/blas/backends/rocblas/rocblas_level1.cpp b/src/blas/backends/rocblas/rocblas_level1.cpp index ab251cc5c..fc2668330 100644 --- a/src/blas/backends/rocblas/rocblas_level1.cpp +++ b/src/blas/backends/rocblas/rocblas_level1.cpp @@ -50,16 +50,12 @@ inline void asum(Func func, sycl::queue& queue, int64_t n, sycl::buffer& // rocblas_set_pointer_mode mode otherwise it causes the segmentation // fault. When it is set to device it is users responsibility to // synchronise as the function is completely asynchronous. - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = sc.get_mem(x_acc); auto res_ = sc.get_mem(res_acc); rocblas_status err; // ASUM does not support negative index rocblas_native_func(func, err, handle, n, x_, std::abs(incx), res_); - // Higher level BLAS functions expect rocblas_pointer_mode_host - // to be set, therfore we need to reset this to the default value - // in order to avoid invalid memory accesses - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); } @@ -183,17 +179,13 @@ inline void rotg(Func func, sycl::queue& queue, sycl::buffer& a, sycl::bu // rocblas_set_pointer_mode mode otherwise it causes the segmentation // fault. When it is set to device it is users responsibility to // synchronise as the function is completely asynchronous. - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto a_ = sc.get_mem(a_acc); auto b_ = sc.get_mem(b_acc); auto c_ = sc.get_mem(c_acc); auto s_ = sc.get_mem(s_acc); rocblas_status err; rocblas_native_func(func, err, handle, a_, b_, c_, s_); - // Higher level BLAS functions expect rocblas_pointer_mode_host - // to be set, therfore we need to reset this to the default value - // in order to avoid invalid memory accesses - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); } @@ -229,16 +221,12 @@ inline void rotm(Func func, sycl::queue& queue, int64_t n, sycl::buffer& x // rocblas_set_pointer_mode mode otherwise it causes the segmentation // fault. When it is set to device it is users responsibility to // synchronise as the function is completely asynchronous. - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = sc.get_mem(x_acc); auto y_ = sc.get_mem(y_acc); auto param_ = sc.get_mem(param_acc); rocblas_status err; rocblas_native_func(func, err, handle, n, x_, incx, y_, incy, param_); - // Higher level BLAS functions expect rocblas_pointer_mode_host - // to be set, therfore we need to reset this to the default value - // in order to avoid invalid memory accesses - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); } @@ -305,16 +293,12 @@ inline void dot(Func func, sycl::queue& queue, int64_t n, sycl::buffer& x, // rocblas_set_pointer_mode mode otherwise it causes the segmentation // fault. When it is set to device it is users responsibility to // synchronise as the function is completely asynchronous. - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = sc.get_mem(x_acc); auto y_ = sc.get_mem(y_acc); auto res_ = sc.get_mem(res_acc); rocblas_status err; rocblas_native_func(func, err, handle, n, x_, incx, y_, incy, res_); - // Higher level BLAS functions expect rocblas_pointer_mode_host - // to be set, therfore we need to reset this to the default value - // in order to avoid invalid memory accesses - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); } @@ -397,16 +381,12 @@ void sdsdot(sycl::queue& queue, int64_t n, float sb, sycl::buffer& x, // rocblas_set_pointer_mode mode otherwise it causes the segmentation // fault. When it is set to device it is users responsibility to // synchronise as the function is completely asynchronous. - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = sc.get_mem(x_acc); auto y_ = sc.get_mem(y_acc); auto res_ = sc.get_mem(res_acc); rocblas_status err; rocblas_native_func(rocblas_sdot, err, handle, n, x_, incx, y_, incy, res_); - // Higher level BLAS functions expect rocblas_pointer_mode_host - // to be set, therfore we need to reset this to the default value - // in order to avoid invalid memory accesses - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); @@ -435,7 +415,7 @@ inline void rotmg(Func func, sycl::queue& queue, sycl::buffer& d1, sycl::b // rocblas_set_pointer_mode mode otherwise it causes the segmentation // fault. When it is set to device it is users responsibility to // synchronise as the function is completely asynchronous. - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto d1_ = sc.get_mem(d1_acc); auto d2_ = sc.get_mem(d2_acc); auto x1_ = sc.get_mem(x1_acc); @@ -443,10 +423,6 @@ inline void rotmg(Func func, sycl::queue& queue, sycl::buffer& d1, sycl::b auto param_ = sc.get_mem(param_acc); rocblas_status err; rocblas_native_func(func, err, handle, d1_, d2_, x1_, y1_, param_); - // Higher level BLAS functions expect rocblas_pointer_mode_host - // to be set, therfore we need to reset this to the default value - // in order to avoid invalid memory accesses - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); } @@ -488,17 +464,13 @@ inline void iamax(Func func, sycl::queue& queue, int64_t n, sycl::buffer& // rocblas_set_pointer_mode mode otherwise it causes the segmentation // fault. When it is set to device it is users responsibility to // synchronise as the function is completely asynchronous. - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = sc.get_mem(x_acc); auto int_res_ = sc.get_mem(int_res_acc); rocblas_status err; // For negative incx, iamax returns 0. This behaviour is similar to that of // reference netlib BLAS. rocblas_native_func(func, err, handle, n, x_, incx, int_res_); - // Higher level BLAS functions expect rocblas_pointer_mode_host - // to be set, therfore we need to reset this to the default value - // in order to avoid invalid memory accesses - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); @@ -585,17 +557,13 @@ inline void iamin(Func func, sycl::queue& queue, int64_t n, sycl::buffer& // rocblas_set_pointer_mode mode otherwise it causes the segmentation // fault. When it is set to device it is users responsibility to // synchronise as the function is completely asynchronous. - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = sc.get_mem(x_acc); auto int_res_ = sc.get_mem(int_res_acc); rocblas_status err; // For negative incx, iamin returns 0. This behaviour is similar to that of // implemented as a reference IAMIN. rocblas_native_func(func, err, handle, n, x_, incx, int_res_); - // Higher level BLAS functions expect rocblas_pointer_mode_host - // to be set, therfore we need to reset this to the default value - // in order to avoid invalid memory accesses - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); @@ -641,16 +609,12 @@ inline void nrm2(Func func, sycl::queue& queue, int64_t n, sycl::buffer& // rocblas_set_pointer_mode mode otherwise it causes the segmentation // fault. When it is set to device it is users responsibility to // synchronise as the function is completely asynchronous. - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = sc.get_mem(x_acc); auto res_ = sc.get_mem(res_acc); rocblas_status err; // NRM2 does not support negative index rocblas_native_func(func, err, handle, n, x_, std::abs(incx), res_); - // Higher level BLAS functions expect rocblas_pointer_mode_host - // to be set, therfore we need to reset this to the default value - // in order to avoid invalid memory accesses - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); } @@ -681,14 +645,13 @@ inline sycl::event asum(Func func, sycl::queue& queue, int64_t n, const T1* x, c cgh.depends_on(dependencies); onemath_rocblas_host_task(cgh, queue, [=](RocblasScopedContextHandler& sc) { auto handle = sc.get_handle(queue); - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = reinterpret_cast(x); auto res_ = reinterpret_cast(result); rocblas_status err; // ASUM does not support negative index rocblas_native_func(func, err, handle, n, x_, std::abs(incx), res_); - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); @@ -1065,14 +1028,13 @@ inline sycl::event iamax(Func func, sycl::queue& queue, int64_t n, const T* x, c cgh.depends_on(dependencies); onemath_rocblas_host_task(cgh, queue, [=](RocblasScopedContextHandler& sc) { auto handle = sc.get_handle(queue); - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = reinterpret_cast(x); auto int_res_p_ = reinterpret_cast(int_res_p); rocblas_status err; // For negative incx, iamax returns 0. This behaviour is similar to that of // reference iamax. rocblas_native_func(func, err, handle, n, x_, incx, int_res_p_); - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); @@ -1149,7 +1111,7 @@ inline sycl::event iamin(Func func, sycl::queue& queue, int64_t n, const T* x, c cgh.depends_on(dependencies); onemath_rocblas_host_task(cgh, queue, [=](RocblasScopedContextHandler& sc) { auto handle = sc.get_handle(queue); - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = reinterpret_cast(x); auto int_res_p_ = reinterpret_cast(int_res_p); @@ -1157,7 +1119,6 @@ inline sycl::event iamin(Func func, sycl::queue& queue, int64_t n, const T* x, c // For negative incx, iamin returns 0. This behaviour is similar to that of // implemented iamin. rocblas_native_func(func, err, handle, n, x_, incx, int_res_p_); - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); }); @@ -1192,14 +1153,13 @@ inline sycl::event nrm2(Func func, sycl::queue& queue, int64_t n, const T1* x, c cgh.depends_on(dependencies); onemath_rocblas_host_task(cgh, queue, [=](RocblasScopedContextHandler& sc) { auto handle = sc.get_handle(queue); - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device); + rocblas_pointer_mode_guard pointer_mode_guard(handle, rocblas_pointer_mode_device); auto x_ = reinterpret_cast(x); auto res_ = reinterpret_cast(result); rocblas_status err; // NRM2 does not support negative index rocblas_native_func(func, err, handle, n, x_, std::abs(incx), res_); - rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host); }); });