TensorOps kernels refactoring #3346

novakovicdj · 2024-10-30T14:23:15Z

This is draft PR for refactoring tensor ops kernels to solver structure, so far only Op1dTensorGeneric kernel is switched

src/tensor/problem_description.cpp

src/include/miopen/tensor/solvers.hpp

src/solver/tensor/Op1dTensorGeneric.cpp

src/tensor/problem_description.cpp

CAHEK7 · 2024-11-01T13:18:20Z

src/include/miopen/tensor/problem_description.hpp

+                       const void* alpha0_,
+                       const void* alpha1_,
+                       const void* beta_,


Check this conversation
https://github.com/ROCm/MIOpen/pull/3346/files#r1824480257

Probably alpha0/1 must not be a part of the PD, ideally beta as well, but right now it has to be there..

Would a bool marking if alpha0/... has a "default" value meaning no additional work required suffice?

src/include/miopen/tensor/problem_description.hpp

src/solver/tensor/Op1dTensorGeneric.cpp

src/solver/tensor/Op2dTensorLite.cpp

DrizztDoUrden · 2024-11-01T15:06:05Z

src/include/miopen/tensor/invoke_params.hpp

+    size_t Aoffset;
+    size_t Boffset;
+    size_t Coffset;


Do we need to handle this internally? IIRC it should be possible to externally pass any subtensor via changing pointer+descriptor. If so this is a duplicated functionality

I think that the main point is the pointer is void * and actual type is an miopen_Type_t enum. That's why you can't just add them without special helpers.

src/include/miopen/tensor/invoke_params.hpp

src/include/miopen/tensor/problem_description.hpp

DrizztDoUrden · 2024-11-01T15:08:18Z

src/include/miopen/tensor/problem_description.hpp

+                       const void* alpha0_,
+                       const void* alpha1_,
+                       const void* beta_,


Would a bool marking if alpha0/... has a "default" value meaning no additional work required suffice?

src/include/miopen/tensor/problem_description.hpp

src/solver/tensor/Op1dTensorGeneric.cpp

src/include/miopen/tensor_ops.hpp

…arate some of them into unique solvers and tidy the code

src/solver/tensorOp/tensor_op_helpers.hpp

src/solver/tensorOp/Op2dTensorLite.cpp

src/solver/tensorOp/Op2dTensorSquash.cpp

src/solver/tensorOp/Op4dTensorLite.cpp

src/solver/tensorOp/OpTensorFwdBias.cpp

shurale-nkn · 2024-11-08T13:52:56Z

Please provide a comparison of the average only CPU time (new solver vs old api) measurements for 100 calls with same problem and the costs associated with the first call of the unique problem configuration.

novakovicdj · 2024-11-12T12:07:58Z

Please provide a comparison of the average only CPU time (new solver vs old api) measurements for 100 calls with same problem and the costs associated with the first call of the unique problem configuration.

Here is a comparison of average host time between old and new structure

Kernel		New structure [ms]	Old structure [ms]	diff [ms]
Op1dTensorGeneric	first run	279.3786	291.3806	-12.002
	other 100 runs	0.2908	0.2549	0.0359
Op2dTensorGeneric	first run	281.8186	283.4622	-1.6436
	other 100 runs	0.356	0.2432	0.1128
Op2dTensorLite	first run	634.2228	662.2278	-28.005
	other 100 runs	0.335	0.2308	0.1042
Op2dTensorSquash	first run	668.978	699.9932	-31.0152
	other 100 runs	0.3481	0.2548	0.0933
Op3dTensorGeneric	first run	642.1512	656.3394	-14.1882
	other 100 runs	0.2659	0.2485	0.0174
OpTensorFwdBias	first run	636.6204	654.8222	-18.2018
	other 100 runs	0.3351	0.2321	0.103
OpTensorFwdBiasGeneric	first run	636.4756	662.4915	-26.0159
	other 100 runs	0.3498	0.2434	0.1064
OpTensorLeadingOnes	first run	644.8348	666.8713	-22.0365
	other 100 runs	0.3466	0.2755	0.0711
OpTensorLeadingOnesGeneric	first run	648.6535	669.6379	-20.9844
	other 100 runs	0.3552	0.2569	0.0983
Op4dTensorLite	first run	641.4747	664.4976	-23.0229
	other 100 runs	0.33	0.2206	0.1094
