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.

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

Jump to bottom

Get rid of legacy 2GiB offset limits in CallGemm*() and transpose*() internal APIs and kernels. #2613

Merged
merged 5 commits into from
Dec 20, 2023
Merged

Get rid of legacy 2GiB offset limits in CallGemm*() and transpose*() internal APIs and kernels. #2613

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
391f40b
fix-issue-2593(01) Get rid of 2GiB offset limit in transpose_packed_M…
atamazov Dec 15, 2023
d955531
fix-issue-2593(02) Get rid of legacy 2GiB offset limit in CallGemm*()
atamazov Dec 15, 2023
3e9c430
fix-issue-2593(03) Get rid of 2GiB offset limit in transpose_NCHW2CNH…
atamazov Dec 15, 2023
b2f33d6
fix-issue-2593(04) Get rid of 2GiB offset limit in transpose_CNHW2NCH…
atamazov Dec 15, 2023
1eb2554
fix-issue-2593(05) Fix issue in transpose_CNHW2NCHW()
atamazov Dec 18, 2023
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
36 changes: 18 additions & 18 deletions src/gemm_v2.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -142,11 +142,11 @@ template <typename T>
rocblas_status miopen_rocblas_gemm_ex3(const miopen::Handle& handle,
const miopen::GemmDescriptor& gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset)
std::size_t c_offset)
{
rocblas_status rb_status =
rocblas_status::rocblas_status_internal_error; // cppcheck-suppress redundantInitialization
Expand Down Expand Up @@ -231,11 +231,11 @@ template <typename T>
rocblas_status miopen_rocblas_gemm_strided_batched_ex3(const miopen::Handle& handle,
const miopen::GemmDescriptor& gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset)
std::size_t c_offset)
{
rocblas_status rb_status = rocblas_status::rocblas_status_internal_error;
// Until there is a batched counter part to the ex3 rocBlas call we need to iterate over the
Expand Down Expand Up @@ -366,11 +366,11 @@ static GemmBackend_t enforce_gemm_backend(miopenDataType_t data_type,
miopenStatus_t CallGemmTimeMeasure(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
bool time_precision,
CallGemmType_t call_gemm_type,
GemmBackend_t gemm_backend)
Expand Down Expand Up @@ -415,11 +415,11 @@ miopenStatus_t CallGemmTimeMeasure(const Handle& handle,
miopenStatus_t CallGemm(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend)
{
MIOPEN_LOG_I2("gemm_desc: " << gemm_desc);
Expand Down Expand Up @@ -670,11 +670,11 @@ miopenStatus_t CallGemm(const Handle& handle,
miopenStatus_t CallGemmStridedBatched(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend)
{
MIOPEN_LOG_I2("gemm_desc: " << gemm_desc);
Expand Down Expand Up @@ -946,11 +946,11 @@ miopenStatus_t CallGemmStridedBatched(const Handle& handle,
miopenStatus_t CallGemmStridedBatchedSequential(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend)
{
MIOPEN_LOG_I2("gemm_desc: " << gemm_desc);
Expand Down
24 changes: 12 additions & 12 deletions src/include/miopen/gemm_v2.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -124,44 +124,44 @@ struct GemmDescriptor
miopenStatus_t CallGemmTimeMeasure(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
bool time_precision,
CallGemmType_t call_gemm_type,
GemmBackend_t gemm_backend = GemmBackend_t::rocblas);

miopenStatus_t CallGemm(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend = GemmBackend_t::rocblas);

miopenStatus_t CallGemmStridedBatched(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend = GemmBackend_t::rocblas);

miopenStatus_t
CallGemmStridedBatchedSequential(const Handle& handle,
GemmDescriptor gemm_desc,
ConstData_t A,
int a_offset,
std::size_t a_offset,
ConstData_t B,
int b_offset,
std::size_t b_offset,
Data_t C,
int c_offset,
std::size_t c_offset,
GemmBackend_t gemm_backend = GemmBackend_t::rocblas);

// GEMM parameters for Convolution (using Im2Col) Fwd
Expand Down
12 changes: 6 additions & 6 deletions src/include/miopen/util.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -74,8 +74,8 @@ float transpose_NCHW2CNHW(const Handle& handle,
int w_out,
ConstData_t in,
Data_t out,
int in_offset,
int out_offset,
std::size_t in_offset,
std::size_t out_offset,
int h_stride,
int w_stride,
miopenDataType_t type);
Expand All @@ -89,8 +89,8 @@ float transpose_CNHW2NCHW(const Handle& handle,
int w_in,
ConstData_t in,
Data_t out,
int in_offset,
int out_offset,
std::size_t in_offset,
std::size_t out_offset,
int h_stride,
int w_stride,
miopenDataType_t type);
Expand All @@ -108,8 +108,8 @@ float transpose_NCHW2Vec(const Handle& handle,
float transpose_packed_MN2NM(const Handle& handle,
int m,
int n,
int in_offset,
int out_offset,
std::size_t in_offset,
std::size_t out_offset,
ConstData_t in,
Data_t out,
miopenDataType_t type);
Expand Down
Loading
Loading