[Enhancement] xdlops NCHW support by transpose #1247

carlushuang · 2021-10-28T12:05:43Z

transpose hip kernel src/kernels/gpu_batched_transpose_kernel/batched_transpose.cpp
internal API to call transpose
ctest for batched transpose test/gpu_nchw_nhwc_transpose.cpp
add NHWC solver support NCHW by using transpose
test ssd

@atamazov @junliume this PR, especially the host part may need careful check. Maybe my implementation is not good and need big code refactoring, so please help review this.

tested on gfx908/gfx90a, in above link

…16_ALT_IMPL to control MIOPEN_DEBUG_FP16_ALT_IMP attribute

… based on attribute MIOPEN_CONVOLUTION_ATTRIB_FP16_ALT_IMPL value

…e attribute. Some error handling. Comments.

…nvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC: Update Solver and Invokers.

…ng. MIOPEN_DEBUG_FP16_ALT_IMP -> MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL.

…attribute accessors.

…tribute from all except ConvolutionAttribute::Set/Get().

…LT attribute via InvokeParams.

…Cm 4.5). Sort list of disabled warnings.

…he sake of clarity of the source code.

…ssues

…small refactor

…Update Solver and Invoker.

…Update Solver and Invokers

…ecks" This reverts commit 5bb733f.

…nd arg

…t_by_transpose

atamazov · 2021-10-28T15:43:17Z

[ci-slip] fix error in ocl backend

Nice typo)) skip

src/include/miopen/magic_div.hpp

atamazov · 2021-11-02T17:40:15Z

src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp


    if(isGfx90aFp16altSupport)
    {
        result.construction_params.push_back(kernel);
        std::ostringstream opts_1(options.str(), std::ios_base::ate);
        GenerateClangDefsym(opts_1, "igemm_bwd_fp16_alt_impl", 1);
        result.construction_params[1].comp_options = opts_1.str();
+        msg << ", fp16_alt:" << ctx.conv_problem.GetConv().attribute.gfx90aFp16alt.GetBwd();


The computations related to msg should happen only if msg is going to be printed.

atamazov · 2021-11-02T17:47:12Z

src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp

    }

-    MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString());
+    MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());


You can use ostringstream syntax in MIOPEN_LOG* macros.

Suggested change

MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());

MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " << config.ToString() << msg.str());

Logging output begins with [GetSolution], which creates an impression that logging happens during GetSolution() call. This is not so, however. The logging happens during invocation of kernels. Let's clearly indicate this:

Suggested change

MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());

MIOPEN_LOG_I2("[INVOKER] ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());

Please always use SolverDbId():

Suggested change

MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());

MIOPEN_LOG_I2(SolverDbId(*this) << ": " ...);

Logging output begins with [GetSolution], which creates an impression that logging happens during GetSolution() call. This is not so, however. The logging happens during invocation of kernels. Let's clearly indicate this:

This IS in GetSolution right? not actually invoker the kernels

This is actually invoker.

Invoker is a lambda produced (instantiated) by another lambda (InvokerFactory). Both can be defined directly in GetSolution (like in this case) or somewhere else (only to avoid duplication of code).

InvokerFactory instance resides in the Solution object. The library runs it only when necessary, i.e. sometime later, after GetSolution(). The resulting Invoker object is instantiated and stored in the Invoker Cache.

Sometime more later, when the library wants to start the kernels, it will run Invoker.

atamazov · 2021-11-02T17:53:25Z

src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp

    }

-    MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString());
+    MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());


Logging output begins with [GetSolution], which creates an impression that logging happens during GetSolution() call. This is not so, however. The logging happens during invocation of kernels. Let's clearly indicate this:

Suggested change

MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());

MIOPEN_LOG_I2("[INVOKER] ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());

atamazov · 2021-11-02T17:54:52Z

src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp

    }

-    MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString());
+    MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());


Please always use SolverDbId():

Suggested change

MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());

MIOPEN_LOG_I2(SolverDbId(*this) << ": " ...);

…t_by_transpose

atamazov

Partial review

src/conv/invokers/impl_gemm_dynamic.cpp

src/include/miopen/datatype.hpp

src/hip/batched_transpose_sol.cpp

atamazov · 2021-11-03T13:35:51Z

src/hip/batched_transpose_sol.cpp

+}
+
+template <typename T>
+static inline float get_normalized_radio(T x, T y)


[Q] what is radio?

it is a radio of for example, width/height, or height/width. I use this radio in side kernel selection logic for some special case (although not optimal, which is hard to achieve)

src/hip/batched_transpose_sol.cpp

atamazov · 2021-11-03T14:02:52Z

src/include/miopen/util_sol.hpp

+#ifndef MIOPEN_UTIL_SOL_HPP_
+#define MIOPEN_UTIL_SOL_HPP_
+
+#include <miopen/batched_transpose_sol.hpp>


