Haocong WANG 789c15d703 [Navi3x] Add Device Operations (#567) ...

* wmma_op + unit test

* add arch limitation to wmma test

* change arch limitation

* Refactor + Add all type unit test(int4 compile failed)

* Add f32_16x16x16_bf16 unit test

* tempsave

* tempsave

* tempsave

* runtime bug, cannot find symbol

* workaround for incorrect HIP warpSize return value

* debugging

* tempsave

* Correctness OK, waiting for optimization

* Tidy up + format

* temp save

* temp save, reproduce the v_bfi_b32 issue

* add inline asm for wmmaop test

* tidy up

* clean some debug purpose code

* discard some codes

* clang format

* clang format

* compiler issue fixed + increase tile size

* navi3x_multipleD+example

* temp save

* workable

* batchedgemm[OK], groupconv[debug]

* groupconv: Sanity check[OK], Performance[Bad]

* navi3x_groupconv_need_optimization

* format

* Add arch limitation to all wmma examples

* fix bug: example30 input conv args

[ROCm/composable_kernel commit: 0cfda84d05]

2023-02-15 11:50:51 -06:00

..

batchnorm_multiblock

BatchNorm backward instance/external API/profiler/tests (#519)

2022-11-30 13:32:20 -06:00

gemm_layernorm

Gemm layernorm welford (#413)

2023-01-16 20:08:25 -06:00

block_to_ctile_map.hpp

Wavelet (inter-wave consumer-producer) GEMM (#310)

2023-01-18 12:00:02 -06:00

gridwise_2d_multiple_reduction_multiblock.hpp

Batchnorm-forward and Batchnorm-infer Implemented using generic kernels (#320)

2022-08-15 10:11:02 -05:00

gridwise_2d_multiple_reduction_threadwise.hpp

Batchnorm-forward and Batchnorm-infer Implemented using generic kernels (#320)

2022-08-15 10:11:02 -05:00

gridwise_2d_reduction_multiblock.hpp

add license in file (#303)

2022-06-24 23:32:43 -05:00

gridwise_2d_reduction_threadwise.hpp

add license in file (#303)

2022-06-24 23:32:43 -05:00

gridwise_batched_gemm_gemm_xdl_cshuffle_v1.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle_v1.hpp

batched_gemm + multiple_d + gemm + multiple_d (#394)

2022-09-14 17:54:18 -05:00

gridwise_batched_gemm_multiple_d_softmax_gemm_xdl_cshuffle_v1.hpp

Remove the workaround for bf16 attention tests. (#586)

2023-02-14 18:06:24 -06:00

gridwise_batched_gemm_softmax_gemm_xdl_cshuffle_v1.hpp

Remove the workaround for bf16 attention tests. (#586)

2023-02-14 18:06:24 -06:00

gridwise_batchnorm_backward_blockwise_welford.hpp

BatchNorm backward instance/external API/profiler/tests (#519)

2022-11-30 13:32:20 -06:00

gridwise_batchnorm_forward_blockwise_welford.hpp

BatchNorm backward implementation (#461)

2022-11-28 20:51:10 -06:00

gridwise_contraction_dlops_v1r2.hpp

add license in file (#303)

2022-06-24 23:32:43 -05:00

gridwise_elementwise_1d.hpp

Add 'Permute' device op & example (#408)

2022-09-19 21:30:25 -05:00

gridwise_elementwise_2d.hpp

Gridwise elementwise 2d (#466)

2022-12-12 09:18:10 -06:00

gridwise_elementwise_layernorm_welford_variance.hpp

Gemm layernorm welford (#413)

2023-01-16 20:08:25 -06:00

gridwise_gemm_bias_add_reduce_xdl_cshuffle_v1.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_dl_multiple_d.hpp

Add multiple d gridwise gemm on Navi21 for ResNet50 (#517)

2022-12-02 11:42:31 -06:00

gridwise_gemm_dl_v1r3.hpp

Add Conv Forward on Navi21 for ResNet50 (#490)

2022-10-31 10:24:25 -06:00

gridwise_gemm_dlops_v1r2.hpp

add license in file (#303)

2022-06-24 23:32:43 -05:00

gridwise_gemm_dlops_v2.hpp

add license in file (#303)

2022-06-24 23:32:43 -05:00

gridwise_gemm_dlops_v3.hpp

add license in file (#303)

2022-06-24 23:32:43 -05:00

gridwise_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_multiple_d_wmma_cshuffle.hpp

[Navi3x] Add Device Operations (#567)

2023-02-15 11:50:51 -06:00

gridwise_gemm_multiple_d_xdl_cshuffle.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_pipeline_selector.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_pipeline_v1.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_pipeline_v2.hpp

GEMM pipeline v2 (#317)

2022-07-08 15:55:14 -05:00

gridwise_gemm_reduce_xdl_cshuffle_v1.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_split_k_multiple_d_xdl_cshuffle.hpp

Example contraction splitk (#430)

2022-10-11 17:54:34 -05:00

gridwise_gemm_waveletmodel.hpp

Wavelet (inter-wave consumer-producer) GEMM (#310)

2023-01-18 12:00:02 -06:00

gridwise_gemm_wmma.hpp

[Navi3x-LWPCK-545] Block-wise GEMM + Real GEMM_WMMA_FP16 (#541)

2023-01-16 20:06:01 -06:00

gridwise_gemm_xdl_cshuffle_v1.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_xdl_layernorm_cshuffle_v1.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_xdl_waveletmodel_cshuffle.hpp

Wavelet (inter-wave consumer-producer) GEMM (#310)

2023-01-18 12:00:02 -06:00

gridwise_gemm_xdlops_bwd_weight.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_xdlops_skip_b_lds_v1.hpp

Fused GEMM+GEMM (#351)

2022-08-13 09:18:58 -05:00

gridwise_gemm_xdlops_v2r3.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_xdlops_v2r4.hpp

add license in file (#303)

2022-06-24 23:32:43 -05:00

gridwise_gemm_xdlops_v2r4r2.hpp

Add examples of batched/grouped/SplitK Gemm for int8/bfp16/fp16/fp32 (#361)

2022-08-23 14:41:56 -05:00

gridwise_gemm_xdlops_v3r1.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_xdlops_v3r2.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_gemm_xdlops_v3r3.hpp

Add pipeline v1/v2 selector, add more instances (#381)

2022-11-02 16:50:48 -06:00

gridwise_normalization_naive_variance.hpp

Refine layernorm naming and test code (#497)

2022-11-02 16:57:28 -06:00

gridwise_normalization_welford_variance.hpp

Gemm layernorm welford (#413)

2023-01-16 20:08:25 -06:00

gridwise_permute.hpp

Add 'Permute' device op & example (#408)

2022-09-19 21:30:25 -05:00

gridwise_set_buffer_value.hpp

Batchnorm-forward implemented using welford method to calculate variance (#403)

2022-10-27 18:52:54 -06:00

gridwise_set_multiple_buffer_value.hpp

Batchnorm-forward and Batchnorm-infer Implemented using generic kernels (#320)

2022-08-15 10:11:02 -05:00

gridwise_softmax.hpp

fix standalone softmax race condition around blockwise reduction (#323)

2022-07-14 22:52:45 -05:00

gridwise_sparse_embeddings_forward_layernorm.hpp

remove unused variable (#564)

2023-01-31 10:34:35 +08:00