Op4dTensorGeneric	first run	650.7638	670.8961	-20.1323
	other 100 runs	0.3563	0.2456	0.1107
Op5dTensorGeneric	first run	655.6774	685.431	-29.7536
	other 100 runs	0.3745	0.2437	0.1308

New structure is faster on average for 20ms for first runs and it is slower for 0.1ms for other 100 calls or 0.001ms per call

shurale-nkn · 2024-11-12T13:05:25Z

Please provide a comparison of the average only CPU time (new solver vs old api) measurements for 100 calls with same problem and the costs associated with the first call of the unique problem configuration.

Here is a comparison of average host time between old and new structure

Kernel New structure [ms] Old structure [ms] diff [ms]
Op1dTensorGeneric first run 279.3786 291.3806 -12.002
other 100 runs 0.2908 0.2549 0.0359
Op2dTensorGeneric first run 281.8186 283.4622 -1.6436
other 100 runs 0.356 0.2432 0.1128
Op2dTensorLite first run 634.2228 662.2278 -28.005
other 100 runs 0.335 0.2308 0.1042
Op2dTensorSquash first run 668.978 699.9932 -31.0152
other 100 runs 0.3481 0.2548 0.0933
Op3dTensorGeneric first run 642.1512 656.3394 -14.1882
other 100 runs 0.2659 0.2485 0.0174
OpTensorFwdBias first run 636.6204 654.8222 -18.2018
other 100 runs 0.3351 0.2321 0.103
OpTensorFwdBiasGeneric first run 636.4756 662.4915 -26.0159
other 100 runs 0.3498 0.2434 0.1064
OpTensorLeadingOnes first run 644.8348 666.8713 -22.0365
other 100 runs 0.3466 0.2755 0.0711
OpTensorLeadingOnesGeneric first run 648.6535 669.6379 -20.9844
other 100 runs 0.3552 0.2569 0.0983
Op4dTensorLite first run 641.4747 664.4976 -23.0229
other 100 runs 0.33 0.2206 0.1094
Op4dTensorGeneric first run 650.7638 670.8961 -20.1323
other 100 runs 0.3563 0.2456 0.1107
Op5dTensorGeneric first run 655.6774 685.431 -29.7536
other 100 runs 0.3745 0.2437 0.1308
New structure is faster on average for 20ms for first runs and it is slower for 0.1ms for other 100 calls or 0.001ms per call

The results are very strange; we need to obtain the experiment protocol. How was the program executed, and what was used for measurement?
so far, according to the table, each subsequent launch is on average 30% slower

novakovicdj · 2024-11-12T15:43:38Z

Please provide a comparison of the average only CPU time (new solver vs old api) measurements for 100 calls with same problem and the costs associated with the first call of the unique problem configuration.

Here is a comparison of average host time between old and new structure
Kernel New structure [ms] Old structure [ms] diff [ms]
Op1dTensorGeneric first run 279.3786 291.3806 -12.002
other 100 runs 0.2908 0.2549 0.0359
Op2dTensorGeneric first run 281.8186 283.4622 -1.6436
other 100 runs 0.356 0.2432 0.1128
Op2dTensorLite first run 634.2228 662.2278 -28.005
other 100 runs 0.335 0.2308 0.1042
Op2dTensorSquash first run 668.978 699.9932 -31.0152
other 100 runs 0.3481 0.2548 0.0933
Op3dTensorGeneric first run 642.1512 656.3394 -14.1882
other 100 runs 0.2659 0.2485 0.0174
OpTensorFwdBias first run 636.6204 654.8222 -18.2018
other 100 runs 0.3351 0.2321 0.103
OpTensorFwdBiasGeneric first run 636.4756 662.4915 -26.0159
other 100 runs 0.3498 0.2434 0.1064
OpTensorLeadingOnes first run 644.8348 666.8713 -22.0365
other 100 runs 0.3466 0.2755 0.0711
OpTensorLeadingOnesGeneric first run 648.6535 669.6379 -20.9844
other 100 runs 0.3552 0.2569 0.0983
Op4dTensorLite first run 641.4747 664.4976 -23.0229
other 100 runs 0.33 0.2206 0.1094
Op4dTensorGeneric first run 650.7638 670.8961 -20.1323
other 100 runs 0.3563 0.2456 0.1107
Op5dTensorGeneric first run 655.6774 685.431 -29.7536
other 100 runs 0.3745 0.2437 0.1308
New structure is faster on average for 20ms for first runs and it is slower for 0.1ms for other 100 calls or 0.001ms per call

