diff --git a/projects/miopen/test/cpu_rnn.hpp b/projects/miopen/test/cpu_rnn.hpp index b5845335b973..0139a431ad84 100644 --- a/projects/miopen/test/cpu_rnn.hpp +++ b/projects/miopen/test/cpu_rnn.hpp @@ -30,6 +30,7 @@ * LSTM CPU verification functions **********************************************/ #include "gemm.hpp" +#include "gtest/gtest_desc_guard.hpp" template void LSTMFwdCPUVerify(const miopen::Handle& handle, @@ -92,7 +93,8 @@ void LSTMFwdCPUVerify(const miopen::Handle& handle, std::vector dropout_states_host; std::vector dropout_reservespace_host; std::vector 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); @@ -102,8 +104,6 @@ void LSTMFwdCPUVerify(const miopen::Handle& handle, std::array drop_in_len = {{batch_n_cpu, hy_h * bi}}; std::array drop_in_str = {{hy_stride, 1}}; std::array 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( @@ -204,9 +204,9 @@ void LSTMFwdCPUVerify(const miopen::Handle& handle, DropoutForwardVerify(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, @@ -710,13 +710,12 @@ void LSTMBwdDataCPUVerify(bool use_dropout_cpu, } // initial dropoput - miopenTensorDescriptor_t dropout_inputTensor{}; + TensorDescGuard dropout_inputTensor; std::vector dropout_reservespace_host; if(use_dropout_cpu) { std::array drop_in_len = {{batch_n_cpu, hy_h * bi}}; std::array drop_in_str = {{hy_stride, 1}}; - miopenCreateTensorDescriptor(&dropout_inputTensor); miopenSetTensorDescriptor( dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data()); @@ -781,9 +780,9 @@ void LSTMBwdDataCPUVerify(bool use_dropout_cpu, if(use_dropout_cpu) { DropoutBackwardVerify(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, @@ -1601,7 +1600,8 @@ void RNNFwdTrainCPUVerify(const miopen::Handle& handle, std::vector dropout_states_host; std::vector dropout_reservespace_host; std::vector 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); @@ -1611,8 +1611,6 @@ void RNNFwdTrainCPUVerify(const miopen::Handle& handle, std::array drop_in_len = {{batch_n, hy_h * bi}}; std::array drop_in_str = {{hy_stride, 1}}; std::array 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( @@ -1708,9 +1706,9 @@ void RNNFwdTrainCPUVerify(const miopen::Handle& handle, DropoutForwardVerify(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, @@ -2070,13 +2068,12 @@ void RNNBwdDataCPUVerify(bool use_dropout, } // initial dropoput - miopenTensorDescriptor_t dropout_inputTensor{}; + TensorDescGuard dropout_inputTensor; std::vector dropout_reservespace_host; if(use_dropout) { std::array drop_in_len = {{batch_n, hy_h * bi}}; std::array drop_in_str = {{hy_stride, 1}}; - miopenCreateTensorDescriptor(&dropout_inputTensor); miopenSetTensorDescriptor( dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data()); @@ -2140,9 +2137,9 @@ void RNNBwdDataCPUVerify(bool use_dropout, if(use_dropout) { DropoutBackwardVerify(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, @@ -2748,7 +2745,8 @@ void GRUFwdCPUVerify(const miopen::Handle& handle, std::vector dropout_states_host; std::vector dropout_reservespace_host; std::vector 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); @@ -2758,8 +2756,6 @@ void GRUFwdCPUVerify(const miopen::Handle& handle, std::array drop_in_len = {{batch_n, hy_h * bi}}; std::array drop_in_str = {{hy_stride, 1}}; std::array 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( @@ -2858,9 +2854,9 @@ void GRUFwdCPUVerify(const miopen::Handle& handle, DropoutForwardVerify(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, @@ -3517,13 +3513,12 @@ void GRUBwdDataCPUVerify(bool use_dropout, } // initial dropoput - miopenTensorDescriptor_t dropout_inputTensor{}; + TensorDescGuard dropout_inputTensor; std::vector dropout_reservespace_host; if(use_dropout) { std::array drop_in_len = {{batch_n, hy_h * bi}}; std::array drop_in_str = {{hy_stride, 1}}; - miopenCreateTensorDescriptor(&dropout_inputTensor); miopenSetTensorDescriptor( dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data()); @@ -3589,9 +3584,9 @@ void GRUBwdDataCPUVerify(bool use_dropout, if(use_dropout) { DropoutBackwardVerify(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, diff --git a/projects/miopen/test/gtest/CMakeLists.txt b/projects/miopen/test/gtest/CMakeLists.txt index 3d68de070512..af74113fa312 100644 --- a/projects/miopen/test/gtest/CMakeLists.txt +++ b/projects/miopen/test/gtest/CMakeLists.txt @@ -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) diff --git a/projects/miopen/test/gtest/gpu_mha_backward.cpp b/projects/miopen/test/gtest/gpu_mha_backward.cpp index 0229d84b1caf..c79e4aac73f3 100644 --- a/projects/miopen/test/gtest/gpu_mha_backward.cpp +++ b/projects/miopen/test/gtest/gpu_mha_backward.cpp @@ -376,6 +376,11 @@ class Test_Bwd_Mha : public testing::TestWithParam checkOutput(miopenTensorMhaDK, "tensor dK", dKDesc_ref); checkOutput(miopenTensorMhaDV, "tensor dV", dVDesc_ref); } + + for(auto& solution : solutions) + { + ASSERT_EQ(miopenDestroySolution(solution), miopenStatusSuccess); + } } void TearDown() override diff --git a/projects/miopen/test/gtest/gpu_mha_forward.cpp b/projects/miopen/test/gtest/gpu_mha_forward.cpp index 89a12313e6e7..cb0cf785d306 100644 --- a/projects/miopen/test/gtest/gpu_mha_forward.cpp +++ b/projects/miopen/test/gtest/gpu_mha_forward.cpp @@ -294,6 +294,11 @@ class Test_Fwd_Mha : public testing::TestWithParam VerifyResults(handle); } + + for(auto& solution : solutions) + { + ASSERT_EQ(miopenDestroySolution(solution), miopenStatusSuccess); + } } virtual void VerifyResults(const Handle& handle) diff --git a/projects/miopen/test/gtest/gru_test.cpp b/projects/miopen/test/gtest/gru_test.cpp index 7594689f2eea..22ee89eeb675 100644 --- a/projects/miopen/test/gtest/gru_test.cpp +++ b/projects/miopen/test/gtest/gru_test.cpp @@ -7,6 +7,8 @@ #include "tensor_holder.hpp" #include "compare_helper.hpp" +#include "gtest_desc_guard.hpp" +#include "gtest_handle_guard.hpp" #include "../dropout_util.hpp" #include "../rnn_util.hpp" #include "../workspace.hpp" @@ -90,7 +92,8 @@ void GRUFwdCPUVerify(const miopen::Handle& handle, std::vector dropout_states_host; std::vector dropout_reservespace_host; std::vector 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); @@ -100,8 +103,6 @@ void GRUFwdCPUVerify(const miopen::Handle& handle, std::array drop_in_len = {{batch_n, hy_h * bi}}; std::array drop_in_str = {{hy_stride, 1}}; std::array 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( @@ -200,9 +201,9 @@ void GRUFwdCPUVerify(const miopen::Handle& handle, DropoutForwardVerify(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, @@ -854,13 +855,12 @@ void GRUBwdDataCPUVerify(bool use_dropout, } // initial dropoput - miopenTensorDescriptor_t dropout_inputTensor{}; + TensorDescGuard dropout_inputTensor; std::vector dropout_reservespace_host; if(use_dropout) { std::array drop_in_len = {{batch_n, hy_h * bi}}; std::array drop_in_str = {{hy_stride, 1}}; - miopenCreateTensorDescriptor(&dropout_inputTensor); miopenSetTensorDescriptor( dropout_inputTensor, miopenFloat, 2, drop_in_len.data(), drop_in_str.data()); @@ -921,9 +921,9 @@ void GRUBwdDataCPUVerify(bool use_dropout, if(use_dropout) { DropoutBackwardVerify(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, @@ -2952,25 +2952,27 @@ class gru_test : public testing::TestWithParam 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(&dropout_state_buf), statesSizeInBytes); + (void)hipMalloc(static_cast(&dropout_state_buf), statesSizeInBytes); miopenSetDropoutDescriptor(DropoutDesc, mio_handle, @@ -3175,6 +3177,15 @@ class gru_test : public testing::TestWithParam // 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); } }; diff --git a/projects/miopen/test/gtest/gtest_desc_guard.hpp b/projects/miopen/test/gtest/gtest_desc_guard.hpp index 8de1b9436ce6..06eabee13890 100644 --- a/projects/miopen/test/gtest/gtest_desc_guard.hpp +++ b/projects/miopen/test/gtest/gtest_desc_guard.hpp @@ -52,3 +52,20 @@ using DropoutDescGuard = DescGuard; + +// Frees the internal DropoutDescriptor allocated by miopen::RNNDescriptor's +// default and 8-arg constructors. The library has no destructor for this +// field, so any test that creates an RNN descriptor and then calls +// miopenSetRNNDescriptor[_V2] on it must invoke this helper: +// 1. Before each Set* call (frees the default-allocated internal one). +// 2. After the run, only when no user-supplied DropoutDescriptor was passed +// via the _V2 path (in that path the internal pointer aliases the +// user-owned descriptor — freeing it would double-free). +inline void DestroyInternalRnnDropoutDesc(miopenRNNDescriptor_t rnnDesc) +{ + miopenDropoutDescriptor_t dropDesc = nullptr; + miopenGetRNNDescriptor_V2( + rnnDesc, nullptr, nullptr, &dropDesc, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr); + if(dropDesc != nullptr) + miopenDestroyDropoutDescriptor(dropDesc); +} diff --git a/projects/miopen/test/gtest/gtest_handle_guard.hpp b/projects/miopen/test/gtest/gtest_handle_guard.hpp new file mode 100644 index 000000000000..65c9d8e41b43 --- /dev/null +++ b/projects/miopen/test/gtest/gtest_handle_guard.hpp @@ -0,0 +1,44 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT +#pragma once + +#include + +#include + +// RAII wrapper for miopenHandle_t. Cannot use the DescGuard template because +// miopenCreateWithStream takes an extra hipStream_t argument. The handle is +// default-constructed empty; call create(stream) to allocate it lazily. +class HandleGuard +{ +public: + HandleGuard() = default; + + explicit HandleGuard(hipStream_t stream) { status = miopenCreateWithStream(&handle, stream); } + + ~HandleGuard() + { + if(handle != nullptr) + miopenDestroy(handle); + } + + void create(hipStream_t stream) + { + if(handle != nullptr) + miopenDestroy(handle); + status = miopenCreateWithStream(&handle, stream); + } + + operator miopenHandle_t() { return handle; } + miopenHandle_t get() { return handle; } + miopenStatus_t getStatus() const { return status; } + + HandleGuard(const HandleGuard&) = delete; + HandleGuard& operator=(const HandleGuard&) = delete; + HandleGuard(HandleGuard&&) = delete; + HandleGuard& operator=(HandleGuard&&) = delete; + +private: + miopenHandle_t handle = nullptr; + miopenStatus_t status = miopenStatusSuccess; +}; diff --git a/projects/miopen/test/gtest/lstm.hpp b/projects/miopen/test/gtest/lstm.hpp index 76a93a45ba98..7bfb3824ed38 100644 --- a/projects/miopen/test/gtest/lstm.hpp +++ b/projects/miopen/test/gtest/lstm.hpp @@ -30,6 +30,8 @@ #include "cpu_rnn.hpp" #include "workspace.hpp" #include "verify.hpp" +#include "gtest_desc_guard.hpp" +#include "gtest_handle_guard.hpp" #include #include @@ -1707,22 +1709,25 @@ struct LSTM_test : Verifier #endif auto&& handle = get_handle(); - miopenRNNDescriptor_t rnnDesc; - miopenCreateRNNDescriptor(&rnnDesc); - miopenDropoutDescriptor_t DropoutDesc; - miopenCreateDropoutDescriptor(&DropoutDesc); + RNNDescGuard rnnDesc; + DropoutDescGuard DropoutDesc; size_t statesSizeInBytes = 0; + 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(useDropout != 0) { - miopenHandle_t mio_handle; - miopenCreateWithStream(&mio_handle, handle.GetStream()); + mio_handle.create(handle.GetStream()); float dropout_rate{0.5f}; unsigned long long dropout_seed{0ULL}; miopenDropoutGetStatesSize(mio_handle, &statesSizeInBytes); - void* dropout_state_buf; (void)hipMalloc(static_cast(&dropout_state_buf), statesSizeInBytes); miopenSetDropoutDescriptor(DropoutDesc, @@ -1967,5 +1972,14 @@ struct LSTM_test : Verifier rnnDesc, input, dyin, hx, rsvgpu, rsvcpu, workSpaceBwdData, batchSeq, hiddenSize, wei_sz, batch_n, seqLength, numLayers, biasMode, dirMode, inputMode, inVecReal, hx_sz, nohx, bool(useDropout), usePadding}); + + // 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(useDropout == 0) + DestroyInternalRnnDropoutDesc(rnnDesc); + + if(useDropout != 0) + (void)hipFree(dropout_state_buf); } }; diff --git a/projects/miopen/test/gtest/mha_find20.cpp b/projects/miopen/test/gtest/mha_find20.cpp index f8e5e3c82fe4..01290df99b75 100644 --- a/projects/miopen/test/gtest/mha_find20.cpp +++ b/projects/miopen/test/gtest/mha_find20.cpp @@ -666,6 +666,12 @@ TEST(GPU_TestMhaFind20_FP32, MhaForward) test.TestSolutionAttributes(solutions); test.TestRunSolutions(handle, solutions); + + for(auto& solution : solutions) + { + EXPECT_EQUAL(miopenDestroySolution(solution), miopenStatusSuccess); + } + test.Finalize(); } @@ -679,5 +685,11 @@ TEST(GPU_TestMhaFind20_FP32, MhaBackward) test.TestSolutionAttributes(solutions); test.TestRunSolutions(handle, solutions); + + for(auto& solution : solutions) + { + EXPECT_EQUAL(miopenDestroySolution(solution), miopenStatusSuccess); + } + test.Finalize(); } diff --git a/projects/miopen/test/gtest/rnn_seq_api.hpp b/projects/miopen/test/gtest/rnn_seq_api.hpp index b96a7537c3fd..4f7cded37687 100644 --- a/projects/miopen/test/gtest/rnn_seq_api.hpp +++ b/projects/miopen/test/gtest/rnn_seq_api.hpp @@ -27,6 +27,7 @@ #include "gtest_common.hpp" #include "compare_helper.hpp" +#include "gtest_desc_guard.hpp" #include "../dropout_util.hpp" #include "get_handle.hpp" @@ -1758,6 +1759,11 @@ struct RNNSeqApiCommon : public ::testing::TestWithParam miopen::DropoutDescriptor dropoutDesc{}; size_t statesSizeInBytes = 0; + void* dropout_state_buf = nullptr; + + // See DestroyInternalRnnDropoutDesc — frees the descriptor allocated + // by RNNDescriptor's default constructor that the upcoming Set* will leak. + DestroyInternalRnnDropoutDesc(&rnnDesc); if(useDropout != 0) { @@ -1765,7 +1771,6 @@ struct RNNSeqApiCommon : public ::testing::TestWithParam unsigned long long dropout_seed = 0ULL; miopenDropoutGetStatesSize(&handle, &statesSizeInBytes); - void* dropout_state_buf; (void)hipMalloc(static_cast(&dropout_state_buf), statesSizeInBytes); miopenSetDropoutDescriptor(&dropoutDesc, @@ -1880,6 +1885,17 @@ struct RNNSeqApiCommon : public ::testing::TestWithParam << std::endl; inference_dropout_issue_notified = true; } + + // Free the DropoutDescriptor that miopenSetRNNDescriptor just allocated. + // In the dropout path, the internal pointer aliases the user-owned stack + // `dropoutDesc` — freeing it would double-free. + if(useDropout == 0) + DestroyInternalRnnDropoutDesc(&rnnDesc); + + if(useDropout != 0) + { + (void)hipFree(dropout_state_buf); + } } public: diff --git a/projects/miopen/test/gtest/rnn_vanilla_common.hpp b/projects/miopen/test/gtest/rnn_vanilla_common.hpp index cf0b3a7cef84..125cfa0a772e 100644 --- a/projects/miopen/test/gtest/rnn_vanilla_common.hpp +++ b/projects/miopen/test/gtest/rnn_vanilla_common.hpp @@ -1211,6 +1211,10 @@ struct RNNVanillaCommon : ::testing::TestWithParam DropoutDescGuard DropoutDesc; size_t statesSizeInBytes = 0; + // See DestroyInternalRnnDropoutDesc — frees the descriptor allocated + // by miopenCreateRNNDescriptor that the upcoming Set* will leak. + DestroyInternalRnnDropoutDesc(rnnDesc); + miopenRNNAlgo_t algoMode = miopenRNNdefault; miopenHandle_t mio_handle = nullptr; #if MIOPEN_BACKEND_HIP @@ -1486,6 +1490,12 @@ struct RNNVanillaCommon : ::testing::TestWithParam // biasMode, dirMode, // inputMode, rnnMode, 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(useDropout == 0) + DestroyInternalRnnDropoutDesc(rnnDesc); + if(useDropout != 0) { #if MIOPEN_BACKEND_HIP diff --git a/projects/miopen/test/gtest/softmax_find20.cpp b/projects/miopen/test/gtest/softmax_find20.cpp index c3f4857c38c8..094a432d4521 100644 --- a/projects/miopen/test/gtest/softmax_find20.cpp +++ b/projects/miopen/test/gtest/softmax_find20.cpp @@ -280,7 +280,13 @@ class SoftmaxFind20Test std::cerr << "Finished testing solution functions." << std::endl; } - void Finalize() { EXPECT_EQUAL(miopenDestroyProblem(problem), miopenStatusSuccess); } + void Finalize(std::vector& solutions) + { + for(auto& solution : solutions) + EXPECT_EQUAL(miopenDestroySolution(solution), miopenStatusSuccess); + + EXPECT_EQUAL(miopenDestroyProblem(problem), miopenStatusSuccess); + } private: void Initialize() @@ -364,7 +370,8 @@ TEST(GPU_SoftmaxFind20_FP32, softmaxForward) test.TestSolutionAttributes(solutions); test.TestRunSolutionsForward(handle, solutions); - test.Finalize(); + + test.Finalize(solutions); } TEST(GPU_SoftmaxFind20_FP32, softmaxBackward_fp32) @@ -377,7 +384,8 @@ TEST(GPU_SoftmaxFind20_FP32, softmaxBackward_fp32) test.TestSolutionAttributes(solutions); test.TestRunSolutionsBackward(handle, solutions); - test.Finalize(); + + test.Finalize(solutions); } TEST(GPU_SoftmaxFind20_FP16, softmaxBackward_log_instance_mode_fp16) @@ -390,7 +398,8 @@ TEST(GPU_SoftmaxFind20_FP16, softmaxBackward_log_instance_mode_fp16) test.TestSolutionAttributes(solutions); test.TestRunSolutionsBackward(handle, solutions); - test.Finalize(); + + test.Finalize(solutions); } TEST(GPU_SoftmaxFind20_FP16, softmaxBackward_log_channel_mode_fp16) @@ -403,7 +412,8 @@ TEST(GPU_SoftmaxFind20_FP16, softmaxBackward_log_channel_mode_fp16) test.TestSolutionAttributes(solutions); test.TestRunSolutionsBackward(handle, solutions); - test.Finalize(); + + test.Finalize(solutions); } TEST(GPU_SoftmaxFind20_BFP16, softmaxBackward_log_instance_mode_bfp16) @@ -416,7 +426,8 @@ TEST(GPU_SoftmaxFind20_BFP16, softmaxBackward_log_instance_mode_bfp16) test.TestSolutionAttributes(solutions); test.TestRunSolutionsBackward(handle, solutions); - test.Finalize(); + + test.Finalize(solutions); } TEST(GPU_SoftmaxFind20_BFP16, softmaxBackward_log_channel_mode_bfp16) @@ -429,5 +440,6 @@ TEST(GPU_SoftmaxFind20_BFP16, softmaxBackward_log_channel_mode_bfp16) test.TestSolutionAttributes(solutions); test.TestRunSolutionsBackward(handle, solutions); - test.Finalize(); + + test.Finalize(solutions); } diff --git a/projects/miopen/test/gtest/w_supertensor.cpp b/projects/miopen/test/gtest/w_supertensor.cpp index 475e7dd52162..a0a38f4c69f0 100644 --- a/projects/miopen/test/gtest/w_supertensor.cpp +++ b/projects/miopen/test/gtest/w_supertensor.cpp @@ -31,6 +31,7 @@ #include #include "gtest_common.hpp" +#include "gtest_desc_guard.hpp" #include "test_parameter_name_generator.hpp" #include "verify.hpp" @@ -40,7 +41,6 @@ using TestCase = std::tuple, NamedParameter, NamedParameter, NamedParameter, - NamedParameter, NamedParameter, NamedParameter, NamedParameter, @@ -411,8 +411,6 @@ struct verify_w_tensor_set EXPECT_EQUAL(status, miopenStatusSuccess); - const auto param_dev_out = handle.Create(paramSize); - status = miopenGetRNNLayerParam(&handle, rnnDesc, layer, @@ -452,8 +450,6 @@ struct verify_w_tensor_set EXPECT_EQUAL(status, miopenStatusSuccess); - const auto bias_dev_out = handle.Create(biasSize); - status = miopenGetRNNLayerBias(&handle, rnnDesc, layer, @@ -493,8 +489,24 @@ struct verify_w_tensor_set inline auto GenCases() { +#if defined(__SANITIZE_ADDRESS__) || (defined(__has_feature) && __has_feature(address_sanitizer)) + // When Address Sanitizer is enabled, the process exceeds the allowable limit for virtual + // memory areas. Another way around this is to set vm.max_map_count to something large in + // the environment (eg. vm.max_map_count=1048576). However it was determined that while + // ASan is enabled, losing a small amount of edge case test coverage was acceptable. + return testing::Combine( + MakeNamedParameterValues("batch_size", 2, 4), + MakeNamedParameterValues("num_layer", 2, 4), + MakeNamedParameterValues("in_size", 4, 8), + MakeNamedParameterValues("wei_hh", 4, 8), + MakeNamedParameterValues("mode", miopenRNNRELU, miopenLSTM, miopenGRU), + MakeNamedParameterValues( + "biasMode", miopenRNNwithBias, miopenRNNNoBias), + MakeNamedParameterValues( + "directionMode", miopenRNNunidirection, miopenRNNbidirection), + MakeNamedParameterValues("inMode", miopenRNNskip, miopenRNNlinear)); +#else return testing::Combine( - MakeNamedParameterValues("seqLen", 1, 2, 4), MakeNamedParameterValues("batch_size", 2, 4, 8), MakeNamedParameterValues("num_layer", 4, 8, 16), MakeNamedParameterValues("in_size", 2, 8, 16), @@ -505,6 +517,7 @@ inline auto GenCases() MakeNamedParameterValues( "directionMode", miopenRNNunidirection, miopenRNNbidirection), MakeNamedParameterValues("inMode", miopenRNNskip, miopenRNNlinear)); +#endif } inline auto GetCases() @@ -580,26 +593,18 @@ struct WSuperTensorTest : public testing::TestWithParam void SetUp() override { prng::reset_seed(); - std::tie( - seqLen, batch_size, num_layer, in_size, wei_hh, mode, biasMode, directionMode, inMode) = + std::tie(batch_size, num_layer, in_size, wei_hh, mode, biasMode, directionMode, inMode) = GetParam(); - auto status = miopenCreateRNNDescriptor(&rnnDesc); - EXPECT_EQ(status, miopenStatusSuccess); - - status = miopenCreateTensorDescriptor(&inputTensor); - EXPECT_EQ(status, miopenStatusSuccess); - - status = miopenCreateTensorDescriptor(&weightTensor); - EXPECT_EQ(status, miopenStatusSuccess); - - status = miopenCreateTensorDescriptor(¶mTensor); - EXPECT_EQ(status, miopenStatusSuccess); - - status = miopenCreateTensorDescriptor(&biasTensor); - EXPECT_EQ(status, miopenStatusSuccess); + ASSERT_EQ(rnnDesc.getStatus(), miopenStatusSuccess); + ASSERT_EQ(inputTensor.getStatus(), miopenStatusSuccess); + ASSERT_EQ(weightTensor.getStatus(), miopenStatusSuccess); + ASSERT_EQ(paramTensor.getStatus(), miopenStatusSuccess); + ASSERT_EQ(biasTensor.getStatus(), miopenStatusSuccess); } + void TearDown() override { DestroyDropoutDesc(); } + void Run() { if(inMode == miopenRNNskip && in_size != wei_hh) @@ -609,6 +614,10 @@ struct WSuperTensorTest : public testing::TestWithParam const std::array in_lens{{batch_size, in_size}}; + // miopenSetRNNDescriptor overwrites the descriptor via copy assignment, + // leaking the internal DropoutDescriptor allocated by miopenCreateRNNDescriptor. + DestroyDropoutDesc(); + auto status = miopenSetRNNDescriptor( rnnDesc, wei_hh, num_layer, inMode, directionMode, mode, biasMode, algo, dataType); EXPECT_EQ(status, miopenStatusSuccess); @@ -623,6 +632,25 @@ struct WSuperTensorTest : public testing::TestWithParam Verify(); } + void DestroyDropoutDesc() + { + miopenDropoutDescriptor_t dropDesc = nullptr; + miopenGetRNNDescriptor_V2(rnnDesc, + nullptr, + nullptr, + &dropDesc, + nullptr, + nullptr, + nullptr, + nullptr, + nullptr, + nullptr); + if(dropDesc != nullptr) + { + miopenDestroyDropoutDescriptor(dropDesc); + } + } + private: template void Verify() @@ -651,7 +679,6 @@ struct WSuperTensorTest : public testing::TestWithParam } private: - int seqLen{}; int batch_size{}; int num_layer{}; int in_size{}; @@ -660,21 +687,20 @@ struct WSuperTensorTest : public testing::TestWithParam miopenRNNBiasMode_t biasMode{}; miopenRNNDirectionMode_t directionMode{}; miopenRNNInputMode_t inMode{}; - miopenRNNDescriptor_t rnnDesc{}; + RNNDescGuard rnnDesc; miopenRNNAlgo_t algo{miopenRNNdefault}; miopenDataType_t dataType{miopenFloat}; - miopenTensorDescriptor_t inputTensor{}; - miopenTensorDescriptor_t weightTensor{}; - miopenTensorDescriptor_t paramTensor{}; - miopenTensorDescriptor_t biasTensor{}; + TensorDescGuard inputTensor; + TensorDescGuard weightTensor; + TensorDescGuard paramTensor; + TensorDescGuard biasTensor; }; struct TestNameGenerator { std::string operator()(const auto& info) { - const auto& [seqLen, - batch_size, + const auto& [batch_size, num_layer, in_size, wei_hh, @@ -685,10 +711,9 @@ struct TestNameGenerator std::stringstream ss; std::string str; - ss << "seqLen_" << seqLen() << "_batch_size_" << batch_size() << "_num_layer_" - << num_layer() << "_in_size_" << in_size() << "_wei_hh_" << wei_hh() << "_mode_" - << mode() << "_directionMode_" << directionMode() << "_inMode_" << inMode() - << "_test_id_" << info.index; + ss << "batch_size_" << batch_size() << "_num_layer_" << num_layer() << "_in_size_" + << in_size() << "_wei_hh_" << wei_hh() << "_mode_" << mode() << "_directionMode_" + << directionMode() << "_inMode_" << inMode() << "_test_id_" << info.index; str = ss.str();