[igemm_dynamic] v4r1 bwd dynamic kernel (#272)

* [dynamic-kernel] add v4r1 generic dynamic kernel and solver, fwd fp32 * update tunable table * fix tidy for -abseil-string-find-startswith * fix tidy for readability-simplify-boolean-expr * add code of v4r1 dynamic fwd kc1x1 case * runnable code for v4r1 igemm 1x1 asm kernel case * modify igemm dynamic kernel call func: if kc1x1 kernel, remove the xy kernel args * add test_conv_for_dynamic_implicit_gemm to test dynamic kernel feature * register invoker for igemm_dynamic solver * fix hip-clang bug to run assembly kernel * put asm file in folder kernels/dynamic_igemm * add v4r1 bwd [skip ci] * add missing header include * split invoker into seperate conv direction * use conv_problem as invoker param, instead of conv ctx * remove kernel name check in invoker Co-authored-by: root <[email protected]> Co-authored-by: shaojiewang <[email protected]> Co-authored-by: Daniel Lowell <[email protected]>
ROCm · Jul 27, 2020 · dce9c70 · dce9c70
1 parent c3bbbc2
commit dce9c70
Show file tree

Hide file tree

Showing 11 changed files with 5,354 additions and 92 deletions.
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
@@ -173,6 +173,7 @@ set( MIOpen_Source
     include/miopen/rnn_util.hpp
     include/miopen/bz2.hpp
     include/miopen/comgr.hpp
+    include/miopen/numeric.hpp
     md_graph.cpp
     mdg_expr.cpp
     conv/invokers/gcn_asm_1x1u.cpp
@@ -223,6 +224,7 @@ set( MIOpen_Source
     solver/conv_asm_implicit_gemm_v4r1_dynamic.cpp
     solver/conv_hip_implicit_gemm_v4r4_gen_xdlops_fwd_fp32.cpp
     solver/conv_hip_implicit_gemm_bwd_v1r1_xdlops_nchw_kcyx_nkhw.cpp
+    solver/conv_asm_implicit_gemm_bwd_v4r1_dynamic.cpp
     )
 
 list(APPEND MIOpen_Source tmp_dir.cpp binary_cache.cpp md5.cpp)