Skip to content

[MIOpen] [ROCM-21512] [ALMIOPEN-1814] Fix MIOpen ASan memory leaks #7008

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

Merged
NolanHannaAMD merged 7 commits into from
May 8, 2026
Merged

[MIOpen] [ROCM-21512] [ALMIOPEN-1814] Fix MIOpen ASan memory leaks #7008

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
901beaa
[MIOpen] Fixing the blacklisted ASan memory leaks and removing them f…
NolanHannaAMD Apr 24, 2026
2e70740
[MIOpen] Reducing the amount of memory the supertensor tests use so t…
NolanHannaAMD Apr 24, 2026
01c374e
Addressing some code review comments.
NolanHannaAMD Apr 29, 2026
3db6bb9
Fixing some MIOpen gtest memory leaks that were detected with Address…
NolanHannaAMD May 1, 2026
2caa599
Merge branch 'develop' into users/nhanna/fix-miopen-asan-memory-leaks
NolanHannaAMD May 6, 2026
52b28f7
Merge branch 'develop' into users/nhanna/fix-miopen-asan-memory-leaks
JonathanLichtnerAMD May 7, 2026
2b72e71
Moving HandleGuard into it's own file as per code review comment.
NolanHannaAMD May 7, 2026
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
49 changes: 22 additions & 27 deletions projects/miopen/test/cpu_rnn.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,7 @@
* LSTM CPU verification functions
**********************************************/
#include "gemm.hpp"
#include "gtest/gtest_desc_guard.hpp"

