ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-15 10:37:44 +00:00
Code Issues Packages Projects Releases Wiki Activity
Files
441a7d9847410dd0323cd482d8b017f3cf85bd08
composable_kernel/include/ck/utility
History
lalala-sh 8f426b1216 Moe gemm activation (#2026) 
* fix useless code and remove usless oob

* clang format

* fix coredump in e2e test

* fix2

* fix clang format

* fix output oob

* impl int64 but result not correct

* int64 index ok now

* input output all ok

* fix uint32

* revert v1 test

* use uint32

* mork to support 13w tokens

* moe sorting fix moebuf

* fix merge

* update moe api fix aiter build

* fix buid

* fuse silu

* silu ok

* acale ok

* add silu

* change code

* gemm2 ok

* gufusion compatible ok, fix warnings

* gu fusion for m32 m64 ok

* support bf16 cshuffle

* i4 gemm2 ok

* i4 gemm2 ok and i4 gemm1 build

* 16x16 run ok

* change flops; change cshuffle dtype

* fuse gelu silu act in moe gemm1

* fp8 with act ready

* int4 act ready

* remove useless changes

* remove useless code change

* fix clang format

* add the arch limit of int4 moe gemm

* fuse moe activation

* fix fp8 16x16

* fix no quant case

* fix bugs

* fix fp8 gufusion bug

* remove useless comments

* refine activation code & complete moe example

* fix int8 bugs

* merge tkw1

---------

Co-authored-by: coderfeli <coderfeli@163.com>
Co-authored-by: feli <felix.li@amd.com>
Co-authored-by: illsilin <Illia.Silin@amd.com>
Co-authored-by: root <root@hjbog-srdc-51.amd.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

[ROCm/composable_kernel commit: 39ba03f25d]
2025-04-23 10:35:34 +08:00
..
amd_address_space.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
amd_buffer_addressing_builtins.hpp
Replace buffer load/store intrinsics with builtins (#1876)
2025-03-05 14:33:28 -08:00
amd_buffer_addressing.hpp
Rebase the PR #1520 to ROCm repo. (#1574)
2025-02-20 18:58:14 -08:00
amd_ck_fp8.hpp
fix codegen issues (#2052)
2025-04-07 07:08:39 -07:00
amd_gemm_dpp.hpp
Redesign the DPP8 GEMM kernel to use warp-wise component (#863)
2023-09-06 11:44:09 -05:00
amd_inline_asm.hpp
Split up data_type header. (#1996)
2025-03-24 15:08:54 -07:00
amd_lds.hpp
Implement direct loads split-K GEMM kernel (#1137)
2024-02-07 01:08:34 +01:00
amd_smfmac.hpp
Added structural sparsity blockwise gemm (#1435)
2024-09-11 15:19:42 +02:00
amd_wave_read_first_lane.hpp
Rebase the PR #1520 to ROCm repo. (#1574)
2025-02-20 18:58:14 -08:00
amd_wmma.hpp
Add bf16 and int8 wmma gemms for Navi3x and Navi4x. (#1671)
2024-11-18 14:07:04 -08:00
amd_xdlops.hpp
MX GEMM - Add MX BF8 example (#2071)
2025-04-16 15:25:02 -06:00
array_multi_index.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
array.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
blkgemmpipe_scheduler.hpp
Fix on the error (#1956)
2025-03-07 13:43:52 -08:00
c_style_pointer_cast.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
common_header.hpp
use old instrinsics with staging compiler (#1970)
2025-03-12 07:29:09 -07:00
container_element_picker.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
container_helper.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
data_type.hpp
Add MX FP4 device conversion tests (#1889)
2025-03-26 19:23:01 -05:00
debug.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
dtype_fp64.hpp
Split up data_type header. (#1996)
2025-03-24 15:08:54 -07:00
dtype_vector.hpp
Fix Windows build. (#2012)
2025-04-01 12:22:10 -07:00
dynamic_buffer.hpp
Moe gemm activation (#2026)
2025-04-23 10:35:34 +08:00
e8m0.hpp
MX GEMM - New GEMM pipeline for MX data types (#2059)
2025-04-15 17:17:07 -06:00
enable_if.hpp
Rebase the PR #1520 to ROCm repo. (#1574)
2025-02-20 18:58:14 -08:00
env.hpp
Split env.hpp header from the ck.hpp header. (#2049)
2025-04-03 15:30:21 -07:00
f8_utils.hpp
Split up data_type header. (#1996)
2025-03-24 15:08:54 -07:00
flush_icache.hpp
Universal gemm flush cache (#1251)
2024-04-25 15:07:14 -05:00
functional2.hpp
add a fast compilation path for static for (0..N) (#2005)
2025-04-01 12:06:25 -07:00
functional3.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
functional4.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
functional.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
generic_memory_space_atomic.hpp
Split up data_type header. (#1996)
2025-03-24 15:08:54 -07:00
get_id.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
get_shift.hpp
Batchnorm splitk single kernel (#771)
2023-07-06 10:58:55 -05:00
ignore.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
inner_product_dpp8.hpp
Redesign the DPP8 GEMM kernel to use warp-wise component (#863)
2023-09-06 11:44:09 -05:00
inner_product.hpp
Fix the DL kernel issues on Navi3x. (#998)
2023-10-19 09:34:39 -07:00
integral_constant.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
is_detected.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
is_known_at_compile_time.hpp
Add optimized copy to ck wrapper (#1126)
2024-01-19 11:29:00 +01:00
loop_scheduler.hpp
Rebase the PR #1520 to ROCm repo. (#1574)
2025-02-20 18:58:14 -08:00
magic_division.hpp
Split up data_type header. (#1996)
2025-03-24 15:08:54 -07:00
math_v2.hpp
Rebase the PR #1520 to ROCm repo. (#1574)
2025-02-20 18:58:14 -08:00
math.hpp
Extend available elementwise operations with conv examples (#995)
2023-10-19 17:23:19 +02:00
multi_index.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
mxf4_utils.hpp
Add MX FP4 device conversion tests (#1889)
2025-03-26 19:23:01 -05:00
mxf6_utils.hpp
Split up data_type header. (#1996)
2025-03-24 15:08:54 -07:00
mxf8_utils.hpp
Add FP16/BF16<->FP8/BF8 conversions (#2035)
2025-04-03 12:42:03 -05:00
mxfp_utils.hpp
MX GEMM - New GEMM pipeline for MX data types (#2059)
2025-04-15 17:17:07 -06:00
number.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
numeric_limits.hpp
Split up data_type header. (#1996)
2025-03-24 15:08:54 -07:00
numeric_utils.hpp
Split up data_type header. (#1996)
2025-03-24 15:08:54 -07:00
random_gen.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
reduction_common.hpp
Batchnorm splitk single kernel (#771)
2023-07-06 10:58:55 -05:00
reduction_enums.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
reduction_functions_accumulate.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
reduction_operator.hpp
Remove unsupported (fp8) type from Add memory operation. (#1521)
2024-09-20 09:40:45 +02:00
scaled_type_convert.hpp
Add FP16/BF16<->FP8/BF8 conversions (#2035)
2025-04-03 12:42:03 -05:00
sequence_helper.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
sequence.hpp
simplify generate_tuple (#2043)
2025-04-08 09:00:51 -07:00
span.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
static_buffer.hpp
Jing's contribution: prototype of mixed precision gemm FP16/BF16xint4 GEMM (#1762)
2025-01-02 11:48:06 +08:00
statically_indexed_array_multi_index.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
statically_indexed_array.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
synchronization.hpp
Merging the gfx12 code into public repo. (#1362)
2024-06-27 00:33:34 -07:00
thread_group.hpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
transpose_vectors.hpp
[GEMM] Gemm universal device operation (#1154)
2024-04-13 21:03:18 -05:00
tuple_helper.hpp
Moe gemm activation (#2026)
2025-04-23 10:35:34 +08:00
tuple.hpp
Codegen hipRTC compilation (#1579)
2025-01-31 09:48:39 -08:00
type_convert.hpp
GEMM Multiply Multiply Fix (#2102)
2025-04-22 01:13:22 -07:00
type.hpp
Rebase the PR #1520 to ROCm repo. (#1574)
2025-02-20 18:58:14 -08:00
workgroup_barrier.hpp
initial stream-k implementation with example (#699)
2023-07-26 14:18:15 -05:00
workgroup_synchronization.hpp
Batchnorm splitk single kernel (#771)
2023-07-06 10:58:55 -05:00