### Install aiter Source: https://github.com/rocm/aiter/blob/main/csrc/ck_batched_gemm_a8w8/README.md Installs the 'aiter' library by navigating to its directory and running the setup script. ```bash cd $aiter_path python3 setup.py develop ``` -------------------------------- ### Install aiter Source: https://github.com/rocm/aiter/blob/main/csrc/ck_batched_gemm_bf16/README.md Installs the aiter library by running the setup script in develop mode. ```python cd $aiter_path python3 setup.py develop ``` -------------------------------- ### Install Development Tools and Pre-commit Hook Source: https://github.com/rocm/aiter/blob/main/CONTRIBUTE.md Installs Python linters/formatters (Black, Ruff) and system dependencies (clang-format), then installs the project's pre-commit hook. ```python pip install black==25.1.0 ruff==0.11.11 ``` ```bash apt install clang-format ``` ```bash bash ./.githooks/install ``` -------------------------------- ### Install aiter Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a8w8_blockscale/README.md Installs the aiter library by navigating to the aiter path and running the setup script. ```python cd $aiter_path python3 setup.py develop ``` -------------------------------- ### Install aiter Source: https://github.com/rocm/aiter/blob/main/aiter/configs/model_configs/README.md Installs the aiter library by navigating to the aiter directory and running the setup script. ```shell cd $aiter_path python3 setup.py develop ``` -------------------------------- ### Install aiter Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a4w4_blockscale/README.md Installs the aiter package by running the setup script in develop mode. ```bash cd $aiter_path python3 setup.py develop ``` -------------------------------- ### Install aiter Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a8w8/README.md Installs the aiter library by running the setup script in development mode. ```bash cd $aiter_path python3 setup.py develop ``` -------------------------------- ### Install System Dependencies Source: https://github.com/rocm/aiter/blob/main/csrc/cpp_itfs/README.MD Installs essential system libraries required for ROCm AI tools, including SSL development files, fmt library development files, and Google Test development files. ```bash apt install libssl-dev libfmt-dev libgtest-dev ``` -------------------------------- ### Install aiter Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a8w8_bpreshuffle/README.md Installs the aiter library by running the setup script in develop mode. ```bash cd $aiter_path python3 setup.py develop ``` -------------------------------- ### Install ROCm AITER Source: https://github.com/rocm/aiter/blob/main/README.md Instructions for cloning the ROCm AITER repository and installing it using the setup script. Includes steps for initializing submodules if `--recursive` was omitted during cloning. ```bash git clone --recursive https://github.com/ROCm/aiter.git cd aiter python3 setup.py develop ``` ```bash git submodule sync && git submodule update --init --recursive ``` -------------------------------- ### Build Aiter Dockerfile Source: https://github.com/rocm/aiter/blob/main/docs/aiter_container_nonroot_setup.md This Dockerfile installs Aiter and its dependencies (pandas, zmq, einops, numpy) and sets up a non-root user 'newuser' to run the container. It clones the Aiter repository and installs it in develop mode. ```dockerfile ARG BASE_DOCKER="rocm/pytorch:latest" FROM $BASE_DOCKER RUN pip install pandas zmq einops && \ pip install numpy==1.26.2 # Create a new user RUN useradd -m newuser # Switch to the new user USER newuser # Set the working directory WORKDIR /home/newuser RUN rm -rf aiter && \ git clone --recursive https://github.com/ROCm/aiter.git && \ cd aiter && \ python3 setup.py develop ``` -------------------------------- ### Start GEMM tuning Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a8w8_blockscale/README.md Initiates the tuning process for GEMM shapes by running a Python script with input and output configuration files. This process may take a few minutes as it builds gemm_a8w8_blockscale_tune via JIT. ```python python3 csrc/ck_gemm_a8w8_blockscale/gemm_a8w8_blockscale_tune.py -i aiter/configs/a8w8_blockscale_untuned_gemm.csv -o aiter/configs/a8w8_blockscale_tuned_gemm.csv ``` -------------------------------- ### Start GEMM tuning Source: https://github.com/rocm/aiter/blob/main/csrc/ck_batched_gemm_a8w8/README.md Initiates the tuning process for batched GEMM a8w8 operations by specifying input and output CSV files for shapes. ```bash python3 csrc/ck_batched_gemm_a8w8/batched_gemm_a8w8_tune.py -i aiter/configs/a8w8_untuned_batched_gemm.csv -o aiter/configs/a8w8_tuned_batched_gemm.csv ``` -------------------------------- ### Start GEMM a8w8 tuning Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a8w8_bpreshuffle/README.md Initiates the tuning process for GEMM a8w8 operations. This command builds the gemm_a8w8_bpreshuffle_tune kernel via JIT and outputs the tuned parameters to a CSV file. ```bash python3 csrc/ck_gemm_a8w8_bpreshuffle/gemm_a8w8_bpreshuffle_tune.py -i aiter/configs/a8w8_bpreshuffle_untuned_gemm.csv -o aiter/configs/a8w8_bpreshuffle_tuned_gemm.csv ``` -------------------------------- ### Start GEMM tuning Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a4w4_blockscale/README.md Initiates the tuning process for GEMM shapes by running a Python script. It requires specifying an input CSV file with GEMM shapes and an output CSV file for the tuning results. The process involves JIT compilation of the tuning kernels. ```bash GEMM_A4W4_BLOCKWISE_HIP_CLANG_PATH=/work/llvm-project/build/bin/ python3 csrc/ck_gemm_a4w4_blockscale/gemm_a4w4_blockscale_tune.py -i aiter/configs/a4w4_blockscale_untuned_gemm.csv -o aiter/configs/a4w4_blockscale_tuned_gemm.csv ``` -------------------------------- ### Start GEMM tuning Source: https://github.com/rocm/aiter/blob/main/csrc/ck_batched_gemm_bf16/README.md Initiates the tuning process for batched GEMM bf16 operations. It takes an input CSV file with GEMM shapes and outputs a CSV file with tuned shapes. This process involves JIT compilation. ```python python3 csrc/ck_batched_gemm_bf16/batched_gemm_bf16_tune.py -i aiter/configs/bf16_untuned_batched_gemm.csv -o aiter/configs/bf16_tuned_batched_gemm.csv ``` -------------------------------- ### Install Python Dependencies Source: https://github.com/rocm/aiter/blob/main/csrc/cpp_itfs/README.MD Installs the Jinja2 Python package, which is a dependency for ROCm AI tools. ```python pip install jinja2 ``` -------------------------------- ### Start GEMM a8w8 Tuning Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a8w8/README.md Initiates the tuning process for GEMM a8w8 operations. It takes an input CSV file with GEMM shapes and outputs a CSV file with tuning results. The `-k` parameter can be used for split K kernels. ```bash python3 csrc/ck_gemm_a8w8/gemm_a8w8_tune.py -i aiter/configs/a8w8_untuned_gemm.csv -o aiter/configs/a8w8_tuned_gemm.csv # For split K kernels: python3 csrc/ck_gemm_a8w8/gemm_a8w8_tune.py -i aiter/configs/a8w8_untuned_gemm.csv -o aiter/configs/a8w8_tuned_gemm.csv -k ``` -------------------------------- ### Prebuild Kernels and Update Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a8w8/README.md Information on how pre-building kernels during installation affects the process. If new tuning results are available, it suggests removing existing build artifacts and reinstalling, though this is time-consuming. ```bash # During installation: PREBUILD_KERNELS=1 # To use new tuning results after installation: rm -rf aiter/jit/*.so python3 setup.py develop ``` -------------------------------- ### Run AITER Operator Tests Source: https://github.com/rocm/aiter/blob/main/README.md Demonstrates how to run operator tests provided by AITER, using `test_layernorm2d.py` as an example. This allows users to verify the functionality of individual operators. ```python python3 op_tests/test_layernorm2d.py ``` -------------------------------- ### Run Aiter Gemm Test Source: https://github.com/rocm/aiter/blob/main/docs/aiter_container_nonroot_setup.md Executes a specific Aiter test case for gemm (General Matrix Multiply) with A8W8 quantization. This command verifies the functionality of the Aiter library, particularly its performance with mixed-precision matrix multiplication. ```python newuser@0d2817135822:~/aiter$ python3 op_tests/test_gemm_a8w8_blockscale.py ``` -------------------------------- ### Build and Run mha Kernel Benchmarks Source: https://github.com/rocm/aiter/blob/main/op_tests/cpp/mha/README.md Scripts to build the device library and executable for benchmarking the aiter mha fwd/bwd kernels. Supports building forward, backward, or both. ```Shell bash build_mha.sh bash build_mha.sh fwd bash build_mha.sh bwd ``` -------------------------------- ### Tune GEMMs using gradlib Source: https://github.com/rocm/aiter/blob/main/gradlib/README.md This snippet shows how to tune GEMMs using the gradlib script. It takes an input file containing untuned GEMM shapes and outputs a file with tuned configurations. ```python python3 gradlib/gradlib/gemm_tuner.py --tuned_file aiter/configs/tuned_gemm.csv --input_file aiter/configs/untuned_gemm.csv ``` -------------------------------- ### Build Docker Image Source: https://github.com/rocm/aiter/blob/main/docs/aiter_container_nonroot_setup.md Command to build the Docker image for Aiter using the specified Dockerfile and tag it as 'rocm-aiter:test'. The --no-cache flag ensures that the build uses the latest layers. ```shell docker build --no-cache -t rocm-aiter:test -f Dockerfile.aiter . ``` -------------------------------- ### Build and test tuned kernels Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a8w8_blockscale/README.md Tests the performance of the tuned GEMM kernels. This involves modifying a test instance and running a Python script, which will build the tuned kernels via JIT. If kernels were previously built, prepend `AITER_REBUILD=1` to the command. ```python python3 op_tests/test_gemm_a8w8_blockscale.py ``` ```python AITER_REBUILD=1 python3 op_tests/test_gemm_a8w8_blockscale.py ``` -------------------------------- ### Check GPU Permissions with rocminfo Source: https://github.com/rocm/aiter/blob/main/docs/aiter_container_nonroot_setup.md This command, executed inside the running container, displays information about the ROCm system and detected GPU agents. It's used to verify that the container has access to the GPU hardware. ```shell newuser@0d2817135822:~$ rocminfo ``` -------------------------------- ### FP8 GEMM Kernel Performance Tuning (Per Block Quantization) Source: https://github.com/rocm/aiter/blob/main/aiter/configs/model_configs/README.md Tunes FP8 GEMM kernels for Llama-70B using per-block quantization. This requires copying a configuration file and following the specified tuning steps. ```shell cp aiter/configs/model_configs/llama70B_untuned_gemm.csv aiter/configs/a8w8_blockscale_untuned_gemm.csv # Refer to CK gemm a8w8 blockscale tune steps in https://github.com/ROCm/aiter/tree/main/csrc/ck_gemm_a8w8_blockscale ``` -------------------------------- ### Run Aiter Container with GPU Access Source: https://github.com/rocm/aiter/blob/main/docs/aiter_container_nonroot_setup.md Command to run the Aiter Docker container interactively. It maps GPU devices (/dev/dri, /dev/kfd) and adds the current user's render and video group IDs to ensure proper permissions for GPU access within the container. ```shell docker run -it --device=/dev/dri --device=/dev/kfd --group-add $(getent group render | cut -d: -f3) --group-add $(getent group video | cut -d: -f3) rocm-aiter:test /bin/bash ``` -------------------------------- ### Build and test tuned kernels Source: https://github.com/rocm/aiter/blob/main/csrc/ck_batched_gemm_a8w8/README.md Tests the performance of the tuned batched GEMM a8w8 kernels. Includes an option to force a rebuild if kernels have been previously built. ```bash python3 op_tests/test_batched_gemm_a8w8.py ``` ```bash AITER_REBUILD=1 python3 op_tests/test_batched_gemm_a8w8.py ``` -------------------------------- ### Build and test tuned kernels Source: https://github.com/rocm/aiter/blob/main/csrc/ck_batched_gemm_bf16/README.md Builds the tuned batched_gemm_bf16 kernels and tests their performance. It may require rebuilding kernels if they were previously built with different configurations. The `AITER_REBUILD=1` flag forces a rebuild. ```python python3 op_tests/test_batched_gemm_bf16.py ``` ```python AITER_REBUILD=1 python3 op_tests/test_batched_gemm_bf16.py ``` -------------------------------- ### FP8 GEMM Kernel Performance Tuning (Plain Weight) Source: https://github.com/rocm/aiter/blob/main/aiter/configs/model_configs/README.md Tunes FP8 GEMM kernels for Llama-70B with per-token or per-tensor quantization and plain weight format. Requires copying a configuration file and following specific tuning steps. ```shell cp aiter/configs/model_configs/llama70B_untuned_gemm.csv aiter/configs/a8w8_untuned_gemm.csv # Refer to CK gemm a8w8 tune steps in https://github.com/ROCm/aiter/tree/main/csrc/ck_gemm_a8w8 ``` -------------------------------- ### Compile and Link aiter mha in C++ Project Source: https://github.com/rocm/aiter/blob/main/op_tests/cpp/mha/README.md Steps to compile the device kernel into a shared object (.so) using a Python script and then link it into a C++ executable. Requires specifying include paths for aiter and composable_kernel headers. ```Python python3 compile.py ``` ```Shell # Example linking command: g++ your_project.cpp -o your_executable -L. -l:libmha_fwd.so -l:libmha_bwd.so -I$TOP_DIR/csrc/include -I$TOP_DIR/3rdparty/composable_kernel/include -I$TOP_DIR/3rdparty/composable_kernel/example/ck_tile/01_fmha/ ``` -------------------------------- ### MXFP4 GEMM Kernel Performance Tuning (Per Block Quantization) Source: https://github.com/rocm/aiter/blob/main/aiter/configs/model_configs/README.md Tunes MXFP4 GEMM kernels for Llama-70B with per-block quantization and weight preshuffling. This involves copying a configuration file and following the relevant tuning steps. ```shell cp aiter/configs/model_configs/llama70B_untuned_gemm.csv aiter/configs/a4w4_blockscale_untuned_gemm.csv # Refer to CK gemm a4w4 blockscale tune steps in https://github.com/ROCm/aiter/tree/main/csrc/ck_gemm_a4w4_blockscale ``` -------------------------------- ### FP8 GEMM Kernel Performance Tuning (Preshuffled Weight) Source: https://github.com/rocm/aiter/blob/main/aiter/configs/model_configs/README.md Tunes FP8 GEMM kernels for Llama-70B with per-token or per-tensor quantization and weight in preshuffling format. Involves copying a configuration file and following tuning instructions. ```shell cp aiter/configs/model_configs/llama70B_untuned_gemm.csv aiter/configs/a8w8_bpreshuffle_untuned_gemm.csv # Refer to CK gemm a8w8 tune preshuffle steps in https://github.com/ROCm/aiter/tree/main/csrc/ck_gemm_a8w8_bpreshuffle ``` -------------------------------- ### General GEMM Kernel Performance Tuning (BF16) Source: https://github.com/rocm/aiter/blob/main/aiter/configs/model_configs/README.md Tunes GEMM kernels for Llama-70B using BF16 data type. This involves copying a configuration file and following tuning steps. ```shell cp aiter/configs/model_configs/llama70B_untuned_gemm_bf16.csv aiter/configs/untuned_gemm.csv # Refer to GEMM performance tune steps in https://github.com/ROCm/aiter/tree/main/gradlib ``` -------------------------------- ### Test tuned GEMM kernels Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a4w4_blockscale/README.md Tests the performance of the tuned GEMM kernels. This command also involves JIT compilation of the tuned kernels. An option to force a rebuild of the kernels is provided if previous builds exist. ```bash GEMM_A4W4_BLOCKWISE_HIP_CLANG_PATH=/work/llvm-project/build/bin/ python3 op_tests/test_gemm_a4w4_blockscale.py ``` ```bash AITER_REBUILD=1 python3 op_tests/test_gemm_a4w4_blockscale.py ``` -------------------------------- ### Test tuned GEMM a8w8 kernels Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a8w8_bpreshuffle/README.md Tests the performance of the tuned GEMM a8w8 kernels. This involves building the kernels using JIT based on the tuned parameters and running performance tests. If kernels were previously built, set AITER_REBUILD=1 to force a rebuild. ```bash python3 op_tests/test_gemm_a8w8.py ``` ```bash AITER_REBUILD=1 python3 op_tests/test_gemm_a8w8.py ``` -------------------------------- ### GPU Input Generation for Triton Tests Source: https://github.com/rocm/aiter/blob/main/op_tests/triton_tests/README.md Generate test inputs directly on the GPU when possible for Triton tests. This approach is approximately two orders of magnitude faster for large test cases compared to generating on the CPU and then moving to the GPU. ```Python # Preferred method for generating large test inputs on GPU torch.randn((M, K), device="cuda") # Less performant method: generate on CPU then move to GPU # torch.randn((M, K)).cuda() ``` -------------------------------- ### Capture GEMM Shapes with gradlib Source: https://github.com/rocm/aiter/blob/main/gradlib/README.md This snippet demonstrates how to capture GEMM shapes for tuning by replacing F.linear with tgemm.mm in the specified Python file. The captured shapes are saved to a CSV file. ```bash AITER_TUNE_GEMM=1 python {workload_tests} ``` -------------------------------- ### Build and Test Tuned Kernels Source: https://github.com/rocm/aiter/blob/main/csrc/ck_gemm_a8w8/README.md Builds the tuned GEMM a8w8 kernels and tests their performance. The `AITER_REBUILD=1` flag can be used to force a rebuild if kernels have been previously built. ```bash python3 op_tests/test_gemm_a8w8.py # To force rebuild: AITER_REBUILD=1 python3 op_tests/test_gemm_a8w8.py ``` -------------------------------- ### Run ROCm Aiter Log Analysis Source: https://github.com/rocm/aiter/blob/main/aiter_logs/readme.md Executes the log analysis script for the ROCm Aiter project. It takes the log directory and analysis type as arguments. ```bash python3 ./aiter_logs/run.py ./aiter_logs/ ar_asm all ``` -------------------------------- ### Tensor Comparison in Triton Tests Source: https://github.com/rocm/aiter/blob/main/op_tests/triton_tests/README.md Use torch.testing.assert_close for tensor comparisons in Triton tests as it is significantly faster than triton.testing.assert_close, which uses numpy internally. Switch to Triton's implementation if facing Out-of-Memory (OOM) issues. ```Python import torch # Recommended for performance torch.testing.assert_close(tensor1, tensor2) # Alternative, potentially slower or memory-intensive # import triton.testing # triton.testing.assert_close(tensor1, tensor2) ``` -------------------------------- ### aiter::mha_bwd Configuration Table Source: https://github.com/rocm/aiter/blob/main/op_tests/cpp/mha/README.md A comprehensive table detailing the supported arguments for the aiter::mha_bwd function. It covers various configurations including data types (fp16, bf16), head dimensions (hdim_q, hdim_v), modes (batch, group), mask types, accumulation modes, and constraints for different hardware platforms and kernel types. ```APIDOC aiter::mha_bwd Configuration: | data_type | hdim_q | hdim_v | mode | mask_type | dq_accumulation | general constraints | shape&stride constraints | kernel type(asm/ck) | mi308 | mi300/325 | mi350/355 | |--------------|--------------|-----------------|----------------|--------------------------|--------------------------|---------------------------------------------------------|----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------|---------------------|-------|-----------|----------------------------------| | fp16 or bf16 | (128,192]/x8 | equal to hdim_q | batch or group | no_mask or causal | atomic_f32 | bias, dbisa, dropout and deterministic is not supported | dq_acc only support BHSD | asm | y | y | n | | fp16 or bf16 | (64,128]/x8 | equal to hdim_q | batch | no_mask or causal | atomic_f32 or atomic_f16 | bias, dbisa, dropout and deterministic is not supported | dq_acc only support BHSD when dq_accumulation is atomic_f32. The shape&stride of q and do must be the same and the shape&stride of k and v must be the same and seqlen_q must be equal to seqlen_k when dq_accumulation is atomic_f16. | asm | y | y | bf16;hd128;sq == sk;sq % 256==0 | | fp16 or bf16 | (64,128]/x8 | equal to hdim_q | group | no_mask or causal | atomic_f32 | bias, dbisa, dropout and deterministic is not supported | dq_acc only support BHSD | asm | y | y | bf16;hd128;sq == sk;sq % 256==0 | | fp16 or bf16 | 64 | equal to hdim_q | batch | no_mask or causal | atomic_f32 or atomic_f16 | bias, dbisa, dropout and deterministic is not supported | dq_acc only support BHSD when dq_accumulation is atomic_f32. The shape&stride of q and do must be the same and the shape&stride of k and v must be the same and seqlen_q must be equal to seqlen_k when dq_accumulation is atomic_f16. | asm | y | y | n | | fp16 or bf16 | 64 | equal to hdim_q | group | no_mask or causal | atomic_f32 | bias, dbisa, dropout and deterministic is not supported | dq_acc only support BHSD | asm | y | y | n | | fp16 or bf16 | [0,32] | [0,32] | batch or group | no_mask or causal or swa | atomic_f32 or atomic_f16 | unconstrained | unconstrained | ck | y | y | y | | fp16 or bf16 | (0,64] | (0,64] | batch or group | no_mask or causal or swa | atomic_f32 or atomic_f16 | unconstrained | unconstrained | ck | y | y | y | ``` -------------------------------- ### aiter::mha_fwd Kernel Configuration Source: https://github.com/rocm/aiter/blob/main/op_tests/cpp/mha/README.md Configuration table detailing supported data types, dimensions, sequence lengths, modes, mask types, and constraints for the `aiter::mha_fwd` kernel across different hardware (mi308, mi300/325, mi350/355). Specifies kernel types (asm, ck) and their respective requirements. ```APIDOC aiter::mha_fwd supported arguments configuration: | data_type | hdim_q | hdim_v | seqlen_q | seqlen_k | mode | mask_type | general constraints | shape&stride constraints | kernel type | mi308 | mi300/325 | mi350/355 | |--------------|---------|---------|---------------|-------------------|----------------|--------------------------|--------------------------------|------------------------------------------------------------------------------------------------|-------------|-------|-----------|-------------------| | bf16 | 128 | 128 | [384,) | equal to seqlen_q | batch or group | no_mask or causal | bias, dropout is not supported | the shape&stride of q, k and v must be the same, the layout of q, k, v, o must be bshd or bhsd | asm | y | y | lse must be true | | fp16 or bf16 | [0,32] | [0,32] | unconstrained | unconstrained | batch or group | no_mask or causal or swa | unconstrained | unconstrained | ck | y | y | y | | fp16 or bf16 | (0,64] | (0,64] | unconstrained | unconstrained | batch or group | no_mask or causal or swa | unconstrained | unconstrained | ck | y | y | y | | fp16 or bf16 | (0,128] | (0,128] | unconstrained | unconstrained | batch or group | no_mask or causal or swa | unconstrained | unconstrained | ck | y | y | y | | fp16 or bf16 | (0,192] | (0,128] | unconstrained | unconstrained | batch or group | no_mask or causal or swa | unconstrained | unconstrained | ck | y | y | y | | fp16 or bf16 | (0,256] | (0,256] | unconstrained | unconstrained | batch or group | no_mask or causal or swa | unconstrained | unconstrained | ck | y | y | y | ``` === COMPLETE CONTENT === This response contains all available snippets from this library. No additional content exists. Do not make further requests.