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 <class Func, class... Types>
 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<T1, 1>&
             // 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<rocDataType1*>(x_acc);
             auto res_ = sc.get_mem<rocDataType2*>(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<T1, 1>& 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<rocDataType1*>(a_acc);
             auto b_ = sc.get_mem<rocDataType1*>(b_acc);
             auto c_ = sc.get_mem<rocDataType2*>(c_acc);
             auto s_ = sc.get_mem<rocDataType1*>(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<T, 1>& 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<rocDataType*>(x_acc);
             auto y_ = sc.get_mem<rocDataType*>(y_acc);
             auto param_ = sc.get_mem<rocDataType*>(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<T, 1>& 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<rocDataType*>(x_acc);
             auto y_ = sc.get_mem<rocDataType*>(y_acc);
             auto res_ = sc.get_mem<rocDataType*>(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<float, 1>& 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<float*>(x_acc);
             auto y_ = sc.get_mem<float*>(y_acc);
             auto res_ = sc.get_mem<float*>(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<T, 1>& 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<rocDataType*>(d1_acc);
             auto d2_ = sc.get_mem<rocDataType*>(d2_acc);
             auto x1_ = sc.get_mem<rocDataType*>(x1_acc);
@@ -443,10 +423,6 @@ inline void rotmg(Func func, sycl::queue& queue, sycl::buffer<T, 1>& d1, sycl::b
             auto param_ = sc.get_mem<rocDataType*>(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<T, 1>&
             // 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<rocDataType*>(x_acc);
             auto int_res_ = sc.get_mem<int*>(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<T, 1>&
             // 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<rocDataType*>(x_acc);
             auto int_res_ = sc.get_mem<int*>(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<T1, 1>&
             // 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<rocDataType1*>(x_acc);
             auto res_ = sc.get_mem<rocDataType2*>(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<const rocDataType1*>(x);
             auto res_ = reinterpret_cast<rocDataType2*>(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<const rocDataType*>(x);
             auto int_res_p_ = reinterpret_cast<int*>(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<const rocDataType*>(x);
             auto int_res_p_ = reinterpret_cast<int*>(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<const rocDataType1*>(x);
             auto res_ = reinterpret_cast<rocDataType2*>(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);
         });
     });