What is the "sol" abbreviation stands for? It would be better if all Solutions in MIOpen comply the Solver/Solution/Invoker arch (#866).

sol is for solution. And I use this name because this is different from that in util.hpp, where there exist some method that use a single function call to launch kernel. At lease this transpose implementation, is not a single function call, but return a kernel object and let the caller to launch

atamazov · 2021-11-03T14:10:51Z

src/include/miopen/util_sol.hpp

+
+namespace miopen {
+
+struct TransposeSolution_NCHW2NHWC : public BatchedTransposeSolution


Why we need these wrappers?

CamelCase

Suggested change

struct TransposeSolution_NCHW2NHWC : public BatchedTransposeSolution

struct TransposeSolutionNchw2Nhwc : public BatchedTransposeSolution

Because NCHW->NHWC, and NHWC->NCHW is exactly the same problem of transpose, only by switching the height/width.
By generalizing this problem in to a batched-transpose problem, we can simplify a lot of (maybe) duplicated code, and this transpose can further be extended.
This wrappers are just to make the transpose easier

src/include/miopen/util_sol.hpp

atamazov · 2021-11-03T14:18:10Z

src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp

@@ -30,6 +30,7 @@
 #include <miopen/gcn_asm_utils.hpp>
 #include <miopen/solver/implicitgemm_util.hpp>
 #include <miopen/conv/asm_implicit_gemm.hpp>
+#include <miopen/util_sol.hpp>


This header name looks strange here. If it provides interface to transpose functionality, then it should be expressed in the header name.

Actually I don't have good idea about the naming. I just want to distinguish that in util.hpp

"ask sol" 🤣

src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp

src/kernels/gpu_batched_transpose_kernel/batched_transpose.cpp

Changes made as requested, requesting re-review :)

junliume

CI has passed, review suggestions are applied.

atamazov · 2021-11-04T23:16:22Z

@junliume That was partial review.

junliume · 2021-11-05T02:54:27Z

@junliume That was partial review.

Let's use post merge issues :) Sorry for pressing hard on these blocker PRs.

carlushuang and others added 27 commits October 16, 2021 09:25

implement set/get attribute API, and add MIOPEN_CONVOLUTION_ATTRIB_FP…

4136bbb

…16_ALT_IMPL to control MIOPEN_DEBUG_FP16_ALT_IMP attribute

Merge remote-tracking branch 'origin/develop' into gfx90a_fp16_alt_impl

e1d563e

get attribute in asm igemm nhwc solver, and condintionally set symbol…

0210e86

… based on attribute MIOPEN_CONVOLUTION_ATTRIB_FP16_ALT_IMPL value

Merge branch 'develop' into gfx90a_fp16_alt_impl

533f954

gfx90a_fp16_alt_impl(01) Add constness to the API. Allow resetting th…

0ba350c

…e attribute. Some error handling. Comments.

gfx90a_fp16_alt_impl(02) WrW: Pass ALT attribute via InvokeParams. Co…

1020032

…nvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC: Update Solver and Invokers.

gfx90a_fp16_alt_impl(03) Accelerate access to attribute. Error handli…

35c81c0

…ng. MIOPEN_DEBUG_FP16_ALT_IMP -> MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL.

gfx90a_fp16_alt_impl(04) [Quality] Avoid duplication of code. Inline …

65ee1f6

…attribute accessors.

gfx90a_fp16_alt_impl(05) [fin] Fix build error

c743d63

gfx90a_fp16_alt_impl(06) Fix typo. [Quality] Hide direct access to at…

1846b1a

…tribute from all except ConvolutionAttribute::Set/Get().

gfx90a_fp16_alt_impl(07) [fin] Fix build error

a50e3b1

gfx90a_fp16_alt_impl(08) [TEMP][CI] Disable all but static checks

5bb733f

gfx90a_fp16_alt_impl(09) Remove useless initializers. Fwd/Bwd: Pass A…

7649295

…LT attribute via InvokeParams.

gfx90a_fp16_alt_impl(10) [clang-tidy] Disable altera-unroll-loops (RO…

4043347

…Cm 4.5). Sort list of disabled warnings.

gfx90a_fp16_alt_impl(11) [clang-tidy] Fix some warnings for ROCm 4.5.

65aa044

gfx90a_fp16_alt_impl(12) Disable tidy checks at couple of lines for t…

a32035a

…he sake of clarity of the source code.

gfx90a_fp16_alt_impl(13) Less clarity, but no more cppcheck or tidy i…

3f57ccf

…ssues

gfx90a_fp16_alt_impl(14) ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC: …

7abd5aa

…small refactor

gfx90a_fp16_alt_impl(15) ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC: …

f57f66e

…Update Solver and Invoker.

gfx90a_fp16_alt_impl(16) ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: …

ba44525

…Update Solver and Invokers

Revert "gfx90a_fp16_alt_impl(08) [TEMP][CI] Disable all but static ch…

7d4bbf4

…ecks" This reverts commit 5bb733f.

fix ostringstream constructor with string problem, by adding eta to 2…

a529457

…nd arg

[ci-skip][Quality] openmode is a member of ios_base, not ostringstream.

a376136

add batched transpose gpu kernel aim to serve nchw<->nhwc convert

494fdc4

Merge remote-tracking branch 'origin/develop' into nchw_xdlops_suppor…

fd8c492

…t_by_transpose

[ci-skip] dos2unix

85f910c

[ci-slip] fix error in ocl backend

dfca312

junliume reviewed Oct 28, 2021

View reviewed changes

src/include/miopen/magic_div.hpp Outdated Show resolved Hide resolved

add NCHW support in asm NHWC solver

accb684

carlushuang requested a review from atamazov November 2, 2021 14:14

shaojiewang mentioned this pull request Nov 2, 2021

[NFC] Post review of #1230: Asm igemm nhwc native bf16 kernel on gfx90a #1248

Merged

atamazov suggested changes Nov 2, 2021

View reviewed changes

atamazov reviewed Nov 2, 2021

View reviewed changes

atamazov suggested changes Nov 2, 2021

View reviewed changes

carlushuang added 4 commits November 3, 2021 09:28

[ci-skip] optimize transpose kernel launch parameters

4c3abc7

Merge remote-tracking branch 'origin/develop' into nchw_xdlops_suppor…

d3d234e

…t_by_transpose

fix msg print

3fdffa8

fix hip-tidy about pointer reinterpret_cast

0a9e9f4

atamazov previously requested changes Nov 3, 2021

View reviewed changes

junliume reviewed Nov 3, 2021

View reviewed changes

src/kernels/gpu_batched_transpose_kernel/batched_transpose.cpp Outdated Show resolved Hide resolved

fix several review comments

5c76906

carlushuang requested review from junliume and atamazov November 3, 2021 15:56

carlushuang added 4 commits November 4, 2021 03:57

fix tidy

fa5be89

suppress missing invalidPointerCast

1edf441

optimize transpose for conv 1st layer C=3 case

6cf5f20

suppress clang warning while compiling kernel

2c348f2

junliume added this to the ROCm 5.0 milestone Nov 4, 2021

junliume added complexity_high urgency_blocker TESTING_CI_PASSED labels Nov 4, 2021

junliume approved these changes Nov 4, 2021

View reviewed changes

junliume changed the title ~~xdlops NCHW support by transpose~~ [Enhancement] xdlops NCHW support by transpose Nov 4, 2021

junliume merged commit f02d959 into develop Nov 4, 2021

carlushuang mentioned this pull request Nov 10, 2021

[TESTS] Add/fix tests for batched transpose #1266

Closed

carlushuang deleted the nchw_xdlops_support_by_transpose branch December 6, 2021 13:48

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[Enhancement] xdlops NCHW support by transpose #1247

[Enhancement] xdlops NCHW support by transpose #1247

carlushuang commented Oct 28, 2021 •

edited

Loading

atamazov commented Oct 28, 2021

atamazov Nov 2, 2021

carlushuang Nov 4, 2021

atamazov Nov 2, 2021

atamazov Nov 2, 2021

atamazov Nov 2, 2021

carlushuang Nov 3, 2021

atamazov Nov 3, 2021 •

edited

Loading

atamazov Nov 2, 2021

atamazov Nov 2, 2021

atamazov left a comment

atamazov Nov 3, 2021

carlushuang Nov 3, 2021

atamazov Nov 3, 2021

carlushuang Nov 3, 2021

atamazov Nov 3, 2021

atamazov Nov 3, 2021

carlushuang Nov 3, 2021

atamazov Nov 3, 2021

carlushuang Nov 3, 2021

atamazov Nov 4, 2021

junliume left a comment

atamazov commented Nov 4, 2021

junliume commented Nov 5, 2021

	MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());
	MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " << config.ToString() << msg.str());

	MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());
	MIOPEN_LOG_I2("[INVOKER] ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());

	MIOPEN_LOG_I2("ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC: " + config.ToString() + msg.str());
	MIOPEN_LOG_I2(SolverDbId(*this) << ": " ...);


		namespace miopen {

		struct TransposeSolution_NCHW2NHWC : public BatchedTransposeSolution

	struct TransposeSolution_NCHW2NHWC : public BatchedTransposeSolution
	struct TransposeSolutionNchw2Nhwc : public BatchedTransposeSolution

[Enhancement] xdlops NCHW support by transpose #1247

[Enhancement] xdlops NCHW support by transpose #1247

Conversation

carlushuang commented Oct 28, 2021 • edited Loading

atamazov commented Oct 28, 2021

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

atamazov Nov 3, 2021 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

atamazov 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

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

junliume left a comment

Choose a reason for hiding this comment

atamazov commented Nov 4, 2021

junliume commented Nov 5, 2021

carlushuang commented Oct 28, 2021 •

edited

Loading

atamazov Nov 3, 2021 •

edited

Loading