template <class T>
void LSTMFwdCPUVerify(const miopen::Handle& handle,
Expand Down Expand Up @@ -92,7 +93,8 @@ void LSTMFwdCPUVerify(const miopen::Handle& handle,
std::vector<rocrand_state_xorwow> dropout_states_host;
std::vector<unsigned char> dropout_reservespace_host;
std::vector<T> dropout_hid_state;
miopenTensorDescriptor_t dropout_inputTensor{}, dropout_outputTensor{};
TensorDescGuard dropout_inputTensor;
TensorDescGuard dropout_outputTensor;
if(use_dropout)
{
size_t states_size = dropoutDesc.stateSizeInBytes / sizeof(rocrand_state_xorwow);
Expand All @@ -102,8 +104,6 @@ void LSTMFwdCPUVerify(const miopen::Handle& handle,
std::array<int, 2> drop_in_len = {{batch_n_cpu, hy_h * bi}};
std::array<int, 2> drop_in_str = {{hy_stride, 1}};
std::array<int, 2> drop_out_str = {{hy_h * bi, 1}};
miopenCreateTensorDescriptor(&dropout_inputTensor);
miopenCreateTensorDescriptor(&dropout_outputTensor);
miopenSetTensorDescriptor(
dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data());
miopenSetTensorDescriptor(
Expand Down Expand Up @@ -204,9 +204,9 @@ void LSTMFwdCPUVerify(const miopen::Handle& handle,

DropoutForwardVerify<T>(handle,
dropoutDesc,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
rsvspace,
miopen::deref(dropout_outputTensor),
miopen::deref(dropout_outputTensor.get()),
dropout_hid_state,
dropout_reservespace_host,
dropout_states_tmp,
Expand Down Expand Up @@ -710,13 +710,12 @@ void LSTMBwdDataCPUVerify(bool use_dropout_cpu,
}

// initial dropoput
miopenTensorDescriptor_t dropout_inputTensor{};
TensorDescGuard dropout_inputTensor;
std::vector<unsigned char> dropout_reservespace_host;
if(use_dropout_cpu)
{
std::array<int, 2> drop_in_len = {{batch_n_cpu, hy_h * bi}};
std::array<int, 2> drop_in_str = {{hy_stride, 1}};
miopenCreateTensorDescriptor(&dropout_inputTensor);
miopenSetTensorDescriptor(
dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data());

Expand Down Expand Up @@ -781,9 +780,9 @@ void LSTMBwdDataCPUVerify(bool use_dropout_cpu,
if(use_dropout_cpu)
{
DropoutBackwardVerify<T>(dropoutDesc,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
wkspace,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
wkspace,
dropout_reservespace_host,
hid_shift + bi * 5 * hy_h,
Expand Down Expand Up @@ -1601,7 +1600,8 @@ void RNNFwdTrainCPUVerify(const miopen::Handle& handle,
std::vector<rocrand_state_xorwow> dropout_states_host;
std::vector<unsigned char> dropout_reservespace_host;
std::vector<T> dropout_hid_state;
miopenTensorDescriptor_t dropout_inputTensor{}, dropout_outputTensor{};
TensorDescGuard dropout_inputTensor;
TensorDescGuard dropout_outputTensor;
if(use_dropout)
{
size_t states_size = dropoutDesc.stateSizeInBytes / sizeof(rocrand_state_xorwow);
Expand All @@ -1611,8 +1611,6 @@ void RNNFwdTrainCPUVerify(const miopen::Handle& handle,
std::array<int, 2> drop_in_len = {{batch_n, hy_h * bi}};
std::array<int, 2> drop_in_str = {{hy_stride, 1}};
std::array<int, 2> drop_out_str = {{hy_h * bi, 1}};
miopenCreateTensorDescriptor(&dropout_inputTensor);
miopenCreateTensorDescriptor(&dropout_outputTensor);
miopenSetTensorDescriptor(
dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data());
miopenSetTensorDescriptor(
Expand Down Expand Up @@ -1708,9 +1706,9 @@ void RNNFwdTrainCPUVerify(const miopen::Handle& handle,

DropoutForwardVerify<T>(handle,
dropoutDesc,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
rsvspace,
miopen::deref(dropout_outputTensor),
miopen::deref(dropout_outputTensor.get()),
dropout_hid_state,
dropout_reservespace_host,
dropout_states_tmp,
Expand Down Expand Up @@ -2070,13 +2068,12 @@ void RNNBwdDataCPUVerify(bool use_dropout,
}

// initial dropoput
miopenTensorDescriptor_t dropout_inputTensor{};
TensorDescGuard dropout_inputTensor;
std::vector<unsigned char> dropout_reservespace_host;
if(use_dropout)
{
std::array<int, 2> drop_in_len = {{batch_n, hy_h * bi}};
std::array<int, 2> drop_in_str = {{hy_stride, 1}};
miopenCreateTensorDescriptor(&dropout_inputTensor);
miopenSetTensorDescriptor(
dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data());

Expand Down Expand Up @@ -2140,9 +2137,9 @@ void RNNBwdDataCPUVerify(bool use_dropout,
if(use_dropout)
{
DropoutBackwardVerify<T>(dropoutDesc,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
wkspace,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
wkspace,
dropout_reservespace_host,
hid_shift,
Expand Down Expand Up @@ -2748,7 +2745,8 @@ void GRUFwdCPUVerify(const miopen::Handle& handle,
std::vector<rocrand_state_xorwow> dropout_states_host;
std::vector<unsigned char> dropout_reservespace_host;
std::vector<T> dropout_hid_state;
miopenTensorDescriptor_t dropout_inputTensor{}, dropout_outputTensor{};
TensorDescGuard dropout_inputTensor;
TensorDescGuard dropout_outputTensor;
if(use_dropout)
{
size_t states_size = dropoutDesc.stateSizeInBytes / sizeof(rocrand_state_xorwow);
Expand All @@ -2758,8 +2756,6 @@ void GRUFwdCPUVerify(const miopen::Handle& handle,
std::array<int, 2> drop_in_len = {{batch_n, hy_h * bi}};
std::array<int, 2> drop_in_str = {{hy_stride, 1}};
std::array<int, 2> drop_out_str = {{hy_h * bi, 1}};
miopenCreateTensorDescriptor(&dropout_inputTensor);
miopenCreateTensorDescriptor(&dropout_outputTensor);
miopenSetTensorDescriptor(
dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data());
miopenSetTensorDescriptor(
Expand Down Expand Up @@ -2858,9 +2854,9 @@ void GRUFwdCPUVerify(const miopen::Handle& handle,

DropoutForwardVerify<T>(handle,
dropoutDesc,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
rsvspace,
miopen::deref(dropout_outputTensor),
miopen::deref(dropout_outputTensor.get()),
dropout_hid_state,
dropout_reservespace_host,
dropout_states_tmp,
Expand Down Expand Up @@ -3517,13 +3513,12 @@ void GRUBwdDataCPUVerify(bool use_dropout,
}

// initial dropoput
miopenTensorDescriptor_t dropout_inputTensor{};
TensorDescGuard dropout_inputTensor;
std::vector<unsigned char> dropout_reservespace_host;
if(use_dropout)
{
std::array<int, 2> drop_in_len = {{batch_n, hy_h * bi}};
std::array<int, 2> drop_in_str = {{hy_stride, 1}};
miopenCreateTensorDescriptor(&dropout_inputTensor);
miopenSetTensorDescriptor(
dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data());

Expand Down Expand Up @@ -3589,9 +3584,9 @@ void GRUBwdDataCPUVerify(bool use_dropout,
if(use_dropout)
{
DropoutBackwardVerify<T>(dropoutDesc,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
wkspace,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
wkspace,
dropout_reservespace_host,
hid_shift + bi * 3 * hy_h,
Expand Down
22 changes: 0 additions & 22 deletions projects/miopen/test/gtest/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,28 +36,6 @@ if(THEROCK_SANITIZER STREQUAL "ASAN")
list(APPEND SKIP_TESTS bad_fusion_plan.cpp cba_find2_infer.cpp cba_infer.cpp ck_builder_shared.cpp ck_builder_xdl.cpp conv_activ_infer.cpp conv_ai_3d_kernel_tuning_utils.cpp conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp conv_hip_igemm_xdlops.cpp conv_igemm_mlir_xdlops_bwd_wrw.cpp conv_igemm_mlir_xdlops_fwd.cpp find_2_conv.cpp find_db.cpp fused_conv_bias_res_add_activ.cpp graphapi_conv_bias_res_add_activ_fwd.cpp group_conv_deterministic_split_k.cpp group_conv2d_fwd.cpp group_conv2d_bwd.cpp group_conv2d_wrw.cpp group_conv3d_fwd.cpp group_conv3d_bwd.cpp group_conv3d_wrw.cpp find_mode_trust_verify.cpp kernel_tuning_net.cpp miopendriver_conv_immed.cpp miopendriver_conv2d_trans.cpp miopendriver_gemm.cpp miopendriver_regression_big_tensor.cpp miopendriver_regression_big_tensor.cpp miopendriver_regression_half_gfx9.cpp miopendriver_regression_half.cpp perf_config_HipImplicitGemm3DGroupFwdXdlops.cpp smoke_solver_ConvCkIgemmFwdV6r1DlopsNchw.cpp unit_conv_solver_ConvAsmImplicitGemmGTCDynamicBwdXdlops.cpp unit_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdXdlops.cpp unit_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC.cpp unit_conv_solver_ConvAsmImplicitGemmGTCDynamicWrwXdlops.cpp unit_conv_solver_ConvCkGroupedConvFwd.cpp unit_conv_solver_ConvHipImplicitGemm3DGroupBwdXdlops.cpp unit_conv_solver_ConvHipImplicitGemm3DGroupFwdXdlops.cpp unit_conv_solver_ConvHipImplicitGemm3DGroupWrwXdlops.cpp unit_conv_solver_ConvHipImplicitGemmBwdDataV1R1Xdlops.cpp unit_conv_solver_ConvHipImplicitGemmBwdDataV4R1Xdlops.cpp unit_conv_solver_ConvHipImplicitGemmFwdXdlops.cpp unit_conv_solver_ConvHipImplicitGemmGroupBwdXdlops.cpp unit_conv_solver_ConvHipImplicitGemmForwardV4R4Xdlops.cpp unit_conv_solver_ConvHipImplicitGemmForwardV4R5Xdlops.cpp unit_conv_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm.cpp unit_conv_solver_ConvHipImplicitGemmWrwV4R4Xdlops.cpp unit_conv_solver_ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm.cpp unit_conv_solver_ConvHipImplicitGemmBwdXdlops.cpp unit_conv_solver_ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC.cpp unit_conv_solver_ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC.cpp unit_conv_solver_ConvHipImplicitGemmGroupFwdXdlops.cpp unit_conv_solver_ConvHipImplicitGemmGroupWrwXdlops.cpp unit_implicitgemm_ck_util.cpp smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16.cpp smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16.cpp)
# temporarily disable due to test failures
list(APPEND SKIP_TESTS unit_conv_solver_ConvWinoRageRxS.cpp)
# Disable tests that leak memory when Address Sanitizer is enabled. These tests need to be investigated and fixed.
list(APPEND SKIP_TESTS
gpu_mha_backward.cpp
gpu_mha_forward.cpp
graphapi_convolution.cpp
graphapi_enginecfg.cpp
graphapi_engineheur.cpp
graphapi_execution_plan.cpp
graphapi_matmul.cpp
graphapi_operationgraph_descriptor.cpp
graphapi_operation_reduction.cpp
graphapi_operation_reshape.cpp
graphapi_operation_matmul.cpp
graphapi_operation_pointwise.cpp
graphapi_operation_rng.cpp
graphapi_pointwise.cpp
graphapi_reduction.cpp
graphapi_rng.cpp
graphapi_tensor.cpp
graphapi_variant_pack.cpp
mha_find20.cpp
w_supertensor.cpp)
endif()

function(add_gtest_negative_filter NEGATIVE_FILTER_TO_ADD)
Expand Down
5 changes: 5 additions & 0 deletions projects/miopen/test/gtest/gpu_mha_backward.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -376,6 +376,11 @@ class Test_Bwd_Mha : public testing::TestWithParam<TestCase>
checkOutput(miopenTensorMhaDK, "tensor dK", dKDesc_ref);
checkOutput(miopenTensorMhaDV, "tensor dV", dVDesc_ref);
}

for(auto& solution : solutions)
{
ASSERT_EQ(miopenDestroySolution(solution), miopenStatusSuccess);
}
}

void TearDown() override
Expand Down
5 changes: 5 additions & 0 deletions projects/miopen/test/gtest/gpu_mha_forward.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -294,6 +294,11 @@ class Test_Fwd_Mha : public testing::TestWithParam<TestCase>

VerifyResults(handle);
}

for(auto& solution : solutions)
{
ASSERT_EQ(miopenDestroySolution(solution), miopenStatusSuccess);
}
}

virtual void VerifyResults(const Handle& handle)
Expand Down
46 changes: 28 additions & 18 deletions projects/miopen/test/gtest/gru_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,7 @@
#include "tensor_holder.hpp"

#include "compare_helper.hpp"
#include "gtest_desc_guard.hpp"
#include "../dropout_util.hpp"
#include "../rnn_util.hpp"
#include "../workspace.hpp"
Expand Down Expand Up @@ -90,7 +91,8 @@ void GRUFwdCPUVerify(const miopen::Handle& handle,
std::vector<rocrand_state_xorwow> dropout_states_host;
std::vector<unsigned char> dropout_reservespace_host;
std::vector<T> dropout_hid_state;
miopenTensorDescriptor_t dropout_inputTensor{}, dropout_outputTensor{};
TensorDescGuard dropout_inputTensor;
TensorDescGuard dropout_outputTensor;
if(use_dropout)
{
size_t states_size = dropoutDesc.stateSizeInBytes / sizeof(rocrand_state_xorwow);
Expand All @@ -100,8 +102,6 @@ void GRUFwdCPUVerify(const miopen::Handle& handle,
std::array<int, 2> drop_in_len = {{batch_n, hy_h * bi}};
std::array<int, 2> drop_in_str = {{hy_stride, 1}};
std::array<int, 2> drop_out_str = {{hy_h * bi, 1}};
miopenCreateTensorDescriptor(&dropout_inputTensor);
miopenCreateTensorDescriptor(&dropout_outputTensor);
miopenSetTensorDescriptor(
dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data());
miopenSetTensorDescriptor(
Expand Down Expand Up @@ -200,9 +200,9 @@ void GRUFwdCPUVerify(const miopen::Handle& handle,

DropoutForwardVerify<T>(handle,
dropoutDesc,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
rsvspace,
miopen::deref(dropout_outputTensor),
miopen::deref(dropout_outputTensor.get()),
dropout_hid_state,
dropout_reservespace_host,
dropout_states_tmp,
Expand Down Expand Up @@ -854,13 +854,12 @@ void GRUBwdDataCPUVerify(bool use_dropout,
}

// initial dropoput
miopenTensorDescriptor_t dropout_inputTensor{};
TensorDescGuard dropout_inputTensor;
std::vector<unsigned char> dropout_reservespace_host;
if(use_dropout)
{
std::array<int, 2> drop_in_len = {{batch_n, hy_h * bi}};
std::array<int, 2> drop_in_str = {{hy_stride, 1}};
miopenCreateTensorDescriptor(&dropout_inputTensor);
miopenSetTensorDescriptor(
dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data());

Expand Down Expand Up @@ -921,9 +920,9 @@ void GRUBwdDataCPUVerify(bool use_dropout,
if(use_dropout)
{
DropoutBackwardVerify<T>(dropoutDesc,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
wkspace,
miopen::deref(dropout_inputTensor),
miopen::deref(dropout_inputTensor.get()),
wkspace,
dropout_reservespace_host,
hid_shift + bi * 3 * hy_h,
Expand Down Expand Up @@ -2944,25 +2943,27 @@ class gru_test : public testing::TestWithParam<TestParam>

batch_n = std::accumulate(param.batchSeq.begin(), param.batchSeq.end(), 0);

miopenRNNDescriptor_t rnnDesc;
miopenCreateRNNDescriptor(&rnnDesc);
RNNDescGuard rnnDesc;
miopenRNNAlgo_t algoMode = miopenRNNdefault;

miopenDropoutDescriptor_t DropoutDesc;
miopenCreateDropoutDescriptor(&DropoutDesc);
DropoutDescGuard DropoutDesc;

HandleGuard mio_handle;
void* dropout_state_buf = nullptr;

// See DestroyInternalRnnDropoutDesc — frees the descriptor allocated
// by miopenCreateRNNDescriptor that the upcoming Set* will leak.
DestroyInternalRnnDropoutDesc(rnnDesc);

if(param.useDropout)
{
miopenHandle_t mio_handle;
miopenCreateWithStream(&mio_handle, handle.GetStream());
mio_handle.create(handle.GetStream());

float dropout_rate = 0.5;
unsigned long long dropout_seed = 0ULL;
miopenDropoutGetStatesSize(mio_handle, &statesSizeInBytes);

void* dropout_state_buf;
[[maybe_unused]] auto err =
hipMalloc(static_cast<void**>(&dropout_state_buf), statesSizeInBytes);
(void)hipMalloc(static_cast<void**>(&dropout_state_buf), statesSizeInBytes);

miopenSetDropoutDescriptor(DropoutDesc,
mio_handle,
Expand Down Expand Up @@ -3167,6 +3168,15 @@ class gru_test : public testing::TestWithParam<TestParam>
// seqLength, numLayers,
// biasMode, dirMode,
// inputMode, inVecReal});

// Free the DropoutDescriptor that miopenSetRNNDescriptor just allocated.
// In the dropout path, the internal pointer aliases the user-owned
// DropoutDescGuard — freeing it would double-free.
if(!param.useDropout)
DestroyInternalRnnDropoutDesc(rnnDesc);

if(param.useDropout)
(void)hipFree(dropout_state_buf);
}
};

Expand Down
Loading
Loading