Commit Graph

1498 Commits

Author SHA1 Message Date
Illia Silin
ddfcce82ab enable compilation for generic navi targets (#1645)
[ROCm/composable_kernel commit: 75c5bfa364]
2024-11-07 14:14:42 -08:00
rocking
b6086eed7c Fix F16 type (#1583)
[ROCm/composable_kernel commit: 3599418aa8]
2024-11-06 11:32:44 -08:00
aledudek
91228f5e50 Generic threshold calculation after merge fixes (#1618)
* Generic threshold calculation add passing num of accums

* Generic threshold - after merge fixes

* Fix cmakelists

---------

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>

[ROCm/composable_kernel commit: dcafb1de15]
2024-11-06 10:44:58 +01:00
Andriy Roshchenko
bc1f12c32f Prevent instantiation of undefined FP8 operators. (#1639)
[ROCm/composable_kernel commit: 365f39aed0]
2024-11-05 13:58:29 -08:00
Illia Silin
ecbe618edd remove gfx940;gfx941 from default target lists (#1640)
[ROCm/composable_kernel commit: 54440cf562]
2024-11-05 13:56:20 -08:00
darren-amd
bee5289f56 Statically Cast Pointer Offset (#1631)
* explicit cast ptr offset

* formating change

[ROCm/composable_kernel commit: d0e3a70a2e]
2024-11-05 09:59:08 -08:00
Illia Silin
827e5ed06d Make sure cmake can handle the xnack+/xnack- targets. (#1633)
* make sure cmake can handle xnack targets

* dont build xdl instances for gfx906:xnack-

* dont build xdl tests for gfx906:xnack-

[ROCm/composable_kernel commit: b6e74be1aa]
2024-11-05 08:53:10 -08:00
Juan Manuel Martinez Caamaño
6e74da9b87 [generate.py] Override blob list if it already exists (#1635)
Before, generate.py appended the list at the end of the output file.
When running the cmake configuration steps multiple times on the
examples, the blob list (such as fwd_blob_list.txt) would grow at every
configuration.
`library/src/tensor_operation_instance/gpu/mha/CMakeLists.txt` worked around
this issue by removing the output file if it exists.

Now, generate.py overrides the content of the output file.
There is no need for the workaround in the CMakeLists.txt;
and the issue is solved for the example projects too.

[ROCm/composable_kernel commit: 464abd235e]
2024-11-05 10:09:52 +01:00
Lin Sun
40df6ce241 Linsun/convint8 fwd instances (#1626)
Add instances for int8 grouped conv2d fwd
---------

Co-authored-by: root <root@dell300x-pla-t28-03.pla.dcgpu>
Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

[ROCm/composable_kernel commit: 0c9012fb70]
2024-11-04 16:33:20 -08:00
Bartłomiej Kocot
fadab8013c Temporary disable part of dynamic op conv instances (#1630)
* Temporary disable part of dynamic op conv instances

* fix

[ROCm/composable_kernel commit: 4f1fdbb6e3]
2024-11-04 13:34:17 -08:00
carlushuang
537ff25c21 [CK_TILE] layernorm have more accurate residual (#1623)
* more accurate residual

* modify comment

* Fix literal case in README.md

---------

Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>

[ROCm/composable_kernel commit: cb6c5d39dc]
2024-11-02 13:30:16 +08:00
Illia Silin
f29c4cebf5 Reduce build time. (#1621)
* disable fp8 gemm_universal on gfx90a and gfx908 by default

* fix cmake syntax

* fix clang format

* add ifdefs in amd_xdlops

* disable fp8 gemm instances on gfx90a by default

* update readme

[ROCm/composable_kernel commit: 03c6448ba3]
2024-11-01 13:52:23 +08:00
rocking
4faf3ab587 [Ck_tile] smoothquant (#1617)
* fix compile error

* fix typo of padding

* Add smoothquant op

* Add smoothquant instance library

* refine type

* add test script

* Re-generate smoothquant.hpp

* Always use 'current year' in copyright

* use Generic2dBlockShape instead

* Add vector = 8 instance back

* Find exe path automatically

* Simplify the api condition

* Remove debugging code

* update year

* Add blank line between function declaration

* explicitly cast return value to dim3

* refine return value

* Fix default warmup and repeat value

* Add comment

* refactor sommthquant cmake

* Add README

* Fix typo

---------

Co-authored-by: Po Yen, Chen <PoYen.Chen@amd.com>

[ROCm/composable_kernel commit: fbd654545a]
2024-11-01 13:51:56 +08:00
carlushuang
04610407b0 [layernorm] hot fix (#1620)
* hot fix ln

* some rename

[ROCm/composable_kernel commit: 550248deec]
2024-11-01 11:52:50 +08:00
carlushuang
776c87ea7e [CK_TILE] layernorm support fused-quant/fused-add (#1604)
* add prenorm/postnorm support, refactor using generate.py

* update README

* update README

* fix format

* update some description and fix format

* update format

* format

* use non-raw for loading

* format and update n4096

* dynamic-quant ready

* update readme

* support fused dynamic-quant

* update fused-quant, with smooth

* update README

* update args

* update some based on comment

[ROCm/composable_kernel commit: c3a4800c5f]
2024-10-31 14:54:53 +08:00
Bartłomiej Kocot
724312aea3 Remove virtual destructors from unary ops (#1610)
* Remove virtual destructors from unary ops

* Fixes

* Fixes

* clang format fixes

[ROCm/composable_kernel commit: 9a8a52130d]
2024-10-30 17:42:50 +01:00
rocking
4157a73c07 clang-format (#1612)
[ROCm/composable_kernel commit: 7d9111545f]
2024-10-30 08:13:30 -07:00
Adam Osewski
c0de51533d [CK-Tile] Universal gemm memory bound pipeline (#1558)
* CK-Tile GEMM with memory bound pipeline.

* Memory bound gemm pipeline.

* Fix not closed namespace.

* Block gemm mem pipeline draft.

* Do not use ck_tile:: within ck_tile namespace.

* Refactoring & Move Layout info to pipeline problem.

* Get hot loop and TailNum information before lunching kernel.

* Fixes in pipeline.

* Add comment to load_tile_raw and change variable naming style.

* Few small changes & formatting.

* Do not use macro.

* Add gtests.

* Use AccDataType for Output of MFMA instruction.

* Formatting.

* Refactor gemm examples.

* Switch over to current block gemm.

* Use currently available pipeline policy.

* Refactoring and review comment.s

* Fixes after merge.

* Add missing include.

* Add load tile overload which accepts output tensor as parameter.

* This give 8% perf boost at the cost of using more registers.

* Rename example.

* Small changes.

* Fix compilation err and lower K.

* Support different layouts for A/B

* Fix vector size for different layouts.

* Rename Alignment into VectorSize

* Unblock tests.

[ROCm/composable_kernel commit: 24d996aae1]
2024-10-30 10:05:15 +01:00
rocking
8d2057326a [Ck tile] support rmsnorm and related fusion (#1605)
* Add reduce2d new api

* Prevent user use cross warp reduction

* Fix bug of std caculation

* Add rmsnorm2d

* Add rmsnorm small example

* Remove static assert to prevent compile fail

* Add script to test performance and correctness

* Add missing cmake change

* refine naming

* refine example of rmsnorm

* Fix bug of rmsnorm

* Refine naming

* Fix cmake

* clang format

* Refine pipeline name

* Add add_rmsnorm2d_rdquant kernel

* Add reduce op

* host verification

* Fix bug of one pass pipeline

* Refine tile size

* Add two pass pipeline

* Rename two pass to three pass

* Fix bug of kSaveX == false

* Add instance library

* Add test script

* Fix bug of x verification

* Add save_x to trait

* Add README

* Move reduce2d into reduce folder

* Fix bug of welford when number of m warp > 1

* remove reduncant comment

* 1. move 06_rmsnorm2d to 10_rmsnorm2d
2. move 07_add_rmsnorm2d_rdquant to 11_add_rmsnorm2d_rdquant

* clang format and add missing header

* Add host validation of add + layernorm2d + rsquant

* Revert "Add host validation of add + layernorm2d + rsquant"

This reverts commit 936cb45797.

* Remove deprecated flag

[ROCm/composable_kernel commit: 3d60953477]
2024-10-30 15:22:56 +08:00
Qianfeng
1fc6f29f35 [CK_TILE] Add fmha fwd headdim96 support (#1608)
* Add ceil_to_qualified_tile_length()

* Rename kK0BlockLength to kQKHeaddim

* Add kSubQKHeaddim concept to support headdim96

* Fix in math.hpp to avoid using __half interfaces

* Add LdsBufferSequence instance for headdim96

* Update in fmha_fwd/fmha_fwd_splitkv codegen to support hd96 testing

* Disable hd96 instance generation in codegen fmha_fwd and fmha_fwd_splitkv to save compiling time

* Reformat one file

* Fix text alignment in fmha_fwd_splitkv.py

---------

Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>

[ROCm/composable_kernel commit: 8632221814]
2024-10-30 14:03:16 +08:00
valarLip
cb33803f3a [CK_TILE] add scatter_gather (#1609)
[ROCm/composable_kernel commit: 4d7e063a0a]
2024-10-29 18:19:29 +08:00
valarLip
3492423270 [CK_TILE] add generic_permute (#1607)
[ROCm/composable_kernel commit: 9fbd72e97e]
2024-10-29 18:05:53 +08:00
Illia Silin
75b1a7a6fe fix compilation errors for gfx12 with clang20 (#1606)
[ROCm/composable_kernel commit: 922e42a039]
2024-10-28 19:02:48 -07:00
carlushuang
a6b29fcb01 topk_softmax (#1592)
* topk_softmax

* remove some file

* fix atomix linear_offset

* address various comment, and change sfc get_index api to static(tuple)

[ROCm/composable_kernel commit: b098b71b05]
2024-10-26 23:52:49 +08:00
Bartłomiej Kocot
930195c384 Add dynamic elementwise op (#1426)
* Add dynamic elementwise op

Co-authored-by: ThruptiRajLakshmanaGowda <thruptiraj.lakshmanagowda@amd.com>

* CI issues fix

* Custom parameter value for dynamic functions - Comments addressed

---------

Co-authored-by: ThruptiRajLakshmanaGowda <thruptiraj.lakshmanagowda@amd.com>
Co-authored-by: ThruptiRajLakshmanaGowda <tlakshma@amd.com>

[ROCm/composable_kernel commit: 31bf253aeb]
2024-10-26 15:22:37 +02:00
Po Yen Chen
c37f8d33f4 [CK_TILE] More fmha splitkv optimizations (#1588)
* Use pre-defined constants for readability

* Use vector write for o_acc tensor

* Remove no-longer used policy method

* Deprecate no-longer used policy/pipeline

* Specify gemm0/gemm1 block warps separately in codegen

* Fix wrong ps_idx creation logic

* Add single-warp block gemm

* Supoprt single-warp gemm0

* Make MakeCBlockTile() as static method

* Use MakeCBlockTile() to get underlying tile distribution

* Use kNumGemm1Warps to compute # threads for gemm1

* Put normal case in the if clause

* Refine fmha splitkv block mapping

* Refine & fix the lse_acc/o_acc layout

* Fix wrong LDS size for K tile

* Use kK0=64 for hdim=128,256 fmha splitkv kernels

* Use kK1=64 for hdim=32,64,128 fmha splitkv kernels

* Undo kK0/kK1 changes

* Use more reasonable GetAlignmentV() computation

* Using store_tile() in fmha splitkv kernel epilogue

[ROCm/composable_kernel commit: 54f0e6f4bb]
2024-10-26 18:35:45 +08:00
valarLip
85cf31cf40 add int8 gemm multiply multiply a8w8 (#1591)
* add int8 gemm multiply multiply a8w8

* uncomment

* clang-format-12

* Add example_gemm_multiply_multiply_xdl_int8

* Remove shell scripts

* update preprocess number for mi308; bring back printout in ckprofiler

* format

---------

Co-authored-by: chenjun <junchen2@amd.com>
Co-authored-by: Haocong WANG <haocwang@amd.com>
Co-authored-by: carlushuang <carlus.huang@amd.com>

[ROCm/composable_kernel commit: 37f7afed1e]
2024-10-26 16:39:34 +08:00
Max Podkorytov
8471aad362 add parsing grouped conv fwd instances
[ROCm/composable_kernel commit: eda5938386]
2024-10-25 08:25:53 -07:00
Rostyslav Geyyer
b590a9e4b3 Update GPU verification (#1596)
* Update inits

* Update static_cast to type_convert

* Add verification option selection

[ROCm/composable_kernel commit: 7d576f1748]
2024-10-25 08:13:46 -07:00
aledudek
2c3399d29b Generic threshold calculation (#1546)
* Calculate generic relative threshold pool3dfwd

* Calculate absolute error threshold pool3d fwd

* Generic threshold calculation take max input for relative error pool3dfwd

* Remove max possible value for error calculation at runtime

* Remove debug print in pool3dfwd

* Pool3d fwd adjusted types in generic threshold calculation

* Generic threshold calculation take into account number of accumulations and accdatatype

* Generic threshold fix final error formula

* Generic threshold calculation - num of accs fix

* Generic threshold calculation - adjust absolute error

* Generic threshold calculation - OutDataType in absolute error

[ROCm/composable_kernel commit: 9385caa306]
2024-10-25 12:46:24 +02:00
dummycoderfe
77b0a4b1d2 hot_fix epsilon pos (#1597)
Co-authored-by: dummycoderfe <noplydummmycoder@163.com>

[ROCm/composable_kernel commit: 9183ce69ca]
2024-10-25 11:17:45 +08:00
Illia Silin
a9fedd4d20 fix the logic of enabling XDL and WMMA instances (#1595)
[ROCm/composable_kernel commit: 8e22e1ae31]
2024-10-23 15:55:39 -07:00
Bartłomiej Kocot
47c726d0ba [POST MERGE PR] Enable grouped conv bwd wei bf16 NGCHW (#1594)
[ROCm/composable_kernel commit: cedccd59c9]
2024-10-23 12:02:33 +02:00
Jatin Chaudhary
2a074cd391 Explicit cast values to half (#1593)
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

[ROCm/composable_kernel commit: 4d5248e2d1]
2024-10-22 11:17:32 -07:00
Bartłomiej Kocot
5f8fefae24 Enable grouped conv bwd wei bf16 NGCHW (#1589)
* Enable grouped conv bwd wei bf16 NGCHW

* fixes

* fixes

* Fixes

* fixes

* fixes

* Fixes

[ROCm/composable_kernel commit: 82fc53835a]
2024-10-22 16:18:28 +02:00
ltqin
45d7cc2f41 update layernorm (#1570)
* port layernorm

* change warp_welford.hpp

* Update warpshuffle

* 1. Add save mean and save std back
2. Move construction of tensor_view and tile_window to operator()

* refine welford max count calculation

* unify layernorm api

* Rename file

* Remove save mean and inv std

* Revert "refine welford max count calculation"

This reverts commit 022365802b.

* Fix order of parameter

* refine welford max count calculation again

* Remove fp32 instances

* Fix bug of padding

* refactor api

* Support bf16

* Extract common function

* Refine arg of operator()

* Add kMThreadPerBlock to template parameter

* clang format

* Refine variable name

* Refine file name

* remove redundant line

* refactor layernorm2d pipeline and add block-per-block utility

* fix name

* rename more

* add more block-per-tile instance

* remove duplicated define

* update instance for 2048, 1024 case

* support up to 2048 now

* opt loading

* add n1536

* Add two pass pipeline

* format

* Fix incorrect type

* parallel compilation

* Use smaller N

* fix 2p pass

* Support Repeat_M in distribution

* Refine nameing

* Add reduce example

---------

Co-authored-by: letaoqin <letaoqin@amd.com>
Co-authored-by: aska-0096 <haocwang@amd.com>
Co-authored-by: rocking <ChunYu.Lai@amd.com>
Co-authored-by: carlushuang <carlus.huang@amd.com>

[ROCm/composable_kernel commit: 0394f8a713]
2024-10-22 09:26:18 +08:00
Rostyslav Geyyer
124989ade1 Update default stride (#1576)
* Update default stride value to -1

* Fix format

* Revert "Fix format"

This reverts commit ae0c3649ec.

---------

Co-authored-by: Harisankar Sadasivan <135730918+hsadasiv@users.noreply.github.com>

[ROCm/composable_kernel commit: 3f710930f6]
2024-10-21 08:45:22 -07:00
spolifroni-amd
d870509879 added link to documentation (#1578)
[ROCm/composable_kernel commit: 794f2d64a8]
2024-10-21 08:35:57 -07:00
dependabot[bot]
6081be772e Bump rocm-docs-core from 1.8.2 to 1.8.3 in /docs/sphinx (#1587)
Bumps [rocm-docs-core](https://github.com/ROCm/rocm-docs-core) from 1.8.2 to 1.8.3.
- [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/ROCm/rocm-docs-core/compare/v1.8.2...v1.8.3)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-type: direct:production
  update-type: version-update:semver-patch
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

[ROCm/composable_kernel commit: d0565e33d6]
2024-10-21 08:34:53 -07:00
Thomas Ning
a415fb6906 Ck profiler instance support (#1575)
* The draft on ckProfiler instance add

* support the ck profiler instance with same data types

* add a small feature on the M and N variable switch.

* Partially solve the incorrect result problem

* fix based on ci cd

[ROCm/composable_kernel commit: 560917b161]
2024-10-21 22:47:48 +08:00
Po Yen Chen
8cc793d26a [CK_TILE] Optimize fmha splitkv & splitkv combine kernels (#1577)
* Use smaller width for lse_accum dist tensor

* Update pipeline comment

* Fix wrong distribution for lse_accum

* Remove duplicate dim in lse_accum dist encoding

* Decide fmha splitkv combine kernel kBlockSize by kM0

* Remove assumption of MPerThread=1

* Add log<4> & log<8> specialization

* Enlarge occupancy array

* Fix vector size for small tile

* Add support for kMaxSplits=8

* Re-format gemm.hpp

* Use 16x16x16 warp gemm for fwd_splitkv

* Centralize policy code changes

* Leave fp8/bf8 tile settings unchanged

[ROCm/composable_kernel commit: 95e722a3b3]
2024-10-21 10:52:11 +08:00
Haocong WANG
3f37d10e43 disable bad instance detected on MI308CPX (#1584)
[ROCm/composable_kernel commit: a285d6f9b5]
2024-10-18 08:46:11 -07:00
Illia Silin
c2d1876f70 add the lsr-drop-solution=1 compiler flag (#1582)
[ROCm/composable_kernel commit: 88e6fa7fdb]
2024-10-18 08:25:54 -07:00
Qianfeng
2587dc434b [CK_TILE] Improve headdim96 performance for fmha-bwd (#1573)
* Add kQKHeaddimForGemmN and kVHeaddimForGemmN in order to support headdim 96

* Remove the using of MakeKRegBlockDescriptor and MakeVRegBlockDescriptor

* Fix in bwd_piple_default_policy

* Remove kQKHeaddim and rename kQKHeaddimForGemmN to kQKHeaddim in the bwd kernel and pipelines

* Replace kVHeaddimForGemmN by kVHeaddim and kDoDvHeaddim

* Update to hd96 tile settings

* Add smoke test scripts for fmha-bwd hd96

* Revert "Add smoke test scripts for fmha-bwd hd96"

This reverts commit 7ca7e1a93d.

* Remove hd96 tile settings in fmha_bwd codegen to save compiling

* Fix lost code line in bwd_pipeline_default_policy

* Merge kDoDvHeaddim/kPadHeadDimDoDv to kVHeaddim/kPadHeadDimV and remove TileFmhaBwdTraits

* Rename KRegSliceBlockDescriptor/VRegSliceBlockDescriptor to KRegBlockDescriptor/VRegBlockDescriptor

* tiny adjustments

---------

Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
Co-authored-by: danyao12 <Dan.Yao@amd.com>

[ROCm/composable_kernel commit: 14c3cfb1c6]
2024-10-16 18:14:32 +08:00
Paul Fultz II
344c33eb4e Build codegen as standalone (#1556)
* Build codegen as standalone

* Add exception for device tests

* Use local filesystem header

* add a codegen test CI stage and daily build

---------

Co-authored-by: illsilin <Illia.Silin@amd.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

[ROCm/composable_kernel commit: 10158b0ffd]
2024-10-15 13:20:42 -07:00
Bartłomiej Kocot
86a6d14c7d [CK_TILE] Add block universal gemm pipeline policy (#1557)
* [CK_TILE] Add block universal gemm pipeline policy

* Fixes

* fixes2

* Fixes3

* fixeS

[ROCm/composable_kernel commit: d02a92cc0d]
2024-10-15 13:53:41 +02:00
Po Yen Chen
35e652e286 Apply ROCm 6.2 WA to ROCm 6.3 and later (#1563)
[ROCm/composable_kernel commit: 9868fd0245]
2024-10-15 18:02:41 +08:00
Rostyslav Geyyer
63c478e14f Add custom type vector support (#1333)
* Add non_native_vector_type

* Add a test

* Add non-native vector type

* Fix CTOR

* Fix non-native vector type of 1

* Fix CTORs

* Use vector_type to cover non-native implementation as well

* Update the test

* Format

* Format

* Fix copyright years

* Remove BoolVecT so far

* Add AsType test cases

* Update assert error message

* Remove redundant type

* Update naming

* Add complex half type with tests

* Add tests for vector reshaping

* Add missing alignas

* Update test/data_type/test_custom_type.cpp

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>

* Compare custom types to built-in types

* Add default constructor test

* Add an alignment test

---------

Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>

[ROCm/composable_kernel commit: 4cf70b36c1]
2024-10-14 11:56:45 -05:00
Bartłomiej Kocot
f51ed2ad28 Add transpose scale amax example (#1547)
* Add transpose scale amax example

* fixes

* Tune reduce instance

[ROCm/composable_kernel commit: f21cda2536]
2024-10-14 17:39:38 +02:00
Thomas Ning
53df021563 decouple the calling from gemm_pipeline (#1571)
* decouple the calling from gemm_pipeline

* clang format

[ROCm/composable_kernel commit: 35c1777d59]
2024-10-14 13:59:26 +08:00