The results are very strange; we need to obtain the experiment protocol. How was the program executed, and what was used for measurement? so far, according to the table, each subsequent launch is on average 30% slower

I have talked with @CAHEK7 and he suggested to remove compilation time from first run time and to enable kernel profiling so here are new results that I have got
note. diff for first run is (new_total_time - new_comp_time) - (old_total_time - old_comp_time)

Kernel	run	New structure [ms] (total)	New structure [ms] (compilation)	Old structure [ms] (total)	Old structure [ms] (compilation)	diff [ms]
Op1dTensorGeneric	First	279.0176	278.813	289.2538	289.0891	0.0399
	Other 100	2.3191		2.5519		-0.2328
Op2dTensorGeneric	First	282.0158	281.8159	285.6689	285.4783	0.0093
	Other 100	2.318		2.3533		-0.0353
Op2dTensorLite	First	645.8057	645.6093	662.8484	662.6658	0.0138
	Other 100	2.3686		2.275		0.0936
Op2dTensorSquash	First	680.7002	680.5031	709.6122	709.4446	0.0295
	Other 100	2.5681		2.2783		0.2898
Op3dTensorGeneric	First	641.8974	641.701	695.5571	695.3875	0.0268
	Other 100	2.471		2.1212		0.3498
OpTensorFwdBias	First	640.256	640.062	680.8951	680.709	0.0079
	Other 100	2.3217		2.4506		-0.1289
OpTensorFwdBiasGeneric	First	671.2975	671.0999	686.9962	686.8118	0.0132
	Other 100	2.5107		2.3875		0.1232
OpTensorLeadingOnes	First	676.697	676.5015	660.4297	660.2519	0.0177
	Other 100	2.5548		2.4725		0.0823
OpTensorLeadingOnesGen	First	656.037	655.834	677.3558	677.18	0.0272
	Other 100	2.7894		2.4013		0.3881
Op4dTensorLite	First	662.9597	662.7599	667.7071	667.5309	0.0236
	Other 100	2.5528		2.2425		0.3103
Op4dTensorGeneric	First	659.8318	659.6297	673.9867	673.8067	0.0221
	Other 100	2.3756		2.3735		0.0021
Op5dTensorGeneric	First	654.421	654.2273	693.2979	693.1147	0.0105
	Other 100	2.2762		2.187		0.0892

This is part of code that I changed in tensor_ops test for this performance testing

` handle.EnableProfiling(true);

    auto start = std::chrono::steady_clock::now();
    miopen::OpTensor2(handle,
                      // miopenTensorOpAdd,
                      // miopenTensorOpMax,
                      // miopenTensorOpMin,
                      miopenTensorOpMul,
                      &alpha0,
                      a.desc,
                      a_dev.get(),
                      &alpha1,
                      b.desc,
                      b_dev.get(),
                      &beta,
                      c.desc,
                      c_dev.get(),
                      Aoffset,
                      Boffset,
                      Coffset,
                      false); // it does not verify non-standard behaviour
    auto end = std::chrono::steady_clock::now();
    auto elapsed_init =
        std::chrono::duration_cast<std::chrono::duration<float, std::milli>>(end - start)
            .count() -
        handle.GetKernelTime();

    printf("First time: %f ms\n", elapsed_init);

    float elapsed_other_full = 0;

    for(int i = 0; i < 100; i++)
    {
        start = std::chrono::steady_clock::now();
        miopen::OpTensor2(handle,
                          // miopenTensorOpAdd,
                          // miopenTensorOpMax,
                          // miopenTensorOpMin,
                          miopenTensorOpMul,
                          &alpha0,
                          a.desc,
                          a_dev.get(),
                          &alpha1,
                          b.desc,
                          b_dev.get(),
                          &beta,
                          c.desc,
                          c_dev.get(),
                          Aoffset,
                          Boffset,
                          Coffset,
                          false); // it does not verify non-standard behaviour
        end = std::chrono::steady_clock::now();
        auto elapsed_other =
            std::chrono::duration_cast<std::chrono::duration<float, std::milli>>(end - start)
                .count() -
            handle.GetKernelTime();
        elapsed_other_full += elapsed_other;
    }

    printf("Other times total: %f ms\nOther times average: %f ms\n",
           elapsed_other_full,
           elapsed_other_full / 100);`

