* memory op changes
* memory op changes
* Fixing TILE_ENGINE_BASIC in Tile Engine
* Removing gfx90a from Tile Engine Run
* [CK TILE ENGINE] increasing ci configs for BASIC case
* Setting RUN_TILE_ENGINE_BASIC_TESTS to ON by default
---------
Co-authored-by: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>
* solve compiler issue
* solve the gfx950 mfma shuffle regression
* refactor jenkinsfile to handle arch name better
* [CK TILE] set divisor to count of thread along k dimension
* fix the compiler error
* solve degradation
* Finish the multiplies fix
* fix the scales
* solve compilation error
* solve the composes
* solve the error of tile sweeper
* fix the test and example
* fix for gfx950
---------
Co-authored-by: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>
Co-authored-by: illsilin_amdeng <Illia.Silin@amd.com>
Co-authored-by: Cong Ma <congma13@amd.com>
* Removing hard-coded trace filename
* Including stage name in notification
* Simplifying capture setup and tagging file names with arch
* Removed test property from notification message
* Fixing regex to get arch name
* Fixing error in notification and modified regex
* generate and visualize build traces for all archs
* generate build traces in all cases
* fix jenkins logic
* fix typo
* use more threads for parsing dependency map
* add script to parse ninja traces and issue warnings
* fix python script syntax and header
* fix python syntax one more time
* fix python syntax
* Parallelization in dataset generation
* Parallelizable tests for fwd, bwd data, bwd weight with datasets
* .gitignore generated datasets
* Test parallelization script with round-robin GPU scheduling
* Parallelization updates to test generation and running
* Dataset paths relative to executable
* Update output from test generation
* Default to one GPU in test generation
* Add small dataset tests to Jenkins
* Update copyright lines
* Update test_data/generate_test_dataset.sh
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
* Move trap disable
* Common get path function
---------
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
* [CK TILE STREAMK] Introduce initial support for tile engine in streamk GEMM.
- This commit lays the groundwork for integrating the tile engine into streamk GEMM.
It focuses on creating benchmark executables for streamk GEMM.
- Additional scripts like test_benchmark.sh and gemm_benchmark.py will be added once
the streamk implementation reaches stability.
* [CK TILE STREAMK] Enable CI to execute tile engine benchmarks for StreamK GEMM
* [CK TILE STREAMK] Refactor: Extract common utility functions.
* [CK TILE STREAMK] Revise tile engine of streamk to align with the updated implementation
* Add pre-commit
* [CK TILE STREAMK] Add 'dp_persistent' and 'reduction_strategy' in output of CK TILE STREAMK
* [CK TILE STREAMK] Fix a bug about value of 'dp_persistent' of CK TILE STREAMK
* [CK TILE STREAMK] Update Jenkinsfile
* [CK TILE Engine] Update StreamK tile engine help message
Remove default value messages as they are automatically printed
* [CK TILE Engine] Update StreamK tile engine
- Remove namespace reboot
* [CK TILE Engine] Update StreamK tile engine
- Fix merge error
* build and run ck_builder tests
* add test_ckb_all to targets
* fix syntax
* fix test path
* Update CMake targets for builder testing in CI (#3290)
Our existing CMake only had build targets. Update CMakeLists.txt to have CTEST targets:
* smoke-builder
* regression-builder
* check-builder
Co-authored-by: John Shumway <jshumway@amd.com>
* use check-builder target
* get rid of test_ckb_all target
* call ninja check-builder separately
---------
Co-authored-by: John Shumway <jshumway@amd.com>
* allow using alternative compiler in all CI stages
* get rid of some redundancies in jenkinsfile
* clean up jenkinsfile a bit more
* further clean up jenkinsfile
* do not force user jenkins in ci dockers
* Renaming old code
* Adding GEMM code with new Architecture
* Partial Progress : Errors
* Partial Progress : Working code
* Changes to element wise function
* Removing Debugging statements
* Working GEMM Multi D code
* Removing Stale Code
* Address Copilot review comments
* Address Copilot review comments
* Changes to validation file
* Changes to common code snippets
* Creating common folder
* Removing duplicate files
* Pointing to right common file
* Pointing to right common file
* Pointing to right common file
* Changing to VERBOSE
* Changing CMAKE messages to verbose
* Updating Cmake with right layout datatype configs
* Working code for GEMM Multi D
* Adding GPU not found pattern
Also, failurePatterns does not need to be global. Moved variable to live in the failure notifications function scope.
* Testing new failure type
* Testing failure
* Removing the forced failure test
* Adding an additional failure pattern
* Pass hdim to tile_example_fmha_fwd in fp8 tests
* Add WMMA support to fwd FMHA pipelines
* Tune tile sizes a bit for less spilling
fp16 256 is still quite slow
* Fix Q grad tile distribution for warp size = 32 and hdim >= 256
With AccDataType = float and warp size = 32, K0 becomes 0, K repeat is required to correcty distribute the tile.
* Use code based on BlockDropout in BlockDropoutBwd
* Fix split KV combine kernel for gfx12 (warp size 32) and make it more universal
* Fix LSE LDS tensor descriptors: kMaxSplits and kM0 were swapped, it worked on gfx9
because they both equal to 8 while on gfx12 they are 8 and 4;
* Fix Oacc LDS tensor descriptor: it was transposed even though its shape=[4 * kM0, kN1],
it worked on gfx9 because 4 * kM == kN1 == 32;
* Removing these hidden dependecies allows to support:
* any number of warps (power-of-2), not only 4;
* kN1 = 16, not only 32;
* any number of splits;
* Rename ids like o_acc_4 and Oacc4 to eliminate confusion: kNumWarps doesn't have to be 4 now
* Replace hard-coded kN1 in dispatch code with the requested tile size
* Add gfx12-specific tile sizes for split KV
* Pass GPU architecture to kernel generation scripts
This is still a temporary solution.
* Build and run FMHA CI tests for gfx12
* Fix issue after merging
* Fix bwd tile sizes
The current pipelines always read only one tile K and V tile, this
requires bk0 == bhdq and bk2 == bhdv (kK0 == kQKHeaddim and
kK2 == kVHeaddim).
* Use hardware f32->f8 on gfx12, remove v_perm
__builtin_amdgcn_perm is not needed because
__builtin_amdgcn_cvt_pk_fp8_f32 allows to specify which word (16 bit of
32-bit dword) is used to store results (two f8 values).
* Update changelog
* Add WMMA support to pagedkv
* Fix scripts after rebasing
* Support 16x16 (MFMA, WMMA) and 32x32 (MFMA) tiles in fwd and bwd BlockDropout
Add comments with dropout implementation details
Fix performance regression of fwd+dropout
* Remove some usage of type punning (reinterpret_cast with ref or ptr) in Philox;
* "scalarize" seed and offset, they may come either from kernel args or from device memory
(presumably loaded with vector loads).
These changes help the compiler to procude more optimal code and reduce register spilling.
Use WarpGemmDispatcher instead of explicit WarpGemmMfma... to get CWarpDstrEncoding
Use code based on BlockDropout in BlockDropoutBwd
Refactor BlockDropout (fwd)
Implement BlockDropout (fwd) for WMMA
Originally BlockDropout only supported 32x32 tiles (IsWG32 = true),
this version supports 16x16 tiles.
If MPerBlock > MWarp * 16, it can generate numbers for two 16x16 tiles, similarly
to BlockDropoutBwd.
Implement BlockDropoutBwd for WMMA
Remove MakeRandValLds* functions unused in BlockDropoutBwd
Remove unused Run overload from BlockDropoutBwd
* Fix regression with philox seed and offset when they exceed 32-bit int
__builtin_amdgcn_readfirstlane works with 32-bit values, seed and offset
are 64-bit so they get truncated.
* Fix names after cherry-picking
* Fix selection of a fallback tile based on bm0
The assumption that the largest bm0 == 128 is not always true for
current fp32 tiles.
* Do not use filters related to qr_async_trload
They disable tiles/pipelines which are valid for gfx12.
* Use different dstr encoding when C is transposed
* Do not call GetQKBlockGemm (and hence WarpGemmDispatcher) in host code
Some WarpGemmDispatcher instantiations are defined only
for specific archs and undefined on host.
Calculations related to sched barriers are moved from Pipeline's public
fields into pipeline's operator().
* Fix incorrect name WarpGemmMfmaFp8Fp8F32M32N32K16SwizzleBTransposedCDistribution
Correct name is WarpGemmMfmaFp8Fp8F32M32N32K32SwizzleBTransposedCDistribution
because it's 32x32x16 with IterateK = 2 so K = 32, also all tiles used
in codegen scripts are 32, 32, 32.
* Generalize usages of WarpGemmDispatcher for MFMA and WMMA
WarpGemmMfmaFp8Fp8F32M32N32K32SwizzleBTransposedCDistribution is still
used explicitly becaus of swizzle factor = 4.
* Mark has_load_tr as maybe_unused
There are no transpose loading for RDNA.
* Remove CK_TILE_USE_MFMA/WMMA from fmha-related code
* Detect BlockSize on host based on warp size of the current device
If kBlockSize == kNumWarps * get_warp_size(), the kernel is launched with
kBlockSize / 2 because on host get_warp_size() == 64 always.
* Fix calculation of grid size for combine kernel with warp size = 32
* Add missing includes and header
* Support multiple archs in one binary for fwd
* Support multiple archs in one binary for fwd_splitkv, fwd_appendkv, pagedkv_prefill
* Support multiple archs in one binary for bwd
* trload kernels are compiled only for gfx950;
* instances with padding are checked after instances without padding so
they can be used as fallbacks (similarly to fwd);
* Extract common code from register_traits
* Revert "Fix regression with philox seed and offset when they exceed 32-bit int"
To simplify merging , the proper fix is in develop already.
* Support new numerical d paddings in trait ordering checks
* Build fp32 tests only on gfx9
* Do not use hardcoded M0 = 64 for dot bwd kernel
* Use textwrap.indent from standard library
* Make fp8 pipelines on gfx12 consistent with gfx9
* Update tests for current pipelines
* Make ninja check more responsive in CI
ninja buffers output so this job looks hanging.
* Support fp8fp32 by limiting O vector size
The fp32 output type requires storing 8 * sizeof(float) = 32 bytes,
which is not implemented (here 8 is the number of C values per lane for
v_wmma_f32_16x16x16...).
* Remove unused cmake options
* Unify including amd_buffer_addressing.hpp/_builtins.hpp
* Temporarily use amd_buffer_addressing.hpp on >=gfx10
amd_buffer_addressing_builtins.hpp uses inline asm for loads/stores
which is not compatible with >=gfx10:
* 1 scalar for exec masks instead of 2,
* gfx12 uses different instruction names etc.
* Update asm in bf16 conversions to work with warp 32
* Do not generate splitkv/appendkv with vlayout=col for consistency with fwd
* Add arch tags to kernels/host funcs, compile for each arch separately
* Add kM0 to fmha_bwd_dot_do_o kernel name to match filename
* Add workaround for miscompilation of bwd with padded hdim
SWDEV-559729: v_wmma instructions can be incorrectly placed in divergent
branches used to store padded tensors (when some lanes are inactive due
to padding). Inline asm with dummy dependencies on VGPRs of the tensors
prevents the compiler doing this.
* Fix add_gtest_executable for absolute paths
Some tests (like gemm_tile_engine) pass absolute paths to source files.
In CI the branch name is a part of the root dir, and if the branch name
contains "wmma", "xdl" etc., files can be incorrectly excluded.
* Run only hdim 128 smoke tests for fp8fp32
There are no instances for hdim 64 and 256.
* Format py with ruff to simplify merging develop
* Fix incorrect var name
* Codegen for gfx9,gfx950 when --targets is not specified
Aiter and Pytorch require changes for passing their targets to the codegen scripts.
With this temporary solution the files are generated but not all of them
have to be really built (depending on the used --offload-arch=).
* Combine arch-related values into ArchTrait
This more centralized approach removes duplication of various formatting templates.
* Try a workaround for Jenkins error "groovyjarjarasm.asm.MethodTooLargeException: Method too large"
Some code is extracted into a function.
* refactor package generation, add dedicated switch
* allow building packages not only on gfx9
* enable last stage to post packages
* stash packages from different arch into separate stashes
* build packages daily automatically
* Partial Progress : CK Tile Engine GEMM
* Partial Progress : CK Tile Engine GEMM
* Partial Progress : Working GEMM Code
* Partial Progress : Working GEMM Code
* Changinf jenkins to remove preshuffle
* Partial Progress : CK TILE ENGINE GEMM Debugging
* Partial Progress : Removing changes that are not GEMM
* Partial Progress : Validation of full block size in GEMM
* Changes in Jenkins to run only fp16 and bf16
* Addressing Review Comments
* Partial Progress : Addressing CI issues
* Partial Progress - Runing GEMM for fp16,bf16 and rcr
* Clang
* Adding fp8 and bf8
* Adding fp8 and bf8
* Adding additional architrcture
* Limited datatypes and layouts
* Adding k_block_per_cu in test config
* Changes to faling CI errors
* Changes to faling CI errors
* Validation for GEMM
* Adding Layout support
* Adding Validations
* Adding layout in jenkins
* Update on Jenkins
* Distribution validation for GEMM
* Resolving merge conflicts
* Solving merge conflicts
* Fixing check for changed files
* Testing CI skip behavior
* Testing CI Trigger
This should skip CI
---------
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
* Update Jenkinsfile
Adding logic to skip CI checks when a commit contains changes to non-relevant files like docs, .md, licenses, and .github workflow files.
* Update Jenkinsfile
* Update Jenkinsfile
* Update Jenkinsfile
Testing skip env var
* Update Jenkinsfile
Fixing syntax
* Update Jenkinsfile
Simplifying CI check logic
* Update Jenkinsfile
Testing skipping logic on stages.
* Update Jenkinsfile
Removing post block. The status for skipped stages are already reported.
* Testing Docs
Testing modifications to files in the docs folder do not trigger a the build and test stages.
* Testing Multifile Trigger
Removed Jenkinsfile from the skip patterns. Reversed change to docs file. This test should not skip CI checks.
* Clean code
Renamed setup stage to be more descriptive.
Added pipeline env variable for consistency.
Moved performance test results stage conditional up a level so the parent stage appropriate reports the status if it is skipped.
* Fixing syntax error
* Updated CRON Flags
Added the FORCE_CI flag to the CRON instructions. This will ensure CI does not skip the job.
* Updating logging
Making logs more explicit.
* Comment update
Cleaning comments.
* Update Jenkinsfile
Reverting performance reports when condition.
* Parallel Test
Testing stage status with parallel stages
* Update Jenkinsfile
* Update Jenkinsfile
Removing stages for quick testing
* Update Jenkinsfile
* Testing skipped parallel stages
Testing the addition of a coordination stage to always pass and give an update to skipped parent stages with parallel sub-stages.
* Testing parallel stages
Adding coordination stage to test if parent check status is correctly updated.
* Simplified performance results stage
Removed parent stage as there are no other parallel stages to execute (yet).
* Testing final clean up stage
* Testing check status update
Testing - forcing status to update after a stage skip.
* Testing results stage skip
* Removing test stage
* Testing pipeline
* Testing post status updates
* Process Test Results Post Event Update
The stage will report success when it skips or is successful.
* Testing non-relevant file change
This should skip build and test in CI
* Reverting test
updating regex file patterns to use strings instead of regex literal syntax.
* Fixing file matching regex
* Testing docs modification
* Fixing default env var value
* Correcting env var assignment
* Pipeline test
Updating docs file. Should skip ci.
* Testing Pipeline
Setting default run ci state.
* Adding debugging
* Removing debugging
* Pipeline test
Should skip pipeline
* Pipeline Test
Mixed files to trigger a CI run
* Adding additional status updates
The parent stage sometimes remains in pending even if the child stage completes when skipped. Added an additional status update for the parent stage.
* Fixing variable name
* Moving stage names
Moved the performance stage names to a single location because they are referenced multiple times. This reduces errors with typos in the future.
* Revert "Moving stage names"
This reverts commit 7cf6743e54.
* Update Jenkinsfile
Handle both truly empty arrays and arrays containing only empty strings.
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
---------
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
This change ensures that the files being selected for clang format validation are exactly the ones tracked by the git repo we are testing. This protects against an known issue where the repo being tested contained "stray files" from a previous test.
The prefix is causing the status updates from
gitStatusWrapper to be unique to the status updates that
are created by the Jenkins server, which creates duplicates
* upgrade default docker to rocm7.0.1
* turn on build and test on gfx950 by default
* use rocm-dev instead of rocm
* link libhiprtc for codegen targets
* resolving codegen compilation errors: removed calls to other std functions, resolved issues with int32_t: needed the correct header, put use of e8m0 into header guards
---------
Co-authored-by: Astha Rai <astha.rai713@gmail.com>
* initial work on adding support of gfx12 in tile_engine for GEMM benchmarking
* add stage("Run TILE_ENGINE_GEMM Tests on gfx1201") to Jenkins config
* make tile_[m/n/k] validation arch dependent
* enable aiter and fmha test stages on gfx950
* use newer compiler for gfx950
* make sure gfx950 runs correct docker
* fix typo
* upgrade base docker for aiter
* change base docker for aiter tests
* do not add group render to ck_aiter image
* add group irc in ck_aiter docker
* do not fix the irc group id to 39
* do not set jenkins uid and gid
* skip group irc for aiter tests
* fix syntax error in dockerfile
* change the base docker for aiter tests
* add irc group back to ck_aiter docker
* Improve random number generation
* use different seed for each input (Q, K, V...);
* use deterministic generation of:
* seqstart_q/k (for group mode);
* block_table (for paged-kvcahe);
* cache_batch_idx (for kvcache);
* Extract arg_parser-related code from run functions to use them as tests
* Split examples into main programs and fmha runners, build instances separately
* Add dummy tests that use instances and runners
* Fix a missed corner case of f32->f8 conversion
When value if < min f8 denormal but > min f8 denormal / 2, it must be
rounded to min f8 denormal (i.e. 0b1), not to 0.
* Fix incorrect fp8 scales for P and O in validation code
DataTypeConfig was incorrectly compared with fp8_t.
* Add host generation of dropout random values and use it for validation
Previously host validation (reference_batched_dropout) used random
numbers generated by BlockDropout of the kernel, meaning that incorrect
generation on device (bad distribution, repeated numbers, too many zeros,
etc.) would not trigger any validation errors.
* Implement tests from smoke_test_bwd.sh
* Return result as enum to distinguish failure and missing instance
* Add tests for bwd features: bias, alibi, dropout
* Implement tests from smoke_test_fwd.sh
* Pass seqlen_q/k as vectors to fwd and bwd runners
* Add tests for fwd features: bias, alibi, dropout
* Add tests for pagedkv and splitkv
* Fix conditions when to use splitkv and pagedkv kernels
splitkv was executed only when use_kvcache which == (need_append_kvcache || use_cache_batch_idx || 0 < page_block_size).
In the SplitKV tests: the regular fwd kernel was executed if use_cache_batch_idx was not requested even when num_splitkv > 1.
In the AppendKV tests: the pagedkv kernel was executed but it often failed to find an instance.
* Add tests for appendkv
* Use is_v_rowmajor = true because there are no instances with column layout anymore
* Split public and private compile options for instances
Tests and examples need to know only about CK_TILE_FMHA_FWD_*_API.
* Improve parsing validation in bias and mask
* Pass bias as string for consistency with mask
* Catch parsing and other exceptions
* Add bwd test for deterministic flag
* Initialize fp8 tensors (-init=ufq) similarly to uf
* Fix splitkv/pagedkv invocation: use padded sk when seqlen_k_ptr is not null
seqlen_k cannot be used to determine padding when seqlen_k_ptr is
provided. The actual seqlen_k is taken from seqlen_k_ptr[b].
Even seqlen_k values (% bn0 == 0) use padded seqlen_k while seqlen_k_ptr
may contain arbitrary values.
In the example or tests this produces incorrect results with appendkv
(for example, -d=32 -s=1 -s_k=64 -s_knew=7 -vlayout=c -b=8).
* Fix use_pagedkv value when kvcache = true but page_block_size = 0
In this case block_table_ptr is nullptr which is accessed in the kernel.
* Clean up bwd tests
* Unify fwd tests for f16/bf16 and fp8
* Use better explicit instantiation declaration for fmha_bwd<2>
* Use the same seed for all tests, allow to override it with env variable
* Undo clang-format of one irrelevant file
For some reason my local clang-format-18 and the one in CI work differently.
* Do not build instances and tests on unsupported archs
* Build instance libraries as OBJECT library
* CI: Enable sccache for HIP
There are source files with LANGUAGE HIP, they need
-DCMAKE_HIP_COMPILER_LAUNCHER=sccache
* Add tests to REGRESSION_TESTS
* Fix OOB accesses in deterministic bwd due to incorrectly assumed kN0
The runner assumes kN0 = (hdim_q <= 128) ? 128 : 64 but there are
smaller tiles (for tr_load or fp32). This can create too small dq_acc_buf.
* Pass CK_TILE_FMHA_FWD_*_API as INTERFACE compile options
The instances don't actually depend on them, only examples and tests do.
Passing these definitions as INTERFACE allows to change FMHA_FWD_ENABLE_APIS
without recompiling instances that are already in ccache.
* Fix formatting and names
