From 3a33ab4ea3960191d04c5b702916a46a6b4fe5c3 Mon Sep 17 00:00:00 2001 From: Torre Zuk Date: Mon, 5 Oct 2026 14:13:37 -0600 Subject: [PATCH 1/2] rot and rot_ex fix ILP64 large N inc handling --- projects/rocblas/clients/gtest/rot_gtest.yaml | 2 + .../src64/blas1/rocblas_rot_kernels_64.cpp | 47 ++++++++++--------- 2 files changed, 28 insertions(+), 21 deletions(-) diff --git a/projects/rocblas/clients/gtest/rot_gtest.yaml b/projects/rocblas/clients/gtest/rot_gtest.yaml index b4181f6bdbc8..53c53bf281a2 100644 --- a/projects/rocblas/clients/gtest/rot_gtest.yaml +++ b/projects/rocblas/clients/gtest/rot_gtest.yaml @@ -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 } - { N: 2, incx: -2147483649, incy: 1, batch_count: 1 } - { N: 2, incx: 1, incy: -2147483649, batch_count: 1 } api: [ C_64 ] diff --git a/projects/rocblas/library/src/src64/blas1/rocblas_rot_kernels_64.cpp b/projects/rocblas/library/src/src64/blas1/rocblas_rot_kernels_64.cpp index 528170210c11..6a3035aff5e4 100644 --- a/projects/rocblas/library/src/src64/blas1/rocblas_rot_kernels_64.cpp +++ b/projects/rocblas/library/src/src64/blas1/rocblas_rot_kernels_64.cpp @@ -89,23 +89,29 @@ rocblas_status rocblas_internal_rot_launcher_64(rocblas_handle handle, for(int64_t n_base = 0; n_base < n_64; n_base += c_i64_grid_X_chunk) { int32_t n = int32_t(std::min(n_64 - n_base, c_i64_grid_X_chunk)); + + int64_t shiftx + = offset_x + (incx_64 < 0 ? -incx_64 * (n_64 - n - n_base) : n_base * incx_64); + int64_t shifty + = offset_y + (incy_64 < 0 ? -incy_64 * (n_64 - n - n_base) : n_base * incy_64); + // 32bit API call as incx/y int64_t - rocblas_status status = rocblas_internal_rot_launcher( - handle, - rocblas_int(n), - x_ptr, - offset_x + n_base * incx_64, - incx_64, - stride_x, - y_ptr, - offset_y + n_base * incy_64, - incy_64, - stride_y, - c_ptr, - c_stride, - s_ptr, - s_stride, - batch_count); + rocblas_status status + = rocblas_internal_rot_launcher(handle, + rocblas_int(n), + x_ptr, + shiftx, + incx_64, + stride_x, + y_ptr, + shifty, + incy_64, + stride_y, + c_ptr, + c_stride, + s_ptr, + s_stride, + batch_count); if(status != rocblas_status_success) return status; } @@ -126,11 +132,10 @@ rocblas_status rocblas_internal_rot_launcher_64(rocblas_handle handle, { int32_t n = int32_t(std::min(n_64 - n_base, c_i64_grid_X_chunk)); - int64_t shiftx = incx_64 < 0 ? -incx_64 * n_base : incx_64 * n_base; - int64_t shifty = incy_64 < 0 ? -incy_64 * n_base : incy_64 * n_base; - - shiftx += offset_x; - shifty += offset_y; + int64_t shiftx + = offset_x + (incx_64 < 0 ? -incx_64 * (n_64 - n - n_base) : n_base * incx_64); + int64_t shifty + = offset_y + (incy_64 < 0 ? -incy_64 * (n_64 - n - n_base) : n_base * incy_64); // new instantiation for 64bit incx/y rocblas_status status From d94b3213dd67a67a4b9d993075e5e20215c8b339 Mon Sep 17 00:00:00 2001 From: Torre Zuk Date: Wed, 7 Oct 2026 14:33:37 -0600 Subject: [PATCH 2/2] minimize code to match current optimizations --- .../library/src/blas1/rocblas_rot_kernels.cpp | 170 +++++++++++++----- .../src64/blas1/rocblas_rot_kernels_64.cpp | 152 ++++++---------- 2 files changed, 177 insertions(+), 145 deletions(-) diff --git a/projects/rocblas/library/src/blas1/rocblas_rot_kernels.cpp b/projects/rocblas/library/src/blas1/rocblas_rot_kernels.cpp index 900c46bd8757..8e49c7b8eec9 100644 --- a/projects/rocblas/library/src/blas1/rocblas_rot_kernels.cpp +++ b/projects/rocblas/library/src/blas1/rocblas_rot_kernels.cpp @@ -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 \ - (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(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*) @@ -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), blocks, @@ -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; } @@ -203,11 +203,11 @@ 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, \ @@ -215,30 +215,106 @@ rocblas_status rocblas_internal_rot_launcher(rocblas_handle handle, 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 diff --git a/projects/rocblas/library/src/src64/blas1/rocblas_rot_kernels_64.cpp b/projects/rocblas/library/src/src64/blas1/rocblas_rot_kernels_64.cpp index 6a3035aff5e4..d8d8d350a8f7 100644 --- a/projects/rocblas/library/src/src64/blas1/rocblas_rot_kernels_64.cpp +++ b/projects/rocblas/library/src/src64/blas1/rocblas_rot_kernels_64.cpp @@ -55,111 +55,67 @@ rocblas_status rocblas_internal_rot_launcher_64(rocblas_handle handle, if(n_64 <= 0 || batch_count_64 <= 0) return rocblas_status_success; - if(std::abs(incx_64) <= c_ILP64_i32_max && std::abs(incy_64) < c_ILP64_i32_max) + if(n_64 <= c_ILP64_i32_max && batch_count_64 < c_i64_grid_YZ_chunk) { - if(n_64 <= c_ILP64_i32_max && batch_count_64 < c_i64_grid_YZ_chunk) - { - // valid to use original 32bit API with truncated 64bit args - return rocblas_internal_rot_launcher(handle, - rocblas_int(n_64), - x, - offset_x, - incx_64, - stride_x, - y, - offset_y, - incy_64, - stride_y, - c, - c_stride, - s, - s_stride, - batch_count_64); - } - - for(int64_t b_base = 0; b_base < batch_count_64; b_base += c_i64_grid_YZ_chunk) - { - auto x_ptr = adjust_ptr_batch(x, b_base, stride_x); - auto y_ptr = adjust_ptr_batch(y, b_base, stride_y); - auto c_ptr = adjust_ptr_batch(c, b_base, c_stride); - auto s_ptr = adjust_ptr_batch(s, b_base, s_stride); - - int32_t batch_count = int32_t(std::min(batch_count_64 - b_base, c_i64_grid_YZ_chunk)); - - for(int64_t n_base = 0; n_base < n_64; n_base += c_i64_grid_X_chunk) - { - int32_t n = int32_t(std::min(n_64 - n_base, c_i64_grid_X_chunk)); - - int64_t shiftx - = offset_x + (incx_64 < 0 ? -incx_64 * (n_64 - n - n_base) : n_base * incx_64); - int64_t shifty - = offset_y + (incy_64 < 0 ? -incy_64 * (n_64 - n - n_base) : n_base * incy_64); - - // 32bit API call as incx/y int64_t - rocblas_status status - = rocblas_internal_rot_launcher(handle, - rocblas_int(n), - x_ptr, - shiftx, - incx_64, - stride_x, - y_ptr, - shifty, - incy_64, - stride_y, - c_ptr, - c_stride, - s_ptr, - s_stride, - batch_count); - if(status != rocblas_status_success) - return status; - } - } + // valid to use original 32bit API with truncated 64bit args + return rocblas_internal_rot_launcher(handle, + rocblas_int(n_64), + x, + offset_x, + incx_64, + stride_x, + y, + offset_y, + incy_64, + stride_y, + c, + c_stride, + s, + s_stride, + batch_count_64); } - else - { - for(int64_t b_base = 0; b_base < batch_count_64; b_base += c_i64_grid_YZ_chunk) - { - auto x_ptr = adjust_ptr_batch(x, b_base, stride_x); - auto y_ptr = adjust_ptr_batch(y, b_base, stride_y); - auto c_ptr = adjust_ptr_batch(c, b_base, c_stride); - auto s_ptr = adjust_ptr_batch(s, b_base, s_stride); - - int32_t batch_count = int32_t(std::min(batch_count_64 - b_base, c_i64_grid_YZ_chunk)); - for(int64_t n_base = 0; n_base < n_64; n_base += c_i64_grid_X_chunk) - { - int32_t n = int32_t(std::min(n_64 - n_base, c_i64_grid_X_chunk)); - - int64_t shiftx - = offset_x + (incx_64 < 0 ? -incx_64 * (n_64 - n - n_base) : n_base * incx_64); - int64_t shifty - = offset_y + (incy_64 < 0 ? -incy_64 * (n_64 - n - n_base) : n_base * incy_64); + for(int64_t b_base = 0; b_base < batch_count_64; b_base += c_i64_grid_YZ_chunk) + { + auto x_ptr = adjust_ptr_batch(x, b_base, stride_x); + auto y_ptr = adjust_ptr_batch(y, b_base, stride_y); + auto c_ptr = adjust_ptr_batch(c, b_base, c_stride); + auto s_ptr = adjust_ptr_batch(s, b_base, s_stride); - // new instantiation for 64bit incx/y - rocblas_status status - = rocblas_internal_rot_launcher(handle, - n, - x_ptr, - shiftx, - incx_64, - stride_x, - y_ptr, - shifty, - incy_64, - stride_y, - c_ptr, - c_stride, - s_ptr, - s_stride, - batch_count); + int32_t batch_count = int32_t(std::min(batch_count_64 - b_base, c_i64_grid_YZ_chunk)); - if(status != rocblas_status_success) - return status; - } + for(int64_t n_base = 0; n_base < n_64; n_base += c_i64_grid_X_chunk) + { + int32_t n = int32_t(std::min(n_64 - n_base, c_i64_grid_X_chunk)); + + int64_t shiftx + = offset_x + (incx_64 < 0 ? -incx_64 * (n_64 - n - n_base) : n_base * incx_64); + int64_t shifty + = offset_y + (incy_64 < 0 ? -incy_64 * (n_64 - n - n_base) : n_base * incy_64); + + // TODO potentially optimize with instantiate incx/y int32 and int64 + // only n is template API_INT, incx/y are int64_t + rocblas_status status + = rocblas_internal_rot_launcher(handle, + n, + x_ptr, + shiftx, + incx_64, + stride_x, + y_ptr, + shifty, + incy_64, + stride_y, + c_ptr, + c_stride, + s_ptr, + s_stride, + batch_count); + if(status != rocblas_status_success) + return status; } } + return rocblas_status_success; }