Commit Graph

1051 Commits

Author SHA1 Message Date
Illia Silin
1a389f2d2e Enable sccache in the default docker and CI. (#1009)
* replace ccache with sccache, pin package versions

* put ccache back temporarily to avoid breaking other CI jobs

* add sccashe_wrapper.sh script

* fix the package version syntax

* fix the pymysql package issue

* run sccache_wrapper before build if ccache server found

* set the paths before calling the sccache_wrapper

* use /tmp instead of /usr/local for cache

* try using sccache --start-server instead of wrapper

* try using redis server with sccache

* define SCCACHE_REDIS

* add redis and ping packages, and redis port

* use the new sccache redis server

* do not use sccache with staging compiler

* fix the condition syntax

* add stunnel to redis

* add tunnel verification

* separate caches for different architectures

* fix syntax for the cache tag

* quse double brackets for conditions

* add bash line to the script

* add a switch for sccache and only use it in build stage

* run check_host function when enabling sccache

* fix the invocation tags for sccache

* fix groovy syntax

* set the invocation tag in groovy

* disable sccache in clang-format stage

* try another syntax for invocation tags

* use local sccache server if can't connect to redis

* fix script syntax

* update README

* refresh readme

* readme updates

* remove the timing and verification caveat from readme

---------

Co-authored-by: Lisa Delaney <lisa.delaney@amd.com>

[ROCm/composable_kernel commit: 4e44a9e8da]
2023-10-30 13:16:29 -07:00
Illia Silin
60b282fb62 Fix the fp8 gemm for large tensors on MI300. (#1011)
* Fix the fp8 conversion

* Try clipping value before conversion

* Fix return

* Simplify with a const

* reduce the gemm input tensor values to reduce round-off error

* replace if-else with lambda

* fix syntax

---------

Co-authored-by: Rostyslav Geyyer <rosty.geyyer@amd.com>

[ROCm/composable_kernel commit: f46a6ffad8]
2023-10-27 21:10:47 -07:00
dependabot[bot]
6cf76b5095 Bump sphinxcontrib-bibtex from 2.5.0 to 2.6.1 in /docs/sphinx (#871)
Bumps [sphinxcontrib-bibtex](https://github.com/mcmtroffaes/sphinxcontrib-bibtex) from 2.5.0 to 2.6.1.
- [Changelog](https://github.com/mcmtroffaes/sphinxcontrib-bibtex/blob/develop/CHANGELOG.rst)
- [Commits](https://github.com/mcmtroffaes/sphinxcontrib-bibtex/compare/2.5.0...2.6.1)

---
updated-dependencies:
- dependency-name: sphinxcontrib-bibtex
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

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

[ROCm/composable_kernel commit: 6fe0bc7e72]
2023-10-26 14:30:32 -06:00
zjing14
c2bd184885 Enabled padding for regular gemm (#1004)
* add mnk padding for fp8

* add padding for row_col layout

* added padding for fp32

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: bec84efbb7]
2023-10-23 16:52:53 -05:00
Bartłomiej Kocot
42608c87ea Fix cmake dtype check (#989)
* Fix instances dtype check

* Fix source dtypes seletor for examples and tests

* Sync with new cmakefile changes

* Remove not needed ifdefs

* Remove not needed ifdefs

[ROCm/composable_kernel commit: ac0e006766]
2023-10-21 22:19:43 +02:00
Rostyslav Geyyer
07d78c032b Fix bf8 conversion issues (#1003)
* Fix the conversion

* Add bf8 functionality

* Enable example on MI200 as well

[ROCm/composable_kernel commit: 1fd27d520f]
2023-10-20 08:00:45 -05:00
Illia Silin
75e7948ac8 Fix the DL kernel issues on Navi3x. (#998)
* apply the patch for dl kernels on gfx11

* build DL kernels on navi32 CI

[ROCm/composable_kernel commit: f7331c603b]
2023-10-19 09:34:39 -07:00
Qianfeng
2caa0a74f0 Misc fixes (#994)
* reinterpret_cast to const char* in dumpBufferToFile to be compatible with both const and non-const input pointers

* Add seed input to GeneratorTensor_4 for normal_distribution generator

* Add GetTypeString() for DeviceElementwiseImpl

* Add HIP_CHECK_ERROR macro

[ROCm/composable_kernel commit: b4fc4d0b8d]
2023-10-19 11:26:04 -05:00
Bartłomiej Kocot
2d230d0f5c Extend available elementwise operations with conv examples (#995)
* Extend available elementwise operations with conv examples

* Fixes

* Remove not needed convert

* Update CMakeFile and dir name

[ROCm/composable_kernel commit: 82f3a835d5]
2023-10-19 17:23:19 +02:00
Po Yen Chen
7afa128670 Avoid force setting ENABLE_PIPELINE_V2_OPT to OFF (#961)
* Avoid force setting ENABLE_PIPELINE_V2_OPT to OFF

* Remove compilation option variable MAX_ILP_OPTS

[ROCm/composable_kernel commit: deef92d5d0]
2023-10-19 23:19:07 +08:00
Bartlomiej Wroblewski
d2049bc4e7 Change 1d,2d,... to 1D,2D,... (#997)
[ROCm/composable_kernel commit: 0abc0f87db]
2023-10-19 16:53:18 +02:00
rocking
815ed3a1f9 Layernorm and groupnorm support to save mean and inverse std in forward (#929)
* save mean and inverse std in normalization

* Save mean and inverse std in splitK

* Vector save mean and inv std

* Modify instance for save mean and std

* simplify the layernorm example

* Save mean and std in groupnorm example

* Save mean and inv std in ckProfiler and test

* Remove compute data type from base class

* Save mean and inv std in client example

* Add changelog

* clang format

* Fix compile error

* Refine naming

* Avoid error in bf16

* revert changelog

[ROCm/composable_kernel commit: 3696fe1c76]
2023-10-19 07:36:29 +08:00
zjing14
5c3f35d358 fixed math-ci error; suspend a warning (#996)
Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 58338bb203]
2023-10-18 16:30:13 -07:00
zjing14
dc94c20258 Clean DTYPES conditions in CMake (#974)
* Add a condition to build fp8 instances

* simplified buffer_load/store

* add bfp8/fp8

* fixed

* remove all f8/bf8 condition include folder

* fixed cmake conditions

* fixed DTYPES=fp16/bfp16

* fix

* fixed buffer_load

* fixed buffer_store

* fix

* clean example cmake files

* fixed ci

* fixed cit

---------

Co-authored-by: Rostyslav Geyyer <rosty.geyyer@amd.com>
Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: bf435140dc]
2023-10-18 11:14:14 -05:00
zjing14
522bbcb766 Add contraction_multi_abd (#972)
* add gridwise_multi_abd

* move element_op into RunRead

* merge element_wise op with data read

* add multiABD example

* allow packed elementwise_op

* changed example

* clean

* clean

* add is_detected

* fix

* minor fix

* add scaleAdd_vec4 example

* init commit for contraction_multi_ABD

* add examples

* add examples of multiA and broadcast

* update example

* fixed comments

* Update cmake-ck-dev.sh

* Update cmake-ck-dev.sh

* Add comments into the example

* Update CMakeLists.txt

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 1cc36ba5fb]
2023-10-17 20:17:58 -05:00
zjing14
fc1434e0f2 added ab_elementwise_op support into splitK Gemm (#956)
* add ab_elementwise

* fixed ci

* fixed a merge issue

* fixed pr comments

* fixed a conflict

* remove 61_example

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: bf0addb575]
2023-10-17 09:24:02 -05:00
Bartłomiej Kocot
49f179f755 Add grouped conv bwd weight wmma (#985)
* Add grouped conv bwd weight wmma

* Update README, changelog, profiler

* Minor fixes

* Fix grouped conv bwd wei dl kernel

* Minor fixes

* Minor stylistic fixes

[ROCm/composable_kernel commit: 16d7c4d2f7]
2023-10-17 10:32:26 +02:00
zjing14
4052b33a3f workaround with float (#992)
Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 39430bfdeb]
2023-10-16 15:42:59 -07:00
Illia Silin
aa8632b821 Add hipTensor build and test to CK CI. (#990)
* add a hipTensor test to CI

* use jenkins git plugin

* change hipTensor folder location in CI

* change the git method for hipTensor

* run tests usign ctest

* check the hipTensor contents

* only build hipTensor on MI100/200

* pull hipTensor as zip archive

* fix jenkins syntax

* add path to the CK installation

* combine build commands into one shell

* change jenkins syntax for CK installer path

* try different syntax

* allow unzip overwrite

* fix jenkins file syntax

* remove any old versions of hipTensor before building

* add option to select hipTensor branch for testing

[ROCm/composable_kernel commit: 707ad0025f]
2023-10-16 08:40:26 -07:00
Rostyslav Geyyer
9d72b08639 Add splitk gemm fp16 @ fp16 with fp8 compute instances (#983)
* Add ComputeType

* Update for compatibility

* Add instances

* Update profiler api

[ROCm/composable_kernel commit: fa753f27ba]
2023-10-13 16:27:11 -05:00
zjing14
1271deb162 add vector_type support into thread_copy_v3r1 (#969)
* add vector_type support into thread_copy_v3r1

* remove unncessary type_convert

* fixed datatype

* fixed dataType

* changed API with is_packx_invocable

* changed example

* add missing cmake file

* fixed ci

* fixed cmake

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 2ce9b56c64]
2023-10-13 15:11:43 -05:00
dependabot[bot]
af29213e42 Bump gitpython from 3.1.31 to 3.1.35 in /docs/sphinx (#898)
Bumps [gitpython](https://github.com/gitpython-developers/GitPython) from 3.1.31 to 3.1.35.
- [Release notes](https://github.com/gitpython-developers/GitPython/releases)
- [Changelog](https://github.com/gitpython-developers/GitPython/blob/main/CHANGES)
- [Commits](https://github.com/gitpython-developers/GitPython/compare/3.1.31...3.1.35)

---
updated-dependencies:
- dependency-name: gitpython
  dependency-type: indirect
...

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

[ROCm/composable_kernel commit: a3c8026518]
2023-10-12 12:43:54 -06:00
zjing14
775a87175c simplified buffer_load/store (#971)
* simplified buffer_load/store

* add bfp8/fp8

* fixed

* fixed buffer_load

* fixed buffer_store

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: f3b02ecfd2]
2023-10-11 20:29:01 -05:00
zjing14
91e1cf6750 Revert "Grouped Gemm with looping over the tiles. (#788)" (#982)
This reverts commit 43fe5037d4ff9d07365e5d3b8f5b31676a8ff9da.

[ROCm/composable_kernel commit: c99323be6e]
2023-10-11 14:27:29 -05:00
Adam Osewski
34b77070f3 Grouped Gemm with looping over the tiles. (#788)
* Introduce LocalBlockToCTileMap.

* Change the signature of CalculateBottomIndex() function which now does
not accept any argument. The B2C map which is already passed as an
argument to the kernel Run function is calculating block's local id
already outside at kernel entry point __global__ function.
The LocalB2C map stores as members local block ID.

* Use LocalBlockToCTile map in device ops.

* First draft of tile loop work distribution.

* Fix typo.

* Simplify kernel arguments.

Calculate descriptors & B2C maps on the device.

* Use looping kernel.

* Fix B2C constructor.

* Fix Navi21 errors.

* Calculate tile start/end in device kernel.

* Change Run API to accept user provided workspace buffer.

* Add new line at EOF.

* Move Gemm KernelArguments to device op interface.

* Remove unused code.

* Update API.

* Launch grid size which is min of occupancy vs tile count

* Get back to use constant memory for gemm descriptors.

* Remove unused code.

* Add default virtual method implementation.

* Update comments to conform with doxygen style.

* Fix doc style and unused parameters.

* Add thread cluster lengths to kernel name.

* Remove old splitk impl and replace it with tile looping one.

* Modify instances.

* set KPerBlock to 64
* maximize wherever possible vector load size.

* Fix instances cluster lengths.

* Change comment style.

* Use 128b store where possible in instances.

* Update test cases, since KPerBlock has doubled.

* Update output stream operator for Sequence.

* Add pipeline version to GroupedGEMM device op type string.

* Fix pipeline version type logging.

* Fix input tensors type after merge.

* Fix compiler error.

* Fix output stream operator for Pipeline version.

* Store using 128b.

* Set of instances with kpb 32/64

* Limit number of instances

* Remove commented out instances.

* Fix function name.

* Limit the number of instances.

Add pipline version to the regular instances

* Change thr cluster layout for reading B tensor.

* disabled failed instances

---------

Co-authored-by: Adam Osewski <aosewski@amd.com>
Co-authored-by: zjing14 <zhangjing14@gmail.com>
Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: a4f72a314a]
2023-10-10 22:21:15 -05:00
Bartłomiej Kocot
a080f37211 Fix MNKPadding in gridwise_gemm_xdlops_v2r3 (#981)
[ROCm/composable_kernel commit: 98c8071475]
2023-10-10 21:48:07 +02:00
zjing14
0053cbab22 Fixed f8_gemm NaN (#975)
* workaround nan problem by changing output to fp16

* enable f8/bf8 gemm tests on MI200

* workaround f16 to f8 conversion

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: ac9595a9f1]
2023-10-10 10:30:26 -05:00
Lauren Wrubleski
aeb94be7b4 Replace CMake return from later CMake (#970)
[ROCm/composable_kernel commit: 5913609168]
2023-10-05 14:58:58 -07:00
Illia Silin
3b3782dd07 Revert "Add support for mixed precision in contraction scale and bilinear" (#967)
* Revert "Add support for mixed precision in contraction scale and bilinear (#936)"

This reverts commit f7aff936cb9d02dc8e53a8a3ea8648e1058253a2.

* revert commits #957 and #960

[ROCm/composable_kernel commit: 4daedf8ca5]
2023-10-05 14:58:23 -07:00
zjing14
41eb6c1899 remove example 60 (#963)
Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 570ff3ddbe]
2023-10-05 09:41:01 -07:00
zjing14
f60e9faeb1 Grouped conv bwd data with fp16 input and bf8fp8 comp (#962)
* Add f8 bf8 gemm example

* Add element-wise ops

* Add intrinsics

* Update reference calculation

* Add an additional type option for xdlops gemm

* Fix build process

* Add bf8 to buffer addressing

* Update blockwise op, split typeA and typeB

* Update for compatibility

* Uppdate naming to f8->fp8

* Update naming

* Format

* Update naming (#937)

* Add a client example

* Add computetypes to device and gridwise ops

* Add instances, update instance factory

* Format

* Fix a flag

* Add ckProfiler mode

* Fix typos

* Add an example

* Add bf8 generator

* add bf8 mfma; fixed type_convert for bf8

* move verfication ahead of timing

* Update reference calculation

* Fix reference

* Narrow down float init range

* Fix bf8 bf8 mfma

* Add bf8 @ fp8 mfma

* Update example

* Update instances

* Update profiler api

* Update for compatibility

* Format

* Remove extra example

* Clean up

* workaround convert

* added instance of f16_bf8f8, and client example

* fixed mfma selector

* format

---------

Co-authored-by: Rostyslav Geyyer <rosty.geyyer@amd.com>
Co-authored-by: Rostyslav Geyyer <46627076+geyyer@users.noreply.github.com>
Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 04f93aadb8]
2023-10-04 18:04:27 -05:00
Rostyslav Geyyer
6a3eedbff0 Add conv bwd weight fp16 comp bf8 fp8 op, instances and example (#945)
* Add f8 bf8 gemm example

* Add element-wise ops

* Add intrinsics

* Update reference calculation

* Add an additional type option for xdlops gemm

* Fix build process

* Add bf8 to buffer addressing

* Update blockwise op, split typeA and typeB

* Update for compatibility

* Uppdate naming to f8->fp8

* Update naming

* Format

* Update naming (#937)

* Add a client example

* Add computetypes to device and gridwise ops

* Add instances, update instance factory

* Format

* Fix a flag

* Add ckProfiler mode

* Fix typos

* Add an example

* Add bf8 generator

* add bf8 mfma; fixed type_convert for bf8

* move verfication ahead of timing

* Update reference calculation

* Fix reference

* Narrow down float init range

* Fix bf8 bf8 mfma

* Add bf8 @ fp8 mfma

* Update example

* Update instances

* Update profiler api

* Update for compatibility

* Format

* Remove extra example

* Clean up

* workaround convert

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 42facfc6b7]
2023-10-04 08:19:08 -05:00
zjing14
ffaff83a2f 3d grouped conv fwd with input/output fp16 and comp fp8 (#931)
* add f8 comp instance

* fixed

* fixed comments

* rename

* fixed dtype

* format

* fixed CI

* fixed ci

* add missing ComputeType

* fixed cit

* fixed

* Update cmake-ck-dev.sh

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: e921e1f08d]
2023-10-03 20:04:26 -05:00
zjing14
a7a9249cef changed test for grouped_gemm to be random (#959)
Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 5311d1b325]
2023-10-03 09:32:58 -05:00
zjing14
33859062bd Fixed contraction issues (#960)
* add missing ComputeType

* fixed

* Update cmake-ck-dev.sh

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: aa46039f2d]
2023-10-03 09:32:44 -05:00
zjing14
498e886c85 add generic instances (#947)
Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: f477fca436]
2023-10-03 09:32:28 -05:00
Rostyslav Geyyer
28a1199b62 Add fp8 @ bf8 gemm support and example (#933)
* Add f8 bf8 gemm example

* Add element-wise ops

* Add intrinsics

* Update reference calculation

* Add an additional type option for xdlops gemm

* Fix build process

* Add bf8 to buffer addressing

* Update blockwise op, split typeA and typeB

* Update for compatibility

* Uppdate naming to f8->fp8

* Update naming

* Format

[ROCm/composable_kernel commit: bd09b5c538]
2023-10-02 16:39:03 -05:00
Illia Silin
6a8658812a get rid of gfx900/906, set rocm5.7 as default (#958)
[ROCm/composable_kernel commit: 59dbb01fd1]
2023-10-02 12:01:11 -07:00
zjing14
50c12c6c43 Contraction multi abd (#957)
* add gridwise_multi_abd

* move element_op into RunRead

* merge element_wise op with data read

* add multiABD example

* allow packed elementwise_op

* changed example

* clean

* clean

* add is_detected

* fix

* minor fix

* add scaleAdd_vec4 example

* init commit for contraction_multi_ABD

* add examples

* add examples of multiA and broadcast

* update example

* fixed comments

* Update cmake-ck-dev.sh

* Update cmake-ck-dev.sh

* Add comments into the example

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 9d58c42103]
2023-10-02 09:18:36 -05:00
Illia Silin
24eac8e7b4 add gfx942 target to the daily ckprofiler package (#955)
[ROCm/composable_kernel commit: 6b5f647371]
2023-09-29 08:55:25 -07:00
Bartlomiej Wroblewski
ce003d6493 Add support for mixed precision in contraction scale and bilinear (#936)
* Extract common functionality to separate files

* Reference contraction: Remove incorrect consts from type_converts

* Reference contraction: Add missing type_convert for dst value

* Reference contraction: Fix incorrect order of B matrix dimensions

* Add support for mixed precision in contraction scale and bilinear

* Move using statements from instances to a common file

* Move using statements from examples to a common file

* Fix the order of B matrix dimensions across examples and profiler

* Fix the computation of error threshold

* Make ComputeDataType an optional argument

* Include possible DataType -> ComputeDataType casting error in the threshold

* Remove commented code

[ROCm/composable_kernel commit: f07485060e]
2023-09-29 10:54:31 -05:00
Bartłomiej Kocot
612cbbdc54 Add grouped conv bwd data wmma (#950)
* Add grouped conv bwd data wmma

* Fix copyrights

* Add instances with smaller NPerBlock

* Update interface test

* Minor stylistic fixes

* Minor stylistic fixes

[ROCm/composable_kernel commit: cb53874002]
2023-09-28 23:10:18 +02:00
Bartłomiej Kocot
254844f582 Add grouped convolution changes to changelog (#952)
* Add grouped convolution changes to changelog

* Fix 0.2.0 ck release rocm version

* Suggested CHANGELOG.md edits

* Update CHANGELOG.md

* Update CHANGELOG.md

* Update CHANGELOG.md

* Update CHANGELOG.md

* Update CHANGELOG.md

* Update CHANGELOG.md

---------

Co-authored-by: Lisa <lisajdelaney@gmail.com>

[ROCm/composable_kernel commit: 271ef645ac]
2023-09-28 18:18:32 +02:00
Illia Silin
96f752aba9 Fix gemm_splitk test, add hip_check_error after kernel calls in kernel_launch. (#951)
* Added error check after kernel launch (#919)

Co-authored-by: Xiaodong Wang <xdwang@meta.com>
Co-authored-by: Xiaodong Wang <xw285@cornell.edu>

* remove M=0 test cases for test_gemm_splitk

---------

Co-authored-by: Xiaodong Wang <xdwang@meta.com>
Co-authored-by: Xiaodong Wang <xw285@cornell.edu>

[ROCm/composable_kernel commit: bc1108bb3e]
2023-09-27 15:19:33 -07:00
Bartlomiej Wroblewski
bf38d27453 Handle type conversions to a const datatype (#944)
* Handle type conversions to a const datatype

* Review: Handle X being const data type as well

* Review: Remove typo

[ROCm/composable_kernel commit: f4af5aed8b]
2023-09-27 15:02:42 -05:00
Bartłomiej Kocot
be5cb244c0 Add column to image kernel (#930)
* Add column to image kernel

* Minor fixes for dtypes and client examples

* Disable tests for disabled dtypes

* Disable add instances functions for disabled data types

* Minor stylistic fixes

* Revert "Disable add instances functions for disabled data types"

This reverts commit 728b869563.

* Instances reduction

* Add comments in device_column_to_image_impl

* Update changelog and Copyrights

* Improve changelog

[ROCm/composable_kernel commit: e2243a4d1e]
2023-09-27 17:19:06 +02:00
zjing14
fb513ac42b Add multiple A/B support (#906)
* add gridwise_multi_abd

* move element_op into RunRead

* merge element_wise op with data read

* add multiABD example

* allow packed elementwise_op

* changed example

* clean

* clean

* add is_detected

* fix

* minor fix

* add scaleAdd_vec4 example

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 11676c7e49]
2023-09-26 21:16:23 -05:00
Illia Silin
f9ce51a187 Use lower case for ckprofiler package. (#948)
* split ckProfiler gfx9 package into gfx90 and gfx94

* use lower case for package names

[ROCm/composable_kernel commit: 420b5a0382]
2023-09-26 17:43:09 -07:00
zjing14
3fe6761718 Fixed Gemmv2r3 kpad (#938)
* added kpad support into v2r3

* add generic instances

* fixed comments

* fixed mnk padding

* Update device_batched_gemm_xdl.hpp

* fixed kpad

---------

Co-authored-by: Jing Zhang <jizha@amd.com>

[ROCm/composable_kernel commit: 48ba6e8a69]
2023-09-26 18:40:00 -05:00
Rostyslav Geyyer
b74d4f5fc6 Add fp8 gemm instances (#920)
* Add fp8 gemm instances

* Update instance naming

[ROCm/composable_kernel commit: 94bfa50256]
2023-09-26 14:59:33 -05:00