Skip to content

fix(rocBLAS) : rot and rot_ex fix ILP64 large N inc handling #13114

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

Merged
TorreZuk merged 2 commits into from
Oct 8, 2026
Merged

fix(rocBLAS) : rot and rot_ex fix ILP64 large N inc handling #13114

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
3a33ab4
rot and rot_ex fix ILP64 large N inc handling
TorreZuk Oct 5, 2026
d94b321
minimize code to match current optimizations
TorreZuk Oct 7, 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
2 changes: 2 additions & 0 deletions projects/rocblas/clients/gtest/rot_gtest.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -175,9 +175,11 @@ Tests:
category: stress
arguments:
- { N: 2147483649, incx: 1, incy: 1, batch_count: 1 }
- { N: 2147483649, incx: -1, incy: -1, batch_count: 1 }
- { N: 2, incx: 2147483649, incy: 1, batch_count: 1 }
- { N: 2, incx: 1, incy: 2147483649, batch_count: 1 }
- { N: 3, incx: 1, incy: 1, stride_x: 3, stride_y: 3, batch_count: *c_grid_yz_require_passes }
- { N: 1, incx: 1, incy: -2147483649, stride_x: 1, stride_y: 1, batch_count: *c_grid_yz_require_passes }
Comment thread
TorreZuk marked this conversation as resolved.
Resolved
- { N: 2, incx: -2147483649, incy: 1, batch_count: 1 }
- { N: 2, incx: 1, incy: -2147483649, batch_count: 1 }
api: [ C_64 ]
Expand Down
170 changes: 123 additions & 47 deletions projects/rocblas/library/src/blas1/rocblas_rot_kernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,33 +70,31 @@ rocblas_status rocblas_rot_check_numerics(const char* function_name,
// If there are any changes in template parameters in the files *rot*.cpp
// instantiations below will need to be manually updated to match the changes.

// clang-format off
#ifdef INSTANTIATE_ROT_CHECK_NUMERICS
#error INSTANTIATE_ROT_CHECK_NUMERICS already defined
#endif

#define INSTANTIATE_ROT_CHECK_NUMERICS(T_) \
template rocblas_status rocblas_rot_check_numerics<T_> \
(const char* function_name, \
rocblas_handle handle, \
int64_t n, \
T_ x, \
rocblas_stride offset_x, \
int64_t inc_x, \
rocblas_stride stride_x, \
T_ y, \
rocblas_stride offset_y, \
int64_t inc_y, \
rocblas_stride stride_y, \
int64_t batch_count, \
const int check_numerics, \
bool is_input);
#define INSTANTIATE_ROT_CHECK_NUMERICS(T_) \
template rocblas_status rocblas_rot_check_numerics<T_>(const char* function_name, \
rocblas_handle handle, \
int64_t n, \
T_ x, \
rocblas_stride offset_x, \
int64_t inc_x, \
rocblas_stride stride_x, \
T_ y, \
rocblas_stride offset_y, \
int64_t inc_y, \
rocblas_stride stride_y, \
int64_t batch_count, \
const int check_numerics, \
bool is_input);

// instantiate for rocblas_Xrot and rocblas_Xrot_strided_batched
INSTANTIATE_ROT_CHECK_NUMERICS(float* )
INSTANTIATE_ROT_CHECK_NUMERICS(double* )
INSTANTIATE_ROT_CHECK_NUMERICS(rocblas_float_complex* )
INSTANTIATE_ROT_CHECK_NUMERICS(rocblas_double_complex* )
INSTANTIATE_ROT_CHECK_NUMERICS(float*)
INSTANTIATE_ROT_CHECK_NUMERICS(double*)
INSTANTIATE_ROT_CHECK_NUMERICS(rocblas_float_complex*)
INSTANTIATE_ROT_CHECK_NUMERICS(rocblas_double_complex*)

// instantiate for rocblas_Xrot__batched
INSTANTIATE_ROT_CHECK_NUMERICS(float* const*)
Expand Down Expand Up @@ -168,7 +166,8 @@ rocblas_status rocblas_internal_rot_launcher(rocblas_handle handle,
c,
c_stride,
s,
s_stride, batch_count);
s_stride,
batch_count);
else // c and s are on host
ROCBLAS_LAUNCH_KERNEL((rocblas_rot_kernel<API_INT, NB, Tex>),
blocks,
Expand All @@ -187,7 +186,8 @@ rocblas_status rocblas_internal_rot_launcher(rocblas_handle handle,
*c,
c_stride,
*s,
s_stride, batch_count);
s_stride,
batch_count);

return rocblas_status_success;
}
Expand All @@ -203,42 +203,118 @@ rocblas_status rocblas_internal_rot_launcher(rocblas_handle handle,
rocblas_int n, \
Tx_ x, \
rocblas_stride offset_x, \
int64_t incx, \
int64_t incx, \
rocblas_stride stride_x, \
Ty_ y, \
rocblas_stride offset_y, \
int64_t incy, \
int64_t incy, \
rocblas_stride stride_y, \
Tc_ * c, \
rocblas_stride c_stride, \
Ts_ * s, \
rocblas_stride s_stride, \
rocblas_int batch_count);