I was switching between OpTensor and OpTensor2 to run tests for old and new structure, also only test case for one kernel was running each time.
I will run more tests and try to figure out why host times for some kernels are slower than for the others.

src/solver/tensorOp/Op5dTensorGeneric.cpp

src/tensorOp/problem_description.cpp

src/include/miopen/tensorOp/problem_description.hpp

src/include/miopen/tensorOp/invoke_params.hpp

…tentially boost performance on host side

novakovicdj · 2024-11-18T13:27:06Z

I did some profiling to compare old and new structure and saw that creation of network_config is slower than before, this is more visible for bigger dimension tensors and it is consequence of the new format of network_config. Network_config creation for 5d tensors is around 4 times slower than in the old structure and around 3 times slower compared to 1d tensor network_config in the new structure. Because of all of that I switched to using string and got speed up of around 2.2 times compared to using stream.

After that I run 500 iterations of old and new structure for all tensor kernels and got the result that the new version is faster for 0.0005ms on average, which is around 20% faster than old structure.

src/kernels/MIOpenTensorKernels.cl

BrianHarrisonAMD · 2024-11-21T22:57:03Z

I am not seeing a lot of testing coverage for OpTensor before the changes.
The only tests I can see that specifically run OpTensor seem to be these, and it looks like they aren't covering all ops.

Would it be possible to add new tests to the gtest suite to ensure correctness for the new solvers being added?

randyspauldingamd · 2024-11-22T04:01:27Z

src/CMakeLists.txt

+    solver/tensorOp/Op2dTensorLite.cpp
+    solver/tensorOp/Op2dTensorSquash.cpp
+    solver/tensorOp/Op3dTensorGeneric.cpp
+    solver/tensorOp/OpTensorFwdBias.cpp


Could you alphabetize these?

src/include/miopen/rnn/solvers.hpp

src/solver/tensorOp/Op2dTensorGeneric.cpp

randyspauldingamd · 2024-11-22T05:27:37Z

src/solver/tensorOp/tensor_op_helpers.hpp

+    if(is4dLite)
+    {
+        // for naive tensor ops
+        const std::string data_type = GetDataType(bTensorDesc.GetType());


Effectively unused

randyspauldingamd · 2024-11-22T05:27:49Z

src/solver/tensorOp/tensor_op_helpers.hpp

+
+        size_t TENS_LEN = cTensorDesc.GetElementSize();
+        size_t RD_BLCK  = (TENS_LEN % 4 == 0) ? 4 : (TENS_LEN % 2 == 0) ? 2 : 1;
+        const std::string READ_TYPE =


Unused--is something missing?

novakovicdj · 2024-11-22T17:19:31Z

I am not seeing a lot of testing coverage for OpTensor before the changes. The only tests I can see that specifically run OpTensor seem to be these, and it looks like they aren't covering all ops.

Would it be possible to add new tests to the gtest suite to ensure correctness for the new solvers being added?

Current test for tensorOp is covering all solvers except for Op2dTensorSquash but I did some changes and tested it locally and it worked fine. There is a plan to switch this test to gtest and then those improvements of testing tensorOps will be implemented.

As a part of this PR I will add some unit tests for Problem Descriptor, so please do not merge this yet

novakovicdj added 3 commits October 30, 2024 16:17

initial changes and support for 1d generic kernel

a2972c0

1d solver file name change

75cecb2

solver name change in cmakelists.txt

035989c

Vsevolod1983 reviewed Oct 31, 2024

View reviewed changes

src/tensor/problem_description.cpp Outdated Show resolved Hide resolved

more changes, 2d generic and 2d lite kernel

cf91070

CAHEK7 requested changes Nov 1, 2024

View reviewed changes

DrizztDoUrden requested changes Nov 1, 2024

View reviewed changes

novakovicdj added 4 commits November 1, 2024 18:32

some changes suggested in the comments

f2a11d6

additional changes

ac13ff3

initial switch to solver structure for all kernels, still need to sep…

cadb264

…arate some of them into unique solvers and tidy the code

fix for two kernels in one solver

63603f0

CAHEK7 requested changes Nov 6, 2024

View reviewed changes

CAHEK7 requested review from DrizztDoUrden and Vsevolod1983 November 6, 2024 16:46

CAHEK7 added quality complexity_middle labels Nov 6, 2024

novakovicdj added 3 commits November 7, 2024 15:29

additional changes

976bd84

clang format

6be98d0

fwd_conv_bias changed

d6ffea5

novakovicdj marked this pull request as ready for review November 7, 2024 15:19

novakovicdj requested review from JehandadKhan and junliume as code owners November 7, 2024 15:19

tidy some part of the code

89dd24c

novakovicdj requested review from BrianHarrisonAMD and bpepers-me as code owners November 8, 2024 07:51

novakovicdj added 2 commits November 8, 2024 09:52

Merge branch 'develop' into tensor_refactoring

9ba8810

fix typos

5a9b5ed

CAHEK7 reviewed Nov 12, 2024

View reviewed changes

src/solver/tensorOp/Op5dTensorGeneric.cpp Outdated Show resolved Hide resolved

CAHEK7 reviewed Nov 14, 2024

View reviewed changes

src/tensorOp/problem_description.cpp Outdated Show resolved Hide resolved

CAHEK7 reviewed Nov 15, 2024

View reviewed changes

src/include/miopen/tensorOp/problem_description.hpp Outdated Show resolved Hide resolved

CAHEK7 reviewed Nov 15, 2024

View reviewed changes

src/include/miopen/tensorOp/problem_description.hpp Outdated Show resolved Hide resolved

CAHEK7 reviewed Nov 15, 2024

View reviewed changes

src/include/miopen/tensorOp/invoke_params.hpp Outdated Show resolved Hide resolved

implementnting suggestions, updating network_config and changes to po…

c9f310a

…tentially boost performance on host side

novakovicdj requested a review from BradPepersAMD as a code owner November 15, 2024 17:37

clang format

496b414

CAHEK7 self-requested a review November 15, 2024 17:43

CAHEK7 mentioned this pull request Nov 15, 2024

Remove our own implementation of std::index_sequence #3389

Open

novakovicdj added 2 commits November 18, 2024 09:49

Merge branch 'develop' into tensor_refactoring

22c6c48

change for new Op3dTensorGeneric kernel usage

cb6fd6e

novakovicdj added 6 commits November 18, 2024 15:34

remove unused variable

6c3d0c2

clang format

bd0bd61

support for half data type for CL kernels

3f14d3a

additional changes for support for half type

042129e

initial removal of tensorocl.cpp

371d43c

code tidying

155b35f

CAHEK7 added the TESTING_CI_PASSED label Nov 21, 2024

CAHEK7 approved these changes Nov 21, 2024

View reviewed changes

BrianHarrisonAMD reviewed Nov 21, 2024

View reviewed changes

src/kernels/MIOpenTensorKernels.cl Show resolved Hide resolved

randyspauldingamd requested changes Nov 22, 2024

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

TensorOps kernels refactoring #3346

TensorOps kernels refactoring #3346

novakovicdj commented Oct 30, 2024

CAHEK7 Nov 1, 2024

DrizztDoUrden Nov 1, 2024

DrizztDoUrden Nov 1, 2024

CAHEK7 Nov 15, 2024

DrizztDoUrden Nov 1, 2024

shurale-nkn commented Nov 8, 2024

novakovicdj commented Nov 12, 2024

shurale-nkn commented Nov 12, 2024

novakovicdj commented Nov 12, 2024 •

edited

Loading

novakovicdj commented Nov 18, 2024

BrianHarrisonAMD commented Nov 21, 2024

randyspauldingamd Nov 22, 2024

randyspauldingamd Nov 22, 2024

randyspauldingamd Nov 22, 2024

novakovicdj commented Nov 22, 2024

TensorOps kernels refactoring #3346

Are you sure you want to change the base?

TensorOps kernels refactoring #3346

Conversation

novakovicdj commented Oct 30, 2024

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 Nov 8, 2024

novakovicdj commented Nov 12, 2024

shurale-nkn commented Nov 12, 2024

novakovicdj commented Nov 12, 2024 • edited Loading

novakovicdj commented Nov 18, 2024

BrianHarrisonAMD commented Nov 21, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

novakovicdj commented Nov 22, 2024

novakovicdj commented Nov 12, 2024 •

edited

Loading