* [CK TILE] Fix bugs in AQuant preshuffle
- Make Aquant works with block Mx64x256. `M` could be 16, 32, 64
- Make Aquant works with warp 16x16x32 and 32x32x16.
* [CK TILE] Rename Preshuffle to PreshuffleQuant
The new name, PreshuffleQuant, explicitly states the function's purpose:
to preshuffle the quantization matrix.
* [CK TILE Block Scale] Use GemmConfig to save tile properties
- Remove specialization of GemmQuantTypeConfig
- Pass GemmConfig around which contains tile properties. Stop using hard
coded tile properties in `gemm_calc_aquant()`
* [CK TILE Block Scale] Rename GemmConfig used in block scale
- Remove unused GemmConfig
- Rename GemmConfig used in block scale
---------
Co-authored-by: ThomasNing <thomas.ning@amd.com>
[ROCm/composable_kernel commit: 245467f359]
* Add CSV-driven convolution test pipeline
- Add test_grouped_convnd_fwd_dataset_xdl.cpp with CSV reader functionality
- Add complete dataset generation toolchain in test_data/
- Add Jenkins integration with RUN_CONV_COMPREHENSIVE_DATASET parameter
- Ready for comprehensive convolution testing with scalable datasets
* Update convolution test dataset generation pipeline
* add 2d, 3d dataset csv files
* Remove CSV test dataset files from repository
* Update generate_test_dataset.sh
* Fix channel division for MIOpen to CK conversion
* Remove unnecessary test files
* Fix clang-format-18 formatting issues
* TEST: Enable comprehensive dataset tests by default
* Fix test_data path in Jenkins - build runs from build directory
* Add Python dependencies and debug output for CSV generation
* Remove Python package installation - not needed
* Add better debugging for generate_test_dataset.sh execution
* Fix Jenkinsfile syntax error - escape dollar signs
* Add PyTorch to Docker image for convolution test dataset generation
- Install PyTorch CPU version for lightweight model execution
- Fixes Jenkins CI failures where CSV files were empty due to missing PyTorch
- Model generation scripts require PyTorch to extract convolution parameters
* Add debugging to understand Jenkins directory structure and CSV file status
- Print current working directory
- List CSV files in test_data directory
- Show line counts of CSV files
- Will help diagnose why tests fail in Jenkins
* Fix clang-format-18 formatting issues
- Applied clang-format-18 to test file
- Fixed brace placement and whitespace issues
* Add detailed debugging for CSV dataset investigation
- Check generated_datasets directory contents
- List all CSV files with line counts
- Show first 5 lines of main CSV file
- Applied clang-format-18 formatting
- This will help identify why CSV files are empty in Jenkins
* keep testing add pytorch installation in shell script
* Use virtual environment for PyTorch installation
- Jenkins user doesn't have permission to write to /.local
- Create virtual environment in current directory (./pytorch_venv)
- Install PyTorch in virtual environment to avoid permission issues
- Use PYTHON_CMD variable to run all Python scripts with correct interpreter
- Virtual environment will be reused if it already exists
* Remove debug code and reduce verbose logging in Jenkins
- Remove bash -x and debug commands from Jenkinsfile execute_args
- Remove all debug system() calls and getcwd from C++ test file
- Remove unistd.h include that was only needed for getcwd
- Remove debug print in CSV parser
- Add set +x to generate_test_dataset.sh to disable command echo
- Redirect Python script stdout to /dev/null for cleaner output
This makes Jenkins logs much cleaner while still showing progress messages.
* install gpu torch
* Clean up and optimize comprehensive dataset test pipeline
- Reorder Jenkinsfile execution: build -> generate data -> run test
- Remove commented-out debug code from generate_test_dataset.sh
- Ensure all files end with proper newline character (POSIX compliance)
- Keep useful status messages while removing development debug prints
- Set MAX_ITERATIONS=0 for unlimited test generation in production
* Add configuration modes to reduce test execution time
- Add --mode option (half/full) to generate_model_configs.py
- half mode (default): ~278 configs (224 2D + 54 3D) -> ~1,058 total tests
- full mode: ~807 configs (672 2D + 135 3D) -> ~3,093 total tests
- Update generate_test_dataset.sh to use CONFIG_MODE environment variable
- Keeps all model types but reduces parameter combinations intelligently
- Fixes Jenkins timeout issue (was running 3,669 tests taking 17+ hours)
- Default half mode should complete in ~4-5 hours instead of 17+ hours
* Add small mode for quick testing of comprehensive dataset
* jenkins pipeline test done
* jenkins test done
* Trigger CI build
* remove test comment and update data generation option as half
---------
Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>
[ROCm/composable_kernel commit: 19d5327c45]
* Remove some unnecessary calls to create_args in basic and universal GEMM tests
* Remove unnecessary include statements in universal GEMM tests
* Improve compilation time of basic GEMM tests by only compiling the precision variants that we need
* Universal GEMM PrecType should be the same as CDataType
* Improve compilation time of universal GEMM tests by only compiling the precision variants that we need
* Revert to constexpr when defining some constants
[ROCm/composable_kernel commit: 5e85c38d7d]
Configure C++ standard with a CMake variable.
Defaults to C++20, but can be set to C++17 to test backwards compatibility.
* Add validation for allowed C++ standards.
* build CK in rehl8 docker with std=c++17
---------
Co-authored-by: illsilin_amdeng <Illia.Silin@amd.com>
[ROCm/composable_kernel commit: 99d27aca17]
* update the performance test logic
* fix unstash perf logs logic
* untangle unstashing fmha logs for different archs
* run process stage after running fmha tests
* fix the processing of perf logs
* fix arguments for run_performance scripts
[ROCm/composable_kernel commit: 6180685688]
Downstream libraries aren't migrated to c++20 yet, so replace a use of c++20 concept with equivalent SFINAE logic. The template checks for both the existence and the truthiness of the static member variable.
[ROCm/composable_kernel commit: c71d7ddd74]
* script for generating list of not referenced files in tests, list is in json format
* script comment added
* added empty line at the end of the script
* format changes
[ROCm/composable_kernel commit: 49c6b05c72]
* base working version for single groupped conv bwd data
* Fix 2d descriptor
* fix groups
* Add 3d support
* fixes
* fixes
* fixes
---------
Co-authored-by: Jakub Piasecki <jakpia21@gmail.com>
[ROCm/composable_kernel commit: 4212bbc170]
* CK Tile Stream K Device Ops
Implementation of CK Tile StreamKHostArgs and StreamKKernel classes. The
StreamKKernel class injects Universal Gemm and includes functions to
facilitate kernel preparation for the GPU.
* Stream K Device Ops Fixes
- Update GetWorkSpaceSize to call TilePartitioner's GetWorkSpaceSize to
ensure we get size needed for accumulation buffers and semaphores.
- Pass in num_sk_blocks into TilePartitioner constructor
- Update documentation
* Add WarpTile dimensions to GetName function in StreamKKernel class
* Fix typos in StreamKHostArgs class description.
Co-authored-by: Christopher Millette <63608002+cgmillette@users.noreply.github.com>
* Apply clang format on updated comment for StreamKHostArgs
* Explicitly specify type for StreamKReductionStrategy enum
* Remove unecessary scopes
* Unify the commenting style to inline comments
* Add explicit casts for occupancy and num_cu in MakeKernelArgs function
Both the static functions Occupancy and NumCU in the StreamKKernel class
use functions from the HIP API that result in the returned occupancy and
num_cu types being type int. The TilePartitioner interface for stream K will
have occupancy and num_cu being type ck_tile::index_t which is int32_t.
Thus, to be safe, this change ensures that both occupancy and num_cu are
cast to int32_t.
* Fix use of kentry due to interface update
PR #2594 updated the interface for the kentry function in
include/ck_tile/host/kernel_launch.hpp. As a result, the static function
Occupancy was updated to work correctly with the new interface.
PR #2594 also changed UniversalGemmKernel's KernelBlockSize static
variable to kBlockSize, so the StreamKKernel class was updated to
reflect this change.
* Switch type of num_sk_blocks from uint32_t to int32_t
This change switches the type of num_sk_blocks to type ck_tile::index_t
which is int32_t. This was done because parallel work for the CK Tile
StreamK TilePartitioner's constructor will have num_sk_blocks as
ck_tile::index_t. Thus, this change will help unify the interfaces to
avoid any type conversion errors.
---------
Co-authored-by: Christopher Millette <63608002+cgmillette@users.noreply.github.com>
[ROCm/composable_kernel commit: 071165919f]
* stash ckprofiler package built for all targets
* build the lib for all instances in newer docker
* make sure packages get posted
[ROCm/composable_kernel commit: 8b55afcd93]
* Start adding other layouts for gemm_ab_scale
* Add some instances
* Create tensor descriptors for A/B scales depending on A/B layout
* Fix formatting
* Revert some comments
* Revert commented instances in CMakeLists.txt
* Add some more instances for col-row gemm
* enable more row,row instances
* Use occupancy=1 for col,row layout to avoid spills
[ROCm/composable_kernel commit: 26d3300930]
* feat(gemm_wp): add two new configs for wp
* delete the unnecessary files
* fix the config error
* update the config
---------
Co-authored-by: AviralGoelAMD <aviral.goel@amd.com>
[ROCm/composable_kernel commit: 5ada85ec04]
* feat(gemm_wp): add two new configs for wp
* delete the unnecessary files
---------
Co-authored-by: ThomasNing <thomas.ning@amd.com>
[ROCm/composable_kernel commit: c06e8b4a66]
* Refactor CK tile permute ctests to gtests
* Refactor CK tile MOE smoothquant ctests to gtests
* fix typo in comment
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
* Update invalid case in else clause for get_precision_string
* Refactor permute gtests to use templated versions of matrix_core_swizzle and permute functions
---------
Co-authored-by: root <root@splinter-126-wr-c2.aus.dcgpu>
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
[ROCm/composable_kernel commit: 10395fc895]
* refactor moe_sorting ctests to use gtest framework
* Refactor ctests for smoothquant to gtests
* fix clang format to use version 18
* Print local_eid in MOE sorting gtests
* Remove extra space in smoothquant output
[ROCm/composable_kernel commit: 70dce4e0c6]
* Rename vector to ThreadTile
* more notes on tile encoding
* remove number<> from tuple of make_tile_window
* add script to stress test the copy example
[ROCm/composable_kernel commit: 8a698c7445]
* Remove some duplicate code in fmha_fwd_appendkv_kernel.hpp
* Simplify two templated operator calls by having the templated types deduced automatically
* Simplify two GemmPipeline calls
* Fix GemmPipelineAgBgCrCompV4::GetName
* Refactor use of ArgParser in CK tile GEMM examples
* Update args in README.md to match the implementation in create_args
* Remove some unnecessary include statements
* Rename two variables
* Factor out common code
* Factor out do_verify
* Add and use type aliases for memory operation integral constants
* In gemm_basic.cpp, use kPadM, kPadN, kPadK, and kBlockPerCu from GemmConfig
---------
Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
[ROCm/composable_kernel commit: 28a97865f5]