// instantiate for rocblas_Xrot and rocblas_Xrot_strided_batched
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, float, float*, float*, float const, float const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, double, double*, double*, double const, double const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, float, rocblas_bfloat16*, rocblas_bfloat16*, rocblas_bfloat16 const, rocblas_bfloat16 const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, float, rocblas_half*, rocblas_half*, rocblas_half const, rocblas_half const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_float_complex, rocblas_float_complex*, rocblas_float_complex*, float const, float const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_float_complex, rocblas_float_complex*, rocblas_float_complex*, float const, rocblas_float_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_float_complex, rocblas_float_complex*, rocblas_float_complex*, rocblas_float_complex const, rocblas_float_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_double_complex, rocblas_double_complex*, rocblas_double_complex*, double const, double const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_double_complex, rocblas_double_complex*, rocblas_double_complex*, double const, rocblas_double_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_double_complex, rocblas_double_complex*, rocblas_double_complex*, rocblas_double_complex const, rocblas_double_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, float, float*, float*, float const, float const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, double, double*, double*, double const, double const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
float,
rocblas_bfloat16*,
rocblas_bfloat16*,
rocblas_bfloat16 const,
rocblas_bfloat16 const)
INSTANTIATE_ROT_LAUNCHER(
ROCBLAS_ROT_NB, float, rocblas_half*, rocblas_half*, rocblas_half const, rocblas_half const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_float_complex,
rocblas_float_complex*,
rocblas_float_complex*,
float const,
float const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_float_complex,
rocblas_float_complex*,
rocblas_float_complex*,
float const,
rocblas_float_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_float_complex,
rocblas_float_complex*,
rocblas_float_complex*,
rocblas_float_complex const,
rocblas_float_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_double_complex,
rocblas_double_complex*,
rocblas_double_complex*,
double const,
double const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_double_complex,
rocblas_double_complex*,
rocblas_double_complex*,
double const,
rocblas_double_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_double_complex,
rocblas_double_complex*,
rocblas_double_complex*,
rocblas_double_complex const,
rocblas_double_complex const)

// instantiate for rocblas_Xrot__batched
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, float, float* const*, float* const*, float const, float const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, double, double* const*, double* const*, double const, double const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, float, rocblas_bfloat16* const*, rocblas_bfloat16* const*, rocblas_bfloat16 const, rocblas_bfloat16 const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, float, rocblas_half* const*, rocblas_half* const*, rocblas_half const, rocblas_half const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_float_complex, rocblas_float_complex* const*, rocblas_float_complex* const*, float const, float const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_float_complex, rocblas_float_complex* const*, rocblas_float_complex* const*, float const, rocblas_float_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_float_complex, rocblas_float_complex* const*, rocblas_float_complex* const*, rocblas_float_complex const, rocblas_float_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_double_complex, rocblas_double_complex* const*, rocblas_double_complex* const*, rocblas_double_complex const, rocblas_double_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_double_complex, rocblas_double_complex* const*, rocblas_double_complex* const*, double const, double const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB, rocblas_double_complex, rocblas_double_complex* const*, rocblas_double_complex* const*, double const, rocblas_double_complex const)
INSTANTIATE_ROT_LAUNCHER(
ROCBLAS_ROT_NB, float, float* const*, float* const*, float const, float const)
INSTANTIATE_ROT_LAUNCHER(
ROCBLAS_ROT_NB, double, double* const*, double* const*, double const, double const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
float,
rocblas_bfloat16* const*,
rocblas_bfloat16* const*,
rocblas_bfloat16 const,
rocblas_bfloat16 const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
float,
rocblas_half* const*,
rocblas_half* const*,
rocblas_half const,
rocblas_half const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_float_complex,
rocblas_float_complex* const*,
rocblas_float_complex* const*,
float const,
float const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_float_complex,
rocblas_float_complex* const*,
rocblas_float_complex* const*,
float const,
rocblas_float_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_float_complex,
rocblas_float_complex* const*,
rocblas_float_complex* const*,
rocblas_float_complex const,
rocblas_float_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_double_complex,
rocblas_double_complex* const*,
rocblas_double_complex* const*,
rocblas_double_complex const,
rocblas_double_complex const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_double_complex,
rocblas_double_complex* const*,
rocblas_double_complex* const*,
double const,
double const)
INSTANTIATE_ROT_LAUNCHER(ROCBLAS_ROT_NB,
rocblas_double_complex,
rocblas_double_complex* const*,
rocblas_double_complex* const*,
double const,
rocblas_double_complex const)

#undef INSTANTIATE_ROT_LAUNCHER
// clang-format on
Loading
Loading