Adds a new ck_tile::runtime for runtime utilities.
Also moves functionality into the builder. Kernel is accessed with the
ckb::Kernel function that is templated by the builder class.
I've only verfied that the kernel compiles.
Some of my choices, like float32 types and having the epilogue set the member, are not valid template parameters. I now have this indentical to
a default GEMM universal kernel.
I also fixed some other small logical mistakes I made.
The code currently outputs the GetName results for some of the classes:
```
Kernel name: gemm_bf16_pipeline_AgBgCrCompV3_16x64x128_256_1x4_0x0x0
Shape: tile_gemm_shape_16x64x128x4_1x4x1_16x16x32
Problem: gemm_problem_256_0x0x0_Intrawave
Pipeline: pipeline_AgBgCrCompV3_16x64x128_256_1x4_0x0x0
```
* something khushbu can help with
* v1 v2 works with flatmm develop
* v0 v1 v2 numerical error gone
* Fixing numerical error, and interchange preshuffle configs to match with flatmm
* Refactor GEMM pipeline configurations and integrate preshuffle support
- Updated preshuffle pipeline definitions to include multiple versions (V1, V2, V3).
- Changed the pipeline constant from CK_TILE_PIPELINE_PRESHUFFLE to CK_TILE_PIPELINE_PRESHUFFLE_V3 in relevant configurations.
- Removed obsolete code and comments
* clang format
* fix vectorloadsize bug
* add the Preshuffle3
* update kwarp calculation in gemm utils
* update vector size A and B correctly in V2 pipeline; Added few more changes to align with dteng's branch
* fix: add CK_GFX950_SUPPORT macro for gfx950 detection
* default disable rotating buffer
* docs(CHANGELOG): update changelog for rocm 7.0
* Revert "docs(CHANGELOG): update changelog for rocm 7.0"
This reverts commit 2bc16fff84.
* Remove unused Preshuffle V3 pipeline and related code; update gemm function to use Preshuffle V2; clean up comments and formatting in various files.
* revert example/ck_tile/flatmm to its original state
* remove comment added by second author
* switch to xor ALDSDescriptor
* modify the MakeALdsDescriptor()
* temporary profiling script
* getting rid of line marker compiler error
* UniversalWeightPreshufflePipelineAgBgCrPolicy now derives from UniversalGemmBasePolicy
* add a minor fix for the config
* typo fix
* Fix formatting in lambda function for WeightPreshufflePipelineAGmemBGmemCRegV2
* revert change in include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp
* revert change in include/ck_tile/core/arch/amd_buffer_addressing.hpp
* reenable the GemmSpatiallyLocalTilePartitioner
* make GemmConfigPreshuffle_1 for v1 pipeline, GemmConfigPreshuffle_2 for v2 pipeline
* remove hardcoded true for preshuffle bool template argument
* rename script
* remove gemm_profilie.sh script
* merge conflict resolve
* clang formatted
* typo fix
* Remove duplicate include of block_gemm_areg_bsmem_creg_v2r1.hpp in gemm.hpp
* Remove commented-out code in UniversalWeightPreshufflePipelineAgBgCrPolicy
* Fix missing newline at end of file in run_gemm_example.inc
* Remove unused barrier call in BlockWeightPreshuffleASmemBSmemCRegV1
* addressing review comments
* removing debug code
* addressing review comments
* Revert "addressing review comments"
This reverts commit 29c45192ba.
* updating tile_engine code
* addressing review comments
---------
Co-authored-by: amd-khushbu <khuagarw@amd.com>
Co-authored-by: ThomasNing <thomas.ning@amd.com>
* Restrict building of gemm_universal_preshuffle_f8 instances to specific targets in CMakeLists.txt
* Add condition to skip gemm_xdl_universal_preshuffle_f8 instances for unsupported targets in CMakeLists.txt
* Add conditions to skip unsupported targets for gemm_universal_preshuffle_f8 and gemm_xdl_universal_preshuffle_f8 instances in CMakeLists.txt
* Refine conditions to exclude gemm_universal_preshuffle_f8 instances for unsupported targets in CMakeLists.txt
---------
Co-authored-by: AviralGoelAMD <aviralgoel@amd.com>
* Split-K autodeduction for DeviceGroupedConvBwdWeight_Xdl_CShuffle and DeviceGroupedConvBwdWeight_Xdl_CShuffleV3.
* Split-K autodeduction for DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle.
* Use simple best occupancy model to calculate the split-K.
* Handle split-K autodeduction in explicit gemm conv.
* Add unit tests for split-K autodeduction.
* Remove oversubscription.
* Small fixes.
* Added split-K autodeduction for DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle.
* Run clang formatting.
* Fix error handling in the conv profiler.
* Add missing documentation for the autodeducted split-K values.
* Add split-K autodeduction to DeviceGroupedConvBwdWeight_Explicit_Xdl solver.
* Fix clang formatting and split-K profiler documentation.
* Rename max_occupancy value variable.
* Calculate grid size for split-K autodeduction directly from input array shapes and template params.
---------
Co-authored-by: Ville Pietilä <>
* Add tests for host convesion f32/f16 to f8
* Add tests for host convesion from f8 to f32/f16
* Fix UB and corner cases in f32/f16 to/from f8 conversion
* There are UBs when very small values are converted to f8: bitshifts
can be larger that type width. Using unsigned long long does not help
because exponent_diff >= 64 in such cases. This causes that values
like 2.117582368e-22 are converted to non-zero f8 in host validation
of FMHA tests, test_f8 crashes with segfault in completely irrelevant
code like GTest internals or produces non-deterministic results etc.
* Fix FNUZ conversion to return NaN for NaN inputs.
* Fix compilation error (due to uint8_t << 8) in OCP e5m2 to f16
conversion.
* Replace some magic numbers with values from numeric_traits
* Build tests only on devices supporting the type
This is a prototype for building a GEMM using templates and C++20 features, especially concepts and consteval.
This PR only creates a dummy GEMM class, but tests some of the design patterns.
It also includes a utilty to alloc host buffers with RAII.
* add a dummy test file
* add kernel launch logic to the test
* transfer all test cases into gtest params
* factor kernel out into test config
* add load transpose pipeline tests
* add padded tests and skip invalid kernels at runtime
* enum class for pipeline type
* add multiwarp test cases
* fix type
* try to solve the problem
---------
Co-authored-by: ThomasNing <thomas.ning@amd.com>
* update gpu_timer for rotating buffer as hipblasLt's implementation
* timing fix
* Updating gpu timer for old ck as well
* Revert "Updating gpu timer for old ck as well"
This reverts commit 958cd1bc99.
* code clean up with runtime argument; function rename
* code cleanup
* general timer fixes
* bug fix
* clang formatted
* addressing reveiew comments
* clang formatted
* Addressing review comments
* CI fix
---------
Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
* Expand the bandwidth of direct_global_to_lds for gfx950
* clang-format
* fix the remod.py and script for clang format
---------
Co-authored-by: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>
* uncomment all the headdim, use optdim to control
* change default back to -1
* uncomment splitkv instance
* Fix typo in receipt 4 for appendkv
* support optdim for bwd, splitkv and appendkv
* Fix 192 key error
---------
Co-authored-by: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>
Co-authored-by: Andy Lugo <Andy.LugoReyes@amd.com>
* unify pipeline signature with existing example
* iwyu
* move stuff around in load-tile-transpose
* cleanups in batched transpose pipeline
* comments
* use same inputs size
* cleaner printf
* print host args
* use 64 block sides in the 37_transpose example
* roll back grid dimension size adjustment for 37_transpose example
* transpose grid for 37_transpose to unify with 35_batched_transpose
* unify grid computation logic
* make policy methods device only (since they are used only on device from the pipeline)
* more host/device attribute cleanups
* copy over problem
* move over pipeline and policy
* add switch to batched transpose api
* make the lds problem more similar to original problem
* factor out logic into traits
* factor out conditional compilation into trait parameter
* propagate pipeline to args
* unhardcode pipeline dispatch parameter
* refactor vector size
* put warp tile out of dispatch
* rename template parameter for trait
* rewrite vector size in terms of problem
* mark policy-internal struct variable as device
* factor out input distribution and thread access pattern from policies
* reword vector size
* use datatype across batched transpose pipelines, problems and kernel
* remove transpose traits from lds pipeline
* add padding to the lds pipeline *interface*
* add comment
* remove ck_tile example #37
* update cmakelists
* add test for new pipeline
* update batched transpose test
* roll back load_tile_transpose changes
* remove comments
* pack dispatch parameters into a config
* padM can be enabled
* adjust lds vector size to enable padding along N
* update test
* clean up logic
* swap m/n input vector size
* adjust perf test script
* sweep over C/W in perf test
* count both read and written bytes into bandwidth (x2 the number)
* clang-format
* widen size range for perf test
* remove 64k x 64k case; it's too large for index
* remove thread tile from dispatch
* Solve merge conflict
* fix compile
* modify the transpose
* solve the test error and clang format
* Add v3 support for Groupd fwd conv+bias+clamp & ckProfiler (#2463)
* Add logging to IsSupported.
* Less casting in AddClamp
* Conv+bias+clamp instances & profiler BF16
* Fix 3D instances & run just 1x for verification.
* :Run just once for verification conv fwd.
* ckProfiler conv fwd clampwq
* Remove exec bit & formatting
* Add support for MultiD for grouped conv fwd v3.
* Enable 2Lds.
* clean
* align instances
* align instances
* profiler fixes
* Fixes
* fix
* fix
---------
Co-authored-by: Adam Osewski <root@quanta-ccs-aus-f01-19.cs-aus.dcgpu>
Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>
* Fixing 0ms and inf GB/s issue in img2col (#2565)
issue :
====
``` sh
$ bin/tile_example_img2col
Perf: 0 ms, inf GB/s
```
solution :
======
Problem occured because config.time_kernel is false by default.
if false, then no need to calculate perf, just print proper message
`image_to_coloumn: pass, No Perf generated due to config.time_kernel=0`
* merge with develop
* solve clang format
---------
Co-authored-by: ThomasNing <thomas.ning@amd.com>
Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
Co-authored-by: Adam Osewski <root@quanta-ccs-aus-f01-19.cs-aus.dcgpu>
Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>
Co-authored-by: rahjain-amd <Rahul.Jain@amd.com>
issue :
====
``` sh
$ bin/tile_example_img2col
Perf: 0 ms, inf GB/s
```
solution :
======
Problem occured because config.time_kernel is false by default.
if false, then no need to calculate perf, just print proper message
`image_to_coloumn: pass, No Perf generated due to config.time_kernel=0`
