Commit Graph

47 Commits

Author SHA1 Message Date
Haocong WANG
ec634a3d32 Add CMake Option "USE_OPT_NAVI3X" (#647)
* Add CMake Option "USE_OPT_NAVI3X"

* remove navi3x opt compile option from cmake script

[ROCm/composable_kernel commit: 4e097ad283]
2023-03-29 14:07:33 -05:00
Lauren Wrubleski
a4d6101e99 Add packages for examples and profiler (#502)
* Add packages for example and profiler

* correct TEST_NAME -> EXAMPLE_NAME

[ROCm/composable_kernel commit: 37f2e91832]
2022-11-10 13:19:33 -06:00
Illia Silin
47484e4eea reduce the number of default targets (#489)
* reduce the number of default targets

* re-write the setting of target flags

* move all options to one place

* add new custom target instances for installing CK

[ROCm/composable_kernel commit: a5059f8f90]
2022-10-27 15:17:56 -06:00
Illia Silin
6fb932714d Add an option to build CK with clang directly (#387)
* replace hipcc compiler with clang++

* build client app with hipcc

* build client app with clang

* add an option to build with hipcc ro clang

* fix the environment for client app

* fix setting up compiler in cmake_build

* change the way the compiler is set

[ROCm/composable_kernel commit: 1e5b59df22]
2022-08-26 12:51:39 -05:00
Adam Osewski
2332d3657d int4 data type (#364)
* Introduce int4 data type.

* Add unit-tests for int4

* Compile int4 UT only when int4 enabled.

* clang-format

Co-authored-by: Adam Osewski <aosewski@amd.com>

[ROCm/composable_kernel commit: e00149ac67]
2022-08-18 14:53:47 -05:00
Anthony Chang
cc5140ac96 Fused attention (#345)
* initial stub for gemm_gemm_xdl_cshuffle

* set up example code

* compiles

* prevent integer overflow

* harmonize interface between ref_gemm and ref_batched_gemm

* batched_gemm_gemm

* fix example

* host tensor gen: diagonal pattern in lowest two-dimensions only

* make c descriptors containing only integral constants

* clean up

* add BlockwiseGemmXdlops_v2 while exploring an unified approach

* implement proper interface

* tidy up example

* fix compilation warnings

* coarsely controlled 2nd gemm padding

* remove rocm-cmake's hard requirement for certain revision

* clang-format

* resolve merge conflict

* fix compilation error on gfx10

* adds acc0 elementwise op to interface

* attention host validation

* add blockwsie softmax v1

* iteratively update softmax+gemm

* transpose both gemm0 and gemm1 xdl output so as to avoid broadcasting softmax max/sum

* add init method for easier debugging

* do away with manual thread cluster calculation

* generalize blockwise softmax interface

* row-wise softmax sum & max

* format

* rename to DeviceBatchedGemmSoftmaxGemm

* add gemm_softmax_gemm instances and tests

* comment

Co-authored-by: ltqin <letao.qin@amd.com>
Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: cac014f173]
2022-08-13 00:16:14 -05:00
Liam Wrubleski
de7b2f75d9 Remove incorrect old packaging statement (#308)
[ROCm/composable_kernel commit: eccf8773a6]
2022-06-30 09:40:03 -05:00
Liam Wrubleski
ef57890712 Switch to standard ROCm packaging (#301)
* Switch to standard ROCm packaging

* Revert .gitignore changes

* install new rocm-cmake version

* update readme

Co-authored-by: illsilin <Illia.Silin@amd.com>
Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: b653c5eb2e]
2022-06-25 09:35:16 -05:00
Chao Liu
9df0a11a51 Absolute include path (#281)
* ad gelu and fast_gelu

* added GeLU and fast GeLU

* clean up

* add gemm+fastgelu example

* add gemm+gelu instances

* update profiler

* clean up

* clean up

* adding gemm+bias+activation

* clean

* adding bias

* clean

* adding gemm multiple d

* debugging

* add gemm bias add fastgelu

* rename, clean

* refactoring; add readme

* refactor

* refactor

* refactor

* refactor

* refactor

* refactor

* fix

* fix

* update example

* update example

* rename

* update example

* add ckProfiler

* clean

* clean

* clean

* clean

* add client app example

* update readme

* delete obselete files

* remove old client app

* delete old file

* cleaning

* clean

* remove half

* fix header path

* fix header path

* fix header path

* fix header path

* fix header path

* fix header path for all examples

* fix header path

* fix header path

* fix header path

* fix header path

* fix header path

* fix header path

* fix header path

* fix header path

* fix header path

* revert client app example

* clean build

* fix build

* temporary disable client test on Jenkins

* clean

* clean

* clean

[ROCm/composable_kernel commit: d1db6a0c3e]
2022-06-24 20:51:04 -05:00
Chao Liu
983972fedf remove options.hpp.in (#240)
[ROCm/composable_kernel commit: 44943e0e21]
2022-05-20 14:40:12 -05:00
Anthony Chang
3a574a0f5c Validate examples in CI (#233)
* validate examples in ctest runs

* format

* fix usage of check_err

* amend

* add example codes to custom target 'check'

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: 9f71ff48e2]
2022-05-13 16:54:44 -05:00
JD
569dd9f47b Add host API (#220)
* Add host API

* manually rebase on develop

* clean

* manually rebase on develop

* exclude tests from all target

* address review comments

* update client app name

* fix missing lib name

* clang-format update

* refactor

* refactor

* refactor

* refactor

* refactor

* fix test issue

* refactor

* refactor

* refactor

* upate cmake and readme

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: cec69bc3bc]
2022-05-12 09:21:01 -05:00
Adam Osewski
159494284d Introduce GoogleTest framework. (#204)
* Use googletest for tests. Add conv2d_fwd UT.

* Add conv1D/3D to gtest UT.

* Fix: not duplicate test with CTest.

* Convert more tests to googltests.

* Fix: GIT_SHALLOW is not allowed for git commit hash.

* Clang-format

* use integer value for GEMM test

Co-authored-by: Adam Osewski <aosewski@amd.com>
Co-authored-by: Chao Liu <chao.liu2@amd.com>
Co-authored-by: Chao Liu <lc.roy86@gmail.com>

[ROCm/composable_kernel commit: 8eca05a633]
2022-04-30 08:50:16 -05:00
Chao Liu
82ad74304e Reorganize files, Part 1 (#119)
* delete obselete files

* move files

* build

* update cmake

* update cmake

* fix build

* reorg examples

* update cmake for example and test

[ROCm/composable_kernel commit: 5d37d7bff4]
2022-03-08 21:46:36 -06:00
JD
019ea09acf Update test CMakeLists to add new tests automatically and add Jenkins stage for tests (#88)
* add docker file and make default target buildable

* add Jenkinsfile

* remove empty env block

* fix package stage

* remove render group from docker run

* clean up Jenkins file

* add cppcheck as dev dependency

* update cmake file

* Add profiler build stage

* add hip_version config file for reduction operator

* correct jenkins var name

* Build release instead of debug

* Update test CMakeLists.txt
reorg test dir
add test stage

* reduce compile threads to prevent compiler crash

* add optional debug stage, update second test

* remove old test target

* fix tests to return proper results and self review

* Fix package name and make test run without args

* change Dockerfile to ues rocm4.3.1

* remove parallelism from build

* Lower paralellism

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: 992f71e371]
2022-03-03 16:59:42 -06:00
JD
ba2852c7a3 Initial Setup for CI (#86)
* add docker file and make default target buildable

* add Jenkinsfile

* remove empty env block

* fix package stage

* remove render group from docker run

* clean up Jenkins file

* add cppcheck as dev dependency

* update cmake file

* Add profiler build stage

* add hip_version config file for reduction operator

* correct jenkins var name

* Build release instead of debug

* clean up

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: 2778e99758]
2022-02-18 21:44:11 -06:00
Chao Liu
8efcb80fa5 GEMM+Bias+ReLU+Add (#76)
* tweak conv for odd C

* update script

* clean up elementwise op

* fix build

* clean up

* added example for gemm+bias+relu+add

* added example for gemm+bias+relu

* add profiler for gemm_s_shuffle; re-org files

* add profiler

* fix build

* clean up

* clean up

* clean up

* fix build

[ROCm/composable_kernel commit: 823657ed12]
2022-02-06 22:32:47 -06:00
Chao Liu
7ead49ca42 added test for magic number division (#58)
[ROCm/composable_kernel commit: 237d4ca03f]
2021-11-30 09:09:28 -06:00
Chao Liu
b9f9ed96ac ckProfiler and device-level XDL GEMM operator (#48)
* add DeviceGemmXdl

* update script

* fix naming issue

* fix comment

* output HostTensorDescriptor

* rename

* padded GEMM for fwd v4r4r4 nhwc

* refactor

* refactor

* refactor

* adding ckProfiler

* adding ckProfiler

* refactor

* fix tuning parameter bug

* add more gemm instances

* add more fp16 GEMM instances

* fix profiler driver

* fix bug in tuning parameter

* add fp32 gemm instances

* small fix

* refactor

* rename

* refactor gemm profiler; adding DeviceConv and conv profiler

* refactor

* fix

* add conv profiler

* refactor

* adding more GEMM and Conv instance

* Create README.md

Add build instruction for ckProfiler

* Create README.md

Add Readme for gemm_xdl example

* Update README.md

Remove build instruction from top most folder

* Update README.md

* clean up

[ROCm/composable_kernel commit: e823d518cb]
2021-11-14 11:28:32 -06:00
Chao Liu
04d90a65c5 refactor
[ROCm/composable_kernel commit: 16effa767c]
2021-08-16 20:36:47 +00:00
Chao Liu
a0df0eb029 fix kernel filename
[ROCm/composable_kernel commit: 2c48039d0e]
2021-08-10 22:15:23 +00:00
Chao Liu
a15f7c025f fix clang warning suppression
[ROCm/composable_kernel commit: ddd49ec9e7]
2021-08-10 06:20:24 +00:00
Chao Liu
72c6d11a3b tidy
[ROCm/composable_kernel commit: 76f3131939]
2021-08-09 18:49:59 -05:00
Chao Liu
7b306ba01c tidy
[ROCm/composable_kernel commit: d18428901e]
2021-08-09 18:20:02 -05:00
Chao Liu
590dde14c8 tidy
[ROCm/composable_kernel commit: 56fc0842b3]
2021-08-09 19:27:49 +00:00
Chao Liu
193f7cbbcf tidy
[ROCm/composable_kernel commit: e62bae7a4a]
2021-08-09 15:11:35 +00:00
Chao Liu
71a28fc501 add tidy
[ROCm/composable_kernel commit: 24c8728942]
2021-08-08 17:41:54 +00:00
Chao Liu
b0ab2054a9 remove online compilation from CK
[ROCm/composable_kernel commit: ae98b52ad8]
2021-08-07 00:51:05 +00:00
Chao Liu
f94e566273 reorganize files to prepare for MIOpen integration (#51)
* change olc cmake

* adding online compile to fwd-v4r5r2

* update scripts

* remane fwd-v4r5r2 to fwd-v6r1

* clean up

[ROCm/composable_kernel commit: 1264925422]
2021-07-18 00:43:05 -05:00
Chao Liu
cbd5de28e1 Deprecate static kernel (#42)
* deprecate static kernels

[ROCm/composable_kernel commit: 81c942cd7e]
2021-07-08 10:40:00 -05:00
zjing14
67dcc552b6 xdlops_v4r4_fwd fp32/fp16 (#34)
* create files for xdlops

* working on blockwise_gemm_xdlops

* add KReduction

* add m/n repeats

* add 2x2 pipeline

* added 128x128 wavegemm

* use StaticBuffer of vector_type

* break vector type to blk_size

* add kpack into xldops_gemm and blockwise_gemm

* abroadcast only

* add fp32 mfma instructions

* adding fp16 mfma

* pack half4_t

* rename kperwave to kpack

* add 32x32x8fp16

* add fp16 mfma

* clean code

* clean code

* V4r4 xdlops kpack (#35)

* add kpack with incorrect results

* bug fix for make_dynamic_naive_tensor_descriptor_aligned_v2

* add 1x1 kernel

* add gridwise_gemm_v2 - single_buffer

* enabled dwordx4 for fp16

Co-authored-by: Chao Liu <chao.liu2@amd.com>

* refactor fwd-v4r4-xdlops

* add v4r4-nhwc-xdlop

* improve some perf of nhwc and nchw by tuning parameters, and change scheuduling in gridwise-gemm loop

* tweak scheduling in gridwise gemm

* add v4r3 with a single output copy

* init commit: output with slice win

* adding sliceWin

* add multiple repeats pattern

* starting adding bwd-v4r1-xdlops

* use tuple as SrcBuffer

* adding bwd-data v4r1 nhwc xdlops

* fix bug in make_dynamic_naive_tensor_descriptor_aligned_v2()

* fix bug in host bwd-data conv

* initial implementation of bwd-data v4r1 nhwc xdlops

* add launch bound flags

* enable launch bound

* add m/nrepeat=4

* tweak bwd-data v4r1 nhwc xdlops

* added bwd-data v4r1 nhwc xlops with output A and weight B

* add fwd-v4r4 nhwc xdlops, A input, B weight, C output

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: 3835318cc3]
2021-07-01 14:33:00 -05:00
Qianfeng
817b2a47c6 Add online compilation for dynamic kernels (#37)
* Add online-compiling facility

* Synchronize from fwd-v4r5 and implement host interfaces to call conv-fwd v4r4/v4r5 using on-line compiling method

* Tiny adjustment to time reporting

* Use object assignment to replace explicit bytes copying in the first kernel of v4r4/v4r5

* Use single thread to assign descriptor object to device memory

* Adjust to the workload assignment of the two kernels of v4r4 (experimental)

* Revert "Adjust to the workload assignment of the two kernels of v4r4 (experimental)"

This reverts commit eb38461456bb0c82b6c0d32cdd616e181907e20c.

* Update to make constexpr for generating descriptor types in kernel 2 of dynamic conv-fwd v4r4

* Update to dynamic conv-fwd v4r4 online-compiling

* Update to dynamic conv-fwd v4r5 online-compiling (result not accurate)

* Tiny update to driver/CMakeLists.txt

* clang-format

* Tiny comments change

* Add env OLC_DUMP_SAVE_TMP_DIR to support saving of temperary dir

* Fwd v4r5 olc perf (#39)

* added hip-clang flags that fix perf issue of online compilation

* fix bug for olc fwd-v4r5-nchw

* Move constexpr and type reference statements out of the function body in conv-fwd v4r4/v4r5 kernel wrapper

* Remove printing in hip_build_utils.cpp

* Update to root CMakeLists.txt

* Revert "Move constexpr and type reference statements out of the function body in conv-fwd v4r4/v4r5 kernel wrapper"

This reverts commit 3d2c5d8ecdd8298b72d127110500ed5b38d9835c.

Co-authored-by: Chao Liu <chao.liu2@amd.com>
Co-authored-by: Chao Liu <lc.roy86@gmail.com>
Co-authored-by: root <root@dc-smc-18.amd.com>

[ROCm/composable_kernel commit: 1685048a67]
2021-06-24 08:34:19 -05:00
Chao Liu
b4dbf677ce Dynamic tensor descriptor (#24)
* support dynamic tensor descriptor

* use buffer load OOB feature for padding case

* add navi support

* add int8x4 inference kernel

Co-authored-by: Chao Liu <chao@ixt-rack-81.local.lan>
Co-authored-by: Jing Zhang <jizhan@amd.com>

[ROCm/composable_kernel commit: fcbb978828]
2021-03-25 13:51:11 -05:00
Chao Liu
0eb214d1cd Code clean up (#20)
* tuning para,

* testing on v100

* add fp16

* remove deprecated tensor descriptor

* sync with miopen

* update build script

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

[ROCm/composable_kernel commit: 5c7cec1115]
2020-06-23 20:31:27 -05:00
Chao Liu
3799741fee backward data (#7)
* enabled atomic add in tensor copy
* added gridwise GEMM
* added backward data conv using GEMM + atomic
* added backward data conv using GEMM, no atomic


[ROCm/composable_kernel commit: 8f5f64960e]
2019-12-03 01:16:12 -06:00
Chao Liu
08f130fed1 Refactor for MIOpen integration (#4)
Refactor, so can bring multi-index transformation and padding support into MIOpen

[ROCm/composable_kernel commit: 52c3fe05be]
2019-10-11 11:37:31 -05:00
Chao Liu
611944080a refactor
[ROCm/composable_kernel commit: 21f7e9f103]
2019-06-19 17:43:56 -05:00
Chao Liu
5f217ebda5 reorginzed files
[ROCm/composable_kernel commit: 1566b31736]
2019-06-13 15:12:12 -05:00
Chao Liu
11c6b2ab9a change build
[ROCm/composable_kernel commit: c82b833d8e]
2019-06-12 10:47:25 -05:00
Chao Liu
14451df1fd reorginze files
[ROCm/composable_kernel commit: 81497a93a0]
2019-06-11 23:49:51 -05:00
Chao Liu
180f7266cf updated build
[ROCm/composable_kernel commit: fc60454e3a]
2019-02-15 02:17:00 -06:00
Chao Liu
153629655f update build
[ROCm/composable_kernel commit: a414e3fdf8]
2019-02-15 02:06:34 -06:00
Chao Liu
e7f6b820cd hip build
[ROCm/composable_kernel commit: 67c6f73ffe]
2019-02-15 00:54:30 -06:00
Chao Liu
662ba16dc6 refactor build, clean up
[ROCm/composable_kernel commit: e80fbbdd71]
2019-02-14 15:10:16 -06:00
Chao Liu
6521ccba67 cpu direct conv
[ROCm/composable_kernel commit: d51b81588f]
2018-10-19 01:26:21 -05:00
Chao Liu
4ef894a0b8 initial build
[ROCm/composable_kernel commit: 06c9f9fe17]
2018-10-14 02:10:36 -05:00
Chao Liu
8bfafec554 start adding convolution
[ROCm/composable_kernel commit: fc98757acd]
2018-10-08 22:49:58 -05:00