Commit Graph

1347 Commits

Author SHA1 Message Date
Sam Wu
34a8eebe34 Update doc reqs 2024-08-02 12:10:44 -06:00
Makarand Maydeo
bbf7375b77 Merge pull request #1398 from ROCm/asan_fix_6.2
fix ASAN target list
rocm-6.2.0
2024-07-19 18:59:05 -04:00
pramenku
a5cd58c927 Update CMakeLists.txt
removed GFX12 target
2024-07-18 18:57:05 +05:30
Illia Silin
f4701e5a6f fix ASAN target list 2024-07-17 20:50:53 -07:00
carlushuang
c6589826aa WA for rocm-6.2+ s constrait for buffer resource (#1346)
* WA for rocm-6.2+ s constrait for buffer resource

* add missing memory clobber
2024-06-27 10:06:40 -07:00
Bartłomiej Kocot
406baa534c Add read_first_lane function for int64 (#1347) 2024-06-27 10:06:40 -07:00
Bartłomiej Kocot
df06d26129 Fix continous dim selection in contraction (#1336)
* Fix continous dim selection in contraction

* Fixes
2024-06-27 10:06:40 -07:00
Bartłomiej Kocot
fd2feb1419 Support large tensors in grouped conv fwd (#1332)
* Support large tensors in grouped conv fwd

* Multi ABD fixes

* Fix calculate element space size
2024-06-27 10:06:40 -07:00
zjing14
042270388c disabled lds direct load inline asm (#1331) 2024-06-27 10:06:40 -07:00
Bartłomiej Kocot
7def3ad9ad Integrate universal gemm with conv forward (#1320)
* Integrate universal gemm with conv fwd

* Fix conv fwd wmma test

* Fix instances

* Remove direct load check
2024-06-27 10:06:40 -07:00
zjing14
1030c945e1 Remove gfx900 and gfx906 from default target device to reduce package size (#1351) 2024-06-27 10:05:39 -07:00
Jun Liu
17cf8179ee Merge branch 'amd-develop-0605' into amd-master 2024-06-05 15:47:50 -07:00
Jun Liu
e4112de730 Merge branch 'develop' into amd-develop 2024-05-22 13:16:50 -07:00
Bartłomiej Kocot
fd72380aeb Optimize grouped conv bwd weight for small M and N (#1303)
* Optimize grouped conv bwd weight for small M and N

* Fixes
2024-05-22 21:01:01 +02:00
Illia Silin
7b027d5643 Select appropriate GPU targets for instances, tests, and examples. (#1304)
* set individual gpu targets for instances, examples, tests

* fix path to hip compiler

* fix path to hip compiler once more

* aggregate device macros in ck_tile config header

* fix the cmake logic for instances

* fix clang format

* add gfx900 and gfx906 to default set of targets
2024-05-22 11:45:27 -07:00
Rostyslav Geyyer
204da9c522 Move grouped conv fwd client examples (#1299)
* Move grouped conv fwd client examples

* Update existing examples

* Format
2024-05-21 09:52:41 -05:00
Illia Silin
06b891c5c2 aggregate device macros in ck_tile config header (#1297) 2024-05-20 08:34:45 -07:00
Jun Liu
a6ef5c391e Merge branch 'develop' into amd-develop 2024-05-17 11:26:16 -07:00
Illia Silin
1274861a9d replace the ENV macro with CK_ENV (#1296) 2024-05-17 10:42:51 -07:00
dependabot[bot]
6637a810d0 Bump rocm-docs-core from 1.1.1 to 1.1.2 in /docs/sphinx (#1293)
Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 1.1.1 to 1.1.2.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v1.1.1...v1.1.2)

---
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>
2024-05-17 07:44:48 -07:00
rocking
aaa8dfdae9 Fix compile error (#1292)
error: no viable conversion from returned value of type '__half' to function return type 'fp16_hip_t' (aka '_Float16')

Co-authored-by: carlushuang <carlus.huang@amd.com>
2024-05-17 17:19:17 +08:00
Illia Silin
c44137838e remove wrong use of nonexistent class members (#1290) 2024-05-15 08:08:17 -07:00
carlushuang
dd0dd13d4e remove operator-deref (#1291) 2024-05-15 08:06:50 -07:00
jakpiase
3e3471d5d2 Add unit tests for grouped gemm two stage (#1256)
* add unit tests for grouped gemm two stage

* add reviewers suggestions

---------

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
2024-05-15 10:03:39 +02:00
Jun Liu
9b3c4ac475 Merge branch 'develop' into amd-develop 2024-05-14 17:06:50 -07:00
Illia Silin
7843a8a7fb re-enable convnd_fwd_xdl_fp64 testing (#1289) 2024-05-10 22:48:28 -07:00
Illia Silin
566b6480a2 Code clean-up (#1285)
* code clean-up

* remove the profiling output samples
2024-05-10 09:41:39 -07:00
carlushuang
fcba889ef4 [CK_TILE] fix some rand number init (#1287)
* add random norm

* normalized default to 0/3

* change squant->auto
2024-05-10 09:03:39 -07:00
Bartłomiej Kocot
8346af9c68 Change output gemm type to AccDataType in two stage conv bwd wei (#1283) 2024-05-10 10:57:42 +02:00
Adam Osewski
a0ae1c6133 Fix MakeArgument (#1284) 2024-05-09 09:42:41 -07:00
Adam Osewski
3c043cd10b Add vector instruction coherency bits for gfx94 targets. (#1268) 2024-05-09 07:30:17 -07:00
Illia Silin
fdbf8ccbd7 fix the output formatting (#1282) 2024-05-08 16:11:54 -07:00
Bartłomiej Kocot
0b6b5d1785 Add two stage grouped conv bwd weight kernel (#1280) 2024-05-08 09:53:24 +02:00
Illia Silin
bf42097646 Enable logging in CK with environment variable. (#1278)
* enable logging using environment variable

* update ck.hpp header

* fix typo

* fix clang format

* Update include/ck/utility/env.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

---------

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>
2024-05-07 16:26:43 -07:00
Jun Liu
1d784873ee Merge branch 'develop' into amd-develop 2024-05-07 14:10:48 -07:00
carlushuang
851c3ed157 [CK_TILE] support alibi (#1269)
* add alibi support

* fix code

* update code based on comment

* Support more hdim

* fix fp8 bias

* support seqlen_k=0 case

* remove unused printf

* fix format

---------

Co-authored-by: rocking <ChunYu.Lai@amd.com>
2024-05-07 22:32:54 +08:00
Sam Wu
6d073d31bb Add ROCm Doc team as codeowners for RTD yaml (#1277)
Also add component owners as codeowners for header directory
2024-05-06 10:07:39 -06:00
Illia Silin
08d51d9bc4 add missing vector header (#1275) 2024-05-02 11:27:59 -07:00
Illia Silin
d25889b12e Downgrade minimum required python version to 3.6 (#1274) 2024-05-01 15:35:31 -07:00
Illia Silin
7797f7c7a1 Downgrade minimum required python version to 3.6 (#1274) 2024-05-01 15:34:56 -07:00
Illia Silin
f0bf1e3125 [CI] Focus CI stages on MI200 nodes for resource optimization (#1273) 2024-05-01 10:07:14 -07:00
Rostyslav Geyyer
a2d0bdd5a9 Add an ignore (#1270) 2024-04-30 20:45:22 -07:00
Sam Wu
43579900a9 Update documentation requirements and configurations (#1272)
* Update documentation requirements

Set rocm-docs-core to v1.1.1

* Update RTD config

Set Python 3.10 for rocm-docs-core >= v1.0.0
2024-04-30 20:44:59 -07:00
Illia Silin
f6b3f4715d [CI][Tests] Add a daily cron job to build CK instances for gfx9;gfx10;gfx11. (#1271)
* add a daily build for instances for gfx9;gfx10;gfx11

* fix jenkins logic for instances only build

* fix the path for instance_only build

* reduce the number of build threads to 32
2024-04-30 14:44:30 -07:00
Adam Osewski
0f7e8ec485 Fix example CMakeLists.txt (#1267)
Add proper dependency target.
2024-04-30 08:28:19 -07:00
Rostyslav Geyyer
6ced3c12ff Mark unneeded instances as "getting deprecated" (#1265)
* Add a flag

* Add flag check and messages

---------

Co-authored-by: root <root@aus-g7-rogeyyer.amd.com>
2024-04-29 12:00:55 -07:00
Jun Liu
f0759faff4 Merge branch 'develop' into amd-develop 2024-04-26 12:28:00 -07:00
Haocong WANG
764164b488 [GEMM] UniversalGemm update (#1262)
* Add bf16 instances

* Add bf16 gemm universal example

* tempsave

* Add guard to navi compilation

* workground on a specific mixed gemm instance ( bring back it when compiler fix upload)

* fix formatting condition statement issue

* solve conflict

---------

Co-authored-by: Jun Liu <Liu.Jun@amd.com>
2024-04-26 12:56:07 -05:00
Rostyslav Geyyer
f044ff71fb Add element op (#1259) 2024-04-26 12:55:45 -05:00
zjing14
5ae893c0d3 ggemm tile_loop multD bf16 int8 (#1258)
* Overload output stream operator for LoopScheduler and PiplineVersion

* Add Run overload accepting grid descriptors MK.

* Add __device__ keyword for CalculateGridSize

* Create device op GroupedGemmMultipleD

* Add GroupedGemm MultipleD Tile Loop implementation.

* Add an example for GroupedGemm MultipleD tile loop.

* Device Op GroupedGEMMTileLoop.

* Bunch of small changes in exmaple.

* CkProfiler

* Remove unused tparam.

* changed the copy function to v7r2

* adding multi_abd

* in-progress

* add post-load oob check

* Fix include statement.

* Fix output stream overloads.

* Do not make descriptors and check validity untill we find group.

* Fix gemm desc initialization.

* debugging

* adjust instances

* add run_lds

* add elemntwise_op

* replace multi_abd_device with v3

* clean up

* clean

* clean

* Revert device op

* Fix compilation for DTYPES=FP16

* Validate tensor transfers paramters.

* Added LDSType

* profiling

* adjust oobcheck

* add missing file

* Validate on host only NK dims if M is not known.

* add

* clean

* refactor

* clean

* add examples

* add fuse

* add fusion and client example

* Fix bug.

* A convenient debug func for selecting threads.

* Fix has main k block loop bug.

* Make sure that b2c has up to date tile offset.

* Output stream operator for Sequence type.

* Cmake file formatting.

* clean

---------

Co-authored-by: Adam Osewski <Adam.Osewski@amd.com>
2024-04-26 10:37:49 -05:00