carlushuang c12fb0a624 [CK_TILE][HOTFIX] WA for address space by disable it completely (#1947) ...

* port all moe changes from ck_moe_gemm branch

* refine codes in the pr

* fix tail odd

* fix clang format

* fix clang format2

* make hot loop scheduler compatible with 16x16 and 32x32

* clang format

* fix per token quant

* rename moe example

* clang format

* WA for address space by disable it completely

* hot fix moe gemm2

---------

Co-authored-by: coderfeli <coderfeli@163.com>
Co-authored-by: feli <felix.li@amd.com>

2025-03-06 12:01:25 +08:00

..

batchnorm_multiblock

[HotFix] add config and version files to pass on build info (#856)

2023-08-23 11:36:17 -07:00

gemm_layernorm

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

normalization

layernorm and groupnorm backward data (#1083)

2023-12-19 04:23:11 +08:00

block_to_ctile_map.hpp

Rebase the PR #1520 to ROCm repo. (#1574)

2025-02-20 18:58:14 -08:00

gridwise_2d_multiple_reduction_multiblock.hpp

Fix for beta!=0 in reduce (#1440)

2024-08-06 09:10:39 -07:00

gridwise_2d_multiple_reduction_threadwise.hpp

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

gridwise_2d_reduction_multiblock.hpp

Fix for beta!=0 in reduce (#1440)

2024-08-06 09:10:39 -07:00

gridwise_2d_reduction_threadwise_multi_d.hpp

Universal gemm splitk using reduce (with multi-d) (#1341)

2024-07-19 22:01:22 +08:00

gridwise_2d_reduction_threadwise.hpp

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

gridwise_batched_gemm_gemm_xdl_cshuffle_v1.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle_v1.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_batched_gemm_multiple_d_softmax_gemm_xdl_cshuffle_v1.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_batched_gemm_softmax_gemm_wmma_cshuffle.hpp

Merging the gfx12 code into public repo. (#1362)

2024-06-27 00:33:34 -07:00

gridwise_batched_gemm_softmax_gemm_xdl_cshuffle_v1.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_batchnorm_backward_blockwise_welford.hpp

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

gridwise_batchnorm_forward_blockwise_welford.hpp

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

gridwise_elementwise_1d_scale.hpp

Refactor elementwise kernels (#1222)

2024-04-19 13:31:17 +02:00

gridwise_elementwise_2d.hpp

Add support for NGCHW in grouped conv bwd wei (#1491)

2024-09-03 10:52:03 +02:00

gridwise_elementwise_layernorm_welford_variance.hpp

enable compilation of INSTANCES_ONLY for Windows (#1082)

2023-12-20 14:34:53 -08:00

gridwise_fpAintB_gemm_wmma.hpp

Merging the gfx12 code into public repo. (#1362)

2024-06-27 00:33:34 -07:00

gridwise_gemm_bias_add_reduce_xdl_cshuffle_v1.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_dl_multiple_d.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

2024-08-06 10:06:10 +02:00

gridwise_gemm_dl_v1r3.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

2024-08-06 10:06:10 +02:00

gridwise_gemm_dpp.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_multiple_abd_xdl_cshuffle.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_multiple_d_wmma_cshuffle.hpp

Statically Cast Pointer Offset (#1631)

2024-11-05 09:59:08 -08:00

gridwise_gemm_multiple_d_xdl_cshuffle_lds_direct_load.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_multiple_d_xdl_cshuffle.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_multiple_d_xdl_splitk_cshuffle.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_pipeline_selector.hpp

Rebase the PR #1520 to ROCm repo. (#1574)

2025-02-20 18:58:14 -08:00

gridwise_gemm_pipeline_v1.hpp

Navi3 rel (#1176)

2024-03-08 17:11:51 -08:00

gridwise_gemm_pipeline_v2.hpp

Fixed Gemmv2r3 kpad (#938)

2023-09-26 18:40:00 -05:00

gridwise_gemm_pipeline_v3.hpp

initial stream-k implementation with example (#699)

2023-07-26 14:18:15 -05:00

gridwise_gemm_pipeline_v4_direct_load.hpp

Add support for double buffering in direct load GEMM kernel (#1052)

2023-12-03 23:08:47 +01:00

gridwise_gemm_reduce_xdl_cshuffle_v1.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_split_k_multiple_d_xdl_cshuffle_v2.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_split_k_multiple_d_xdl_cshuffle.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_waveletmodel.hpp

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

gridwise_gemm_wmma.hpp

Merging the gfx12 code into public repo. (#1362)

2024-06-27 00:33:34 -07:00

gridwise_gemm_xdl_cshuffle_bwd_weight_v3.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdl_cshuffle_streamk_v3.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdl_cshuffle_v1.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdl_cshuffle_v2.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdl_cshuffle_v3_b_scale.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdl_cshuffle_v3_multi_abd.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdl_cshuffle_v3_multi_d_ab_scale.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp

[A8W8 GEMM] Optimized weight-preshuffled implementation & add quantization datatype for CK TILE rms_norm (#1862)

2025-02-20 14:00:27 -08:00

gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdl_cshuffle_v3.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdl_layernorm_cshuffle_v1.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdl_waveletmodel_cshuffle.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdlops_bwd_weight.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdlops_skip_b_lds_v1.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_splitk_lds_direct_load.hpp

aggregate device macros in ck_tile config header (#1297)

2024-05-20 08:34:45 -07:00

gridwise_gemm_xdlops_streamk.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_v2r3.hpp

Change flag to CK_GFX90A_DENORM_WORKAROUND (#1817)

2025-01-28 09:58:39 -05:00

gridwise_gemm_xdlops_v2r4.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_v2r4r2.hpp

replace the ENV macro with CK_ENV (#1296)

2024-05-17 10:42:51 -07:00

gridwise_gemm_xdlops_v3r1.hpp

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gridwise_gemm_xdlops_v3r2.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_v3r3.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_moe_gemm.hpp

[CK_TILE][HOTFIX] WA for address space by disable it completely (#1947)

2025-03-06 12:01:25 +08:00

gridwise_permute.hpp

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

gridwise_put_element_1d.hpp

Average pool backward deviceOP and example (#797)

2023-08-10 12:04:35 +08:00

gridwise_set_buffer_value.hpp

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

gridwise_set_multiple_buffer_value.hpp

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

gridwise_softmax.hpp

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

gridwise_sparse_embeddings_forward_layernorm_builtins.hpp

Replace buffer load/store intrinsics with builtins (#1876)

2025-03-05 14:33:28 -08:00

gridwise_sparse_embeddings_forward_layernorm.hpp

[HotFix] add config and version files to pass on build info (#856)

2023-08-23 11:36:17 -07:00

gridwise_tensor_rearrange.hpp

Statically Cast Pointer Offset (#1631)

2024-11-05 09:59:08 -08:00