Rnn Relu refactoring #2541

kikimych · 2023-11-20T18:25:40Z

Extracted Relu|Tanh activation functions related code from RNNForwardTrainingPackedTensors to RNNForwardTrainingTanhRelu, from RNNBackwardDataPackedTensors to RNNBackwardDataPackedTensorsRelu

into rnn_training_forward_refactor

…kimych/MIOpen into rnn_training_forward_refactor

Пожалуйста, введите сообщение коммита для ваших изменений. Строки,

junliume · 2023-12-08T20:24:22Z

@shurale-nkn @JehandadKhan please review

shurale-nkn

Partial review

src/ocl/rnnocl.cpp

shurale-nkn · 2023-12-12T19:48:09Z

@junliume please add @atamazov to reviewers

atamazov · 2023-12-12T20:09:28Z

@kikimych Is it time to review? I see this is still a draft.

src/ocl/rnnocl.cpp

shurale-nkn · 2023-12-18T21:45:01Z

src/ocl/rnnocl.cpp

-        const int m = direction == rnn_direction::Forward ? fbatches.at(time)
-                      : time == 0                         ? rbatches.at(0)
-                                                          : rbatches.at(time - 1);
+        const int cur_time  = direction == rnn_direction::Forward ? time : seq_len - 1 - time;


@kikimych
I remember you said that you would make a structure for direct and reverse access to elements.

Readability
auto batch_id_abs = time == seq_len - 1 ? bacc_per_time[directed_time.cur_time()] : bacc_per_time[directed_time.cur_time()] + batches.at(directed_time.next_time());
vs
auto batch_id_abs = time == seq_len - 1 ? bacc_per_time[cur_time] : bacc_per_time[cur_time] + batches.at(next_time);
for example

The question was about all code where bacc_per_time or batches is used.
why not class with API : bacc_per_time.at(time, direction);

auto batch_id_abs = time == seq_len - 1 ? bacc_per_time[cur_time] : bacc_per_time[cur_time] + batches.at(next_time)
vs
auto batch_id_abs = time == seq_len - 1 ? bacc_per_time(time, direction) : bacc_per_time(time,direction) + batches.at(time+1, direction)

so? what conclusion should we get based on this single line?
how this line would look was quite obvious.

After 2 weeks both cases are not obvious at all

src/ocl/rnnocl.cpp

shurale-nkn · 2023-12-18T22:07:18Z

src/ocl/rnnocl.cpp


-            RNNTensorPaddingConverter::ConvertTensorData(
-                handle, dyDesc[0], in_n, dy, packedDYIn, true);
+    auto propagate_dhx = [*this, seqLen, &propagate_dhx_prev, &dhx](int layer) {


*this
copy-capture for current class? really?

&dhx
reference to ptr?

src/ocl/rnnocl.cpp

shurale-nkn · 2023-12-20T18:44:20Z

src/ocl/rnnocl.cpp

+                                  reserveSpace,
+                                  &activDesc,
+                                  propagate_hidden_output,
+                                  propagate_hidden_prev](int layer, int time, int direction) {


shurale-nkn · 2023-12-20T18:53:21Z

src/ocl/rnnocl.cpp

-        const int m = direction == rnn_direction::Forward ? fbatches.at(time)
-                      : time == 0                         ? rbatches.at(0)
-                                                          : rbatches.at(time - 1);
+        const int cur_time  = direction == rnn_direction::Forward ? time : seq_len - 1 - time;


The question was about all code where bacc_per_time or batches is used.
why not class with API : bacc_per_time.at(time, direction);

src/ocl/rnnocl.cpp

shurale-nkn · 2023-12-20T19:01:38Z

src/ocl/rnnocl.cpp

+                                  reserveSpace,
+                                  &activDesc,
+                                  propagate_hidden_output,
+                                  propagate_hidden_prev](int layer, int time, int direction) {


so there is no copy of lambda members? ❔

shurale-nkn

format failed

shurale-nkn · 2023-12-25T22:27:57Z

src/ocl/rnnocl.cpp

@@ -2528,6 +3104,34 @@ void RNNDescriptor::RNNForwardTrainingPackedTensors(
        }
        return;
    }
+
+    if((rnnMode == miopenRNNRELU || rnnMode == miopenRNNTANH) && !use_dropout &&
+       inputMode != miopenRNNskip && !(miopen::IsDisabled(ENV(MIOPEN_RNNFWD_exp))))


why not miopenRNNskip?

shurale-nkn · 2023-12-25T22:31:06Z

src/include/miopen/rnn_util.hpp

+    int first_layer_offset() const { return (in_vec_sz + h_vec_sz) * weight_stride; }
+};
+
+struct RNNOffsets


shurale-nkn · 2023-12-25T22:38:37Z

src/ocl/rnnocl.cpp

-        const int m = direction == rnn_direction::Forward ? fbatches.at(time)
-                      : time == 0                         ? rbatches.at(0)
-                                                          : rbatches.at(time - 1);
+        const int cur_time  = direction == rnn_direction::Forward ? time : seq_len - 1 - time;


so? what conclusion should we get based on this single line?
how this line would look was quite obvious.

CAHEK7 · 2024-01-03T17:09:45Z

src/include/miopen/rnn_util.hpp

+enum rnn_direction
+{
+    Forward  = 0,
+    Backward = 1
+};


It's mostly a question for @atamazov - do we use enum class or have plans to switch to enum class everywhere?

Let's prefer enum class. But switching to it everywhere is impossible IIRC.

… Added RnnBatches class

shurale-nkn · 2024-01-12T17:46:57Z

src/include/miopen/rnn_util.hpp

+    Backward = 1
+};
+
+struct RnnBatches


struct X{ int at(int time, RnnDirection direction){ return RnnDirection::Forward ? batches[time] : batches[ (batches.size() - 1) - time]; } ... std::vector<int> batches; }

this is true for any direction
.next(time, direction) == .at(time+1, direction)
.prev(time, direction) == .at(time-1, direction)

next_time() and prev_time() are redundant functions.

explanation

FWD next_time:
cur_time(time, direction) + 1 => (time+1)

BWD next_time:
cur_time(time, direction) - 1 => (batches.size() - time - 1) - 1 => (batches.size() - (time+1) - 1)

BWD prev_time:
cur_time(time, direction) + 1 => (batches.size() - time - 1) + 1 => (batches.size() - (time-1) - 1)

next() and prev() can be used as syntactic sugar if you want, but than remove mix using of +1 and next() in code

shurale-nkn · 2024-02-05T22:02:41Z

CI reports about a merge conflict that cannot be resolved automatically in rnnocl.cpp

kikimych added 18 commits June 14, 2023 18:01

Added RNN_RELU and RNN_TANH forward training refactored functions

13da580

Added RNNForwardTrainingGRU method

c3d8ab4

Merge branch 'develop' of https://github.com/ROCmSoftwarePlatform/MIOpen

263c278

into rnn_training_forward_refactor

Added RNN Backward Data refactor

f247d7d

Added bidirection support for RNNBackwardData

7bb3be5

Unified relu and gru offset interfaces

3c889de

Lstm minor refactoring

cf96713

Lstm minor refactoring

c2bfee2

Reverted back RNNForwardTraining_MS method

cdce2f5

Merge branch 'rnn_training_forward_refactor' of https://github.com/ki…

b114b9b

…kimych/MIOpen into rnn_training_forward_refactor

Rnn relu bidirectional support draft

12e4a43

Refactored Relu forward

d2a166b

RNNBackwardDataPackedTensorsRelu refactor

f0d1d0a

Пожалуйста, введите сообщение коммита для ваших изменений. Строки,

Simplified RNNBackwardDataPackedTensorsRelu

e8de935

Relu forward minor refactor

ee524b6

Removed hidden_size from offset helpers

5cc2e09

Enabled RNNForwardTrainingTanhRelu by default

7ca8c32

Merge branch 'develop' into rnn_training_forward_refactor

25dc932

kikimych marked this pull request as draft November 20, 2023 18:26

kikimych force-pushed the relu_refactor branch from 8a6beb3 to b230467 Compare November 23, 2023 13:43

RNN Relu and Tanh activation functions refactoring

e9f20d5

kikimych force-pushed the relu_refactor branch from b230467 to e9f20d5 Compare November 24, 2023 06:55

kikimych added 2 commits December 1, 2023 21:13

-no-hx --no-dhy --no-hy --no-dhx modes fix

8ebdcf4

Merge branch 'develop' into relu_refactor

c1ac5af

junliume requested a review from shurale-nkn December 8, 2023 20:24

junliume requested a review from JehandadKhan December 8, 2023 20:24

Merge branch 'develop' into relu_refactor

92cc941

shurale-nkn requested changes Dec 11, 2023

View reviewed changes

kikimych marked this pull request as ready for review December 12, 2023 20:21

kikimych added 3 commits December 18, 2023 18:25

Pull request 2541 review fixes

1745012

Added profiling

d795893

Removed ocl profiling

b93c0d3

shurale-nkn requested changes Dec 18, 2023

View reviewed changes

review fixes

87329c6

shurale-nkn requested changes Dec 20, 2023

View reviewed changes

kikimych added 2 commits December 20, 2023 23:30

review fixes

7ae3116

minor fix

15a891f

shurale-nkn requested changes Dec 25, 2023

View reviewed changes

CAHEK7 reviewed Jan 3, 2024

View reviewed changes

kikimych added 3 commits January 11, 2024 18:29

Review fixes. Made RnnDirection enum class. Removed RnnOffsets class.…

ac66ac0

… Added RnnBatches class

Enabled miopenRNNskip mode in RNNForwardTrainingTanhRelu

6f25e7d

Descriptors renaming

a145ecb

shurale-nkn reviewed Jan 12, 2024

View reviewed changes

RNNBackwardData function renaming

ade5f3a

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Rnn Relu refactoring #2541

Rnn Relu refactoring #2541

kikimych commented Nov 20, 2023

junliume commented Dec 8, 2023

shurale-nkn left a comment

shurale-nkn commented Dec 12, 2023

atamazov commented Dec 12, 2023

shurale-nkn Dec 18, 2023

kikimych Dec 19, 2023

shurale-nkn Dec 20, 2023

kikimych Dec 21, 2023

shurale-nkn Dec 25, 2023

kikimych Jan 9, 2024

shurale-nkn Dec 18, 2023 •

edited

Loading

kikimych Dec 19, 2023

shurale-nkn Dec 20, 2023

shurale-nkn Dec 20, 2023

shurale-nkn Dec 20, 2023

shurale-nkn left a comment

shurale-nkn Dec 25, 2023

shurale-nkn Dec 25, 2023

shurale-nkn Dec 25, 2023

CAHEK7 Jan 3, 2024

atamazov Jan 4, 2024

shurale-nkn Jan 12, 2024

shurale-nkn commented Feb 5, 2024

Rnn Relu refactoring #2541

Are you sure you want to change the base?

Rnn Relu refactoring #2541

Conversation

kikimych commented Nov 20, 2023

junliume commented Dec 8, 2023

shurale-nkn left a comment

Choose a reason for hiding this comment

shurale-nkn commented Dec 12, 2023

atamazov commented Dec 12, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

shurale-nkn Dec 18, 2023 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

shurale-nkn left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

shurale-nkn commented Feb 5, 2024

shurale-nkn Dec 18, 2023 •

edited

Loading