Skip to content

[blas][rocblas] Restore the pointer mode when a rocBLAS call throws #764

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
zjin-lcf wants to merge 1 commit into
base: develop
Choose a base branch
from
Open

[blas][rocblas] Restore the pointer mode when a rocBLAS call throws #764

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 24 additions & 0 deletions src/blas/backends/rocblas/rocblas_helper.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down
66 changes: 13 additions & 53 deletions src/blas/backends/rocblas/rocblas_level1.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
});
});
}
Expand Down Expand Up @@ -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);
});
});
}
Expand Down Expand Up @@ -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);
});
});
}
Expand Down Expand Up @@ -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);
});
});
}
Expand Down Expand Up @@ -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);
});
});

Expand Down Expand Up @@ -435,18 +415,14 @@ 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);
auto y1_ = sc.get_mem<rocDataType*>(y1_acc);
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);
});
});
}
Expand Down Expand Up @@ -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);
});
});

Expand Down Expand Up @@ -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);
});
});

Expand Down Expand Up @@ -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);
});
});
}
Expand Down Expand Up @@ -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);
});
});

Expand Down Expand Up @@ -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);
});
});

Expand Down Expand Up @@ -1149,15 +1111,14 @@ 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);
rocblas_status err;
// 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);
});
});

Expand Down Expand Up @@ -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);
});
});

Expand Down
Loading