Skip to content

[blas] Implement omatcopy2 for cuBLAS and rocBLAS #766

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 3 commits into
base: develop
Choose a base branch
from
Open

[blas] Implement omatcopy2 for cuBLAS and rocBLAS #766

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
facac5c
[blas] Implement omatcopy2 for the cuBLAS and rocBLAS backends
zjin-lcf Aug 15, 2026
2d018e1
[blas] Tune the omatcopy2 tile geometry for NVIDIA
zjin-lcf Aug 15, 2026
3fee43b
[blas] Tune omatcopy2 geometry for AMD architectures
zjin-lcf Aug 15, 2026
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
112 changes: 50 additions & 62 deletions src/blas/backends/cublas/cublas_extensions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,13 +18,19 @@
**************************************************************************/
#include "cublas_helper.hpp"
#include "cublas_task.hpp"
#include "blas/backends/omatcopy2_kernels.hpp"
#include "oneapi/math/exceptions.hpp"
#include "oneapi/math/blas/detail/cublas/onemath_blas_cublas.hpp"

namespace oneapi {
namespace math {
namespace blas {
namespace cublas {

// This backend only ever runs on NVIDIA devices, so the omatcopy2 kernels are
// always built with the tile extents measured there.
constexpr auto omatcopy2_target = omatcopy2_kernels::target::nvidia;

namespace column_major {

// Buffer APIs
Expand Down Expand Up @@ -121,25 +127,24 @@ OMATCOPY_LAUNCHER(std::complex<double>, cublasZgeam)

#undef OMATCOPY_LAUNCHER

template <typename Func, typename T>
void omatcopy2(const char* func_name, Func func, sycl::queue& queue, transpose trans, int64_t m,
int64_t n, T alpha, sycl::buffer<T, 1>& a, int64_t lda, std::int64_t stridea,
sycl::buffer<T, 1>& b, int64_t ldb, std::int64_t strideb) {
throw unimplemented("blas", "omatcopy2", "");
}

#define OMATCOPY2_LAUNCHER(TYPE, CUBLAS_ROUTINE) \
void omatcopy2(sycl::queue& queue, transpose trans, int64_t m, int64_t n, TYPE alpha, \
sycl::buffer<TYPE, 1>& a, int64_t lda, int64_t stridea, \
sycl::buffer<TYPE, 1>& b, int64_t ldb, int64_t strideb) { \
omatcopy2(#CUBLAS_ROUTINE, CUBLAS_ROUTINE, queue, trans, m, n, alpha, a, stridea, lda, b, \
ldb, strideb); \
// Unit element strides make omatcopy2 equivalent to omatcopy, but routing them
// to geam is not worth it: cuBLAS has to be called from a host task and with a
// stream synchronize, whereas these kernels are plain asynchronous SYCL. On an
// A100 the kernels beat geam everywhere from 1024x1024 up, by 2-231% when calls
// are pipelined and by up to 29% even when the caller waits after every one.
#define OMATCOPY2_LAUNCHER(TYPE) \
void omatcopy2(sycl::queue& queue, transpose trans, int64_t m, int64_t n, TYPE alpha, \
sycl::buffer<TYPE, 1>& a, int64_t lda, int64_t stridea, \
sycl::buffer<TYPE, 1>& b, int64_t ldb, int64_t strideb) { \
omatcopy2_kernels::omatcopy2_buffer<omatcopy2_target>( \
queue, oneapi::math::layout::col_major, trans, m, n, alpha, a, lda, stridea, b, ldb, \
strideb); \
}

OMATCOPY2_LAUNCHER(float, "unimplemented")
OMATCOPY2_LAUNCHER(double, "unimplemented")
OMATCOPY2_LAUNCHER(std::complex<float>, "unimplemented")
OMATCOPY2_LAUNCHER(std::complex<double>, "unimplemented")
OMATCOPY2_LAUNCHER(float)
OMATCOPY2_LAUNCHER(double)
OMATCOPY2_LAUNCHER(std::complex<float>)
OMATCOPY2_LAUNCHER(std::complex<double>)
#undef OMATCOPY2_LAUNCHER

void imatcopy(sycl::queue& queue, transpose trans, int64_t m, int64_t n, float alpha,
Expand Down Expand Up @@ -302,25 +307,19 @@ OMATCOPY_LAUNCHER_USM(std::complex<double>, cublasZgeam)

#undef OMATCOPY_LAUNCHER_USM

template <typename Func, typename T>
sycl::event omatcopy2(const char* func_name, Func func, sycl::queue& queue, transpose trans,
int64_t m, int64_t n, T alpha, const T* a, int64_t lda, int64_t stridea, T* b,
int64_t ldb, int64_t strideb, const std::vector<sycl::event>& dependencies) {
throw unimplemented("blas", "omatcopy2", "");
}

#define OMATCOPY2_LAUNCHER_USM(TYPE, CUBLAS_ROUTINE) \
#define OMATCOPY2_LAUNCHER_USM(TYPE) \
sycl::event omatcopy2(sycl::queue& queue, transpose trans, int64_t m, int64_t n, TYPE alpha, \
const TYPE* a, int64_t lda, int64_t stridea, TYPE* b, int64_t ldb, \
int64_t strideb, const std::vector<sycl::event>& dependencies) { \
return omatcopy2(#CUBLAS_ROUTINE, CUBLAS_ROUTINE, queue, trans, m, n, alpha, a, stridea, \
lda, b, ldb, strideb, dependencies); \
return omatcopy2_kernels::omatcopy2_usm<omatcopy2_target>( \
queue, oneapi::math::layout::col_major, trans, m, n, alpha, a, lda, stridea, b, ldb, \
strideb, dependencies); \
}

OMATCOPY2_LAUNCHER_USM(float, "unimplemented")
OMATCOPY2_LAUNCHER_USM(double, "unimplemented")
OMATCOPY2_LAUNCHER_USM(std::complex<float>, "unimplemented")
OMATCOPY2_LAUNCHER_USM(std::complex<double>, "unimplemented")
OMATCOPY2_LAUNCHER_USM(float)
OMATCOPY2_LAUNCHER_USM(double)
OMATCOPY2_LAUNCHER_USM(std::complex<float>)
OMATCOPY2_LAUNCHER_USM(std::complex<double>)
#undef OMATCOPY2_LAUNCHER_USM

sycl::event imatcopy(sycl::queue& queue, transpose trans, int64_t m, int64_t n, float alpha,
Expand Down Expand Up @@ -484,25 +483,20 @@ OMATCOPY_LAUNCHER(std::complex<double>, cublasZgeam)

#undef OMATCOPY_LAUNCHER

template <typename Func, typename T>
void omatcopy2(const char* func_name, Func func, sycl::queue& queue, transpose trans, int64_t m,
int64_t n, T alpha, sycl::buffer<T, 1>& a, int64_t lda, std::int64_t stridea,
sycl::buffer<T, 1>& b, int64_t ldb, std::int64_t strideb) {
throw unimplemented("blas", "omatcopy2", "");
}

#define OMATCOPY2_LAUNCHER(TYPE, CUBLAS_ROUTINE) \
void omatcopy2(sycl::queue& queue, transpose trans, int64_t m, int64_t n, TYPE alpha, \
sycl::buffer<TYPE, 1>& a, int64_t lda, int64_t stridea, \
sycl::buffer<TYPE, 1>& b, int64_t ldb, int64_t strideb) { \
omatcopy2(#CUBLAS_ROUTINE, CUBLAS_ROUTINE, queue, trans, m, n, alpha, a, stridea, lda, b, \
ldb, strideb); \
// See the column-major overloads for why geam is not used at unit stride.
#define OMATCOPY2_LAUNCHER(TYPE) \
void omatcopy2(sycl::queue& queue, transpose trans, int64_t m, int64_t n, TYPE alpha, \
sycl::buffer<TYPE, 1>& a, int64_t lda, int64_t stridea, \
sycl::buffer<TYPE, 1>& b, int64_t ldb, int64_t strideb) { \
omatcopy2_kernels::omatcopy2_buffer<omatcopy2_target>( \
queue, oneapi::math::layout::row_major, trans, m, n, alpha, a, lda, stridea, b, ldb, \
strideb); \
}

OMATCOPY2_LAUNCHER(float, "unimplemented")
OMATCOPY2_LAUNCHER(double, "unimplemented")
OMATCOPY2_LAUNCHER(std::complex<float>, "unimplemented")
OMATCOPY2_LAUNCHER(std::complex<double>, "unimplemented")
OMATCOPY2_LAUNCHER(float)
OMATCOPY2_LAUNCHER(double)
OMATCOPY2_LAUNCHER(std::complex<float>)
OMATCOPY2_LAUNCHER(std::complex<double>)
#undef OMATCOPY2_LAUNCHER

void imatcopy(sycl::queue& queue, transpose trans, int64_t m, int64_t n, float alpha,
Expand Down Expand Up @@ -665,25 +659,19 @@ OMATCOPY_LAUNCHER_USM(std::complex<double>, cublasZgeam)

#undef OMATCOPY_LAUNCHER_USM

template <typename Func, typename T>
sycl::event omatcopy2(const char* func_name, Func func, sycl::queue& queue, transpose trans,
int64_t m, int64_t n, T alpha, const T* a, int64_t lda, int64_t stridea, T* b,
int64_t ldb, int64_t strideb, const std::vector<sycl::event>& dependencies) {
throw unimplemented("blas", "omatcopy2", "");
}

#define OMATCOPY2_LAUNCHER_USM(TYPE, CUBLAS_ROUTINE) \
#define OMATCOPY2_LAUNCHER_USM(TYPE) \
sycl::event omatcopy2(sycl::queue& queue, transpose trans, int64_t m, int64_t n, TYPE alpha, \
const TYPE* a, int64_t lda, int64_t stridea, TYPE* b, int64_t ldb, \
int64_t strideb, const std::vector<sycl::event>& dependencies) { \
return omatcopy2(#CUBLAS_ROUTINE, CUBLAS_ROUTINE, queue, trans, m, n, alpha, a, stridea, \
lda, b, ldb, strideb, dependencies); \
return omatcopy2_kernels::omatcopy2_usm<omatcopy2_target>( \
queue, oneapi::math::layout::row_major, trans, m, n, alpha, a, lda, stridea, b, ldb, \
strideb, dependencies); \
}

OMATCOPY2_LAUNCHER_USM(float, "unimplemented")
OMATCOPY2_LAUNCHER_USM(double, "unimplemented")
OMATCOPY2_LAUNCHER_USM(std::complex<float>, "unimplemented")
OMATCOPY2_LAUNCHER_USM(std::complex<double>, "unimplemented")
OMATCOPY2_LAUNCHER_USM(float)
OMATCOPY2_LAUNCHER_USM(double)
OMATCOPY2_LAUNCHER_USM(std::complex<float>)
OMATCOPY2_LAUNCHER_USM(std::complex<double>)
#undef OMATCOPY2_LAUNCHER_USM

sycl::event imatcopy(sycl::queue& queue, transpose trans, int64_t m, int64_t n, float alpha,
Expand Down
Loading
Loading