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

Fix the dynamic reduction GPU Memory Acccess failure with double on ROCM 4.3 #1131

Merged
merged 5 commits into from
Sep 7, 2021
Merged

Fix the dynamic reduction GPU Memory Acccess failure with double on ROCM 4.3 #1131

Changes from 1 commit
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
5f4deed
Fix the calculation of ws_buf2_bytes_offset for dynamic reduction in …
qianfengz Aug 31, 2021
2822d9d
Just remove IsDynamicReductionEnabled()
qianfengz Sep 1, 2021
867b235
Tiny fix in ReduceTensorDescriptor::GetWorkspaceSize()
qianfengz Sep 1, 2021
4f81393
Update to the calculation of ws_buf2_bytes_offset
qianfengz Sep 1, 2021
fdd2c74
Merge branch 'develop' into dynamic_reduction_fix
qianfengz Sep 1, 2021
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
12 changes: 5 additions & 7 deletions src/reducetensor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,7 +46,6 @@
MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_DYNAMIC_REDUCTION);

#define WORKAROUND_MIOPEN_ISSUE_557 1
#define WORKAROUND_ISSUE_1123 ((HIP_PACKAGE_VERSION_FLAT == 4003000000ULL) && MIOPEN_BACKEND_HIP)

namespace miopen {

Expand All @@ -62,11 +61,7 @@ namespace detail {

static bool IsDynamicReductionEnabled()
junliume marked this conversation as resolved.
Show resolved Hide resolved
{
#if WORKAROUND_ISSUE_1123
return miopen::IsEnabled(MIOPEN_DEBUG_DYNAMIC_REDUCTION{});
#else
return !miopen::IsDisabled(MIOPEN_DEBUG_DYNAMIC_REDUCTION{});
#endif
}

struct get_tunable_reduction_kernel_constants
Expand Down Expand Up @@ -568,10 +563,13 @@ void ReduceTensorDescriptor::ReduceTensor(const Handle& handle,

if(need_indices && workspace != nullptr)
{
std::size_t aTypeSize = detail::GetDataTypeSize(aDesc.GetType());
std::size_t aTypeSize = detail::GetDataTypeSize(aDesc.GetType());
auto ws_buf12_sizeInBytes = detail::IsDynamicReductionEnabled()
? workspaceSizeInBytes - 4096
: workspaceSizeInBytes;

long byteOffset =
static_cast<long>((workspaceSizeInBytes / (aTypeSize + sizeof(int))) * aTypeSize);
static_cast<long>((ws_buf12_sizeInBytes / (aTypeSize + sizeof(int))) * aTypeSize);
junliume marked this conversation as resolved.
Show resolved Hide resolved

ws_buf2_bytes_offset = ((byteOffset + 63) / 64) * 64;
};
Expand Down