RNN Dynamic algo single stream (#3406)

* FWD dynamic lstm * BWD dynamic lstm * Back weights dynamic lstm * test update * file separation * add new type in miopenRNNAlgo_t --------- Co-authored-by: Alex Eremin <[email protected]>
ROCm · Jan 7, 2025 · 370e378 · 370e378
1 parent fc5b012
commit 370e378
Show file tree

Hide file tree

Showing 18 changed files with 2,306 additions and 1,087 deletions.
diff --git a/include/miopen/miopen.h b/include/miopen/miopen.h
@@ -3748,10 +3748,12 @@ typedef enum
  */
 typedef enum
 {
-    miopenRNNdefault = 0, /*!< Use dedicated gate-operation kernel for LSTM and fundamental
-                             algorithm for vanilla RNN & GRU */
-    miopenRNNfundamental =
-        1, /*!< Function by basic tesnsor operations, supported for vanilla RNN, LSTM, GRU */
+    miopenRNNdefault = 0,        /*!< Use dedicated gate-operation kernel for LSTM and fundamental
+                                    algorithm for vanilla RNN & GRU */
+    miopenRNNfundamental = 1,    /*!< Deprecated, low performance. Function by basic tesnsor
+                                    operations, supported for vanilla RNN, LSTM, GRU */
+    miopenRNNroundedDynamic = 2, /*!< The algorithm rounds some RNN parametrs upwards
+                                    to utilize the most optimal GEMM kernel in the computation.*/
 } miopenRNNAlgo_t;
 
 /*! @enum miopenRNNDirectionMode_t

diff --git a/src/include/miopen/rnn.hpp b/src/include/miopen/rnn.hpp
@@ -155,6 +155,13 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor
                                miopenRNNFWDMode_t fwdMode) const;
     size_t GetMaxReserveSize(Handle& handle, const SeqTensorDescriptor& xDesc) const;
 
+    std::tuple<size_t, size_t> GetTmpSpaceSizeDynamicAlgo(Handle& handle,
+                                                          const SeqTensorDescriptor& xDesc,
+                                                          miopenRNNFWDMode_t fwdMode) const;
+    bool CheckDynamicAlgoSelection(Handle& handle,
+                                   const SeqTensorDescriptor& xDesc,
+                                   miopenRNNFWDMode_t fwdMode) const;
+
     size_t
     GetParamsSize(Handle& handle, const TensorDescriptor& xDesc, miopenDataType_t dtype) const;
     size_t GetParamsSize(size_t inputVector) const;
@@ -534,6 +541,8 @@ struct MIOPEN_INTERNALS_EXPORT RNNDescriptor : miopenRNNDescriptor
                                          Data_t hy,
                                          const TensorDescriptor& cyDesc,
                                          Data_t cy,
+                                         Data_t workSpace,
+                                         size_t workSpaceSize,
                                          Data_t reserveSpace,
                                          size_t reserveSpaceSize) const;