Skip to content
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.

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

Hotfix: Fix offset calculation to prevent overflow if offset is really large #1298

Merged
merged 11 commits into from
Mar 28, 2023
27 changes: 27 additions & 0 deletions clients/gtest/trsm_gtest.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -60,6 +60,15 @@ Definitions:
- &large_memory_matrix_size_range
- { M: 8320, N: 128, lda: 8320, ldb: 8320 }

- &size_t_left_matrix_size_range
# - { M: 4, N: 46435, lda: 4, ldb: 46435 }
- { M: 46345, N: 4, lda: 46345, ldb: 46345 }
# - { M: 47000, N: 4, lda: 47000, ldb: 47000 } # calls rocblas_internal_gemm_template with batch_count=367, stride_a=6016128

- &size_t_right_matrix_size_range
- { M: 4, N: 46345, lda: 46345, ldb: 4 }
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These tests timed out for me on gfx90a, but passed when I increased the timeout threshold. I'm testing all other configs as well just to be safe.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@daineAMD Naveen's speedup for the initialization of the triangle matrix in https://github.com/ROCmSoftwarePlatform/rocBLAS-internal/pull/1551 . This PR is in the develop branch but it is not in release/rocm-rel-5.5. Tests will run faster on the develop branch than on release/rocm-rel-5.5.

# - { M: 4, N: 47000, lda: 47000, ldb: 4 } # calls rocblas_internal_gemm_template with batch_count=367, stride_a=6016128

- &substitution_size_range_thorough
- { M: 1, N: 1, lda: 100, ldb: 100 }
- { M: 1, N: 32, lda: 100, ldb: 100 }
Expand Down Expand Up @@ -486,6 +495,24 @@ Tests:
matrix_size: *testset2_matrix_size_range
alpha: [ 1 ]

- name: trsm_size_t_left
category: nightly
function: trsm
precision: *single_precision
arguments:
- { side: L, uplo: L, transA: N, diag: N }
matrix_size: *size_t_left_matrix_size_range
alpha: [2]

- name: trsm_size_t_right
category: nightly
function: trsm
precision: *single_precision
arguments:
- { side: R, uplo: L, transA: N, diag: N }
matrix_size: *size_t_right_matrix_size_range
alpha: [2]

- name: trsm_large
category: nightly
function: trsm
Expand Down
40 changes: 27 additions & 13 deletions library/src/blas3/Tensile/gemm_tensile.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -139,19 +139,33 @@ inline rocblas_status call_tensile(rocblas_handle handle,
}
#endif

RocblasContractionProblem<T> problem{handle, trans_a,
trans_b, m,
n, k,
alpha, A,
nullptr, ld_a,
stride_a, offset_a,
B, nullptr,
ld_b, stride_b,
offset_b, beta,
C, nullptr,
ld_c, stride_c,
offset_c, batch_count,
true, rocblas_gemm_flags_none};
// pre apply offsets for non-batched and strided
RocblasContractionProblem<T> problem{handle,
trans_a,
trans_b,
m,
n,
k,
alpha,
A + offset_a,
nullptr,
ld_a,
stride_a,
0 /* offset_a */,
B + offset_b,
nullptr,
ld_b,
stride_b,
0 /* offset_b */,
beta,
C + offset_c,
nullptr,
ld_c,
stride_c,
0 /* offset_c */,
batch_count,
true,
rocblas_gemm_flags_none};

return runContractionProblem(problem);
}
Expand Down
233 changes: 120 additions & 113 deletions library/src/blas3/rocblas_trsm.hpp

Large diffs are not rendered by default.

Loading