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/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 528170210c11..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,106 +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)); - // 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); - 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 = 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; + 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; }