### Build and Install with Ninja

Source: https://context7.com/rocm/aotriton/llms.txt

Builds and installs the project using Ninja. Use 'ninja install' to retain debug symbols if needed.

```bash
ninja install/strip
```

--------------------------------

### Install for Testing

Source: https://context7.com/rocm/aotriton/llms.txt

Installs the project using Ninja after configuration for testing.

```bash
ninja install
```

--------------------------------

### Known Issues Section Example

Source: https://github.com/rocm/aotriton/blob/main/PR.instructions.md

Example of documenting limitations, pending updates, and platform-specific issues.

```markdown
## Known Issues

* [db] The operator tuning database is not updated.
* [test] The number of functionals supported by FWD AITER ASM kernels are
  fairly limited.
* [aiter] The following AITER ASM functions need more investigation:
  - "Group mode" AITER ASM kernels to support varlen
  - MQA/GQA support
```

--------------------------------

### Install Configuration Header

Source: https://github.com/rocm/aotriton/blob/main/CMakeLists.txt

Configures and installs the AOTriton configuration header file. This ensures build-time configurations are available to the installed library.

```cmake
include(GNUInstallDirs)
configure_file(include/aotriton/config.h.in include/aotriton/config.h)
install(FILES ${CMAKE_CURRENT_BINARY_DIR}/include/aotriton/config.h
  DESTINATION ${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_INCLUDEDIR}/aotriton)
```

--------------------------------

### Build and Install AOTriton

Source: https://github.com/rocm/aotriton/blob/main/docs/How To Run Tests.md

Use CMake to configure and build AOTriton. Ensure AOTRITON_NAME_SUFFIX is set to prevent symbol conflicts. The Ninja generator is recommended for faster builds. Install the built artifacts to a specified directory.

```bash
mkdir -p build-test
cd build-test
# AOTRITON_NAME_SUFFIX is essential to avoid symbol conflicts with AOTriton bundled in PyTorch
cmake .. -DCMAKE_INSTALL_PREFIX=./install_dir -DCMAKE_BUILD_TYPE=Release -DAOTRITON_GPU_BUILD_TIMEOUT=0 -G Ninja -DAOTRITON_NAME_SUFFIX=123
# Optionally only build for one arch
# cmake .. -DCMAKE_INSTALL_PREFIX=./install_dir -DCMAKE_BUILD_TYPE=Release -DAOTRITON_GPU_BUILD_TIMEOUT=0 -G Ninja -DAOTRITON_NAME_SUFFIX=123 -DAOTRITON_TARGET_ARCH=gfx942
ninja install
```

--------------------------------

### Overview Section Examples

Source: https://github.com/rocm/aotriton/blob/main/PR.instructions.md

Examples of high-level PR summaries detailing motivation and performance impact.

```markdown
# Overview

This PR enables kernel pipelining and XCD (compute die) remapping
optimizations for gfx950 GPUs (MI350 series), significantly improving
performance of Flash Attention kernels.

This update delivers 904 TFLOPS performance on MI355X with mainline Triton JIT,
improved from 753 TFLOPS with the same Triton JIT process without pipelining.
```

```markdown
# Overview

This PR changes the API for varlen inputs and expects compact logsumexp
(LSE) tensor, following Tri's FlashAttention API.

Previously for varlen inputs with `B` sequences, `H` heads, AOTriton
expects a regular-sized LSE Tensor: `(B*H, Max_sequence_length)`.
However this approach requires padding at last dimension and wasting
memory when the sequence lengths change drastically within a batch.
```

--------------------------------

### Install Python Dependencies

Source: https://context7.com/rocm/aotriton/llms.txt

Installs necessary Python packages from the requirements.txt file. Ensure you are in the project's root directory.

```bash
pip install -r requirements.txt
```

--------------------------------

### Notes Section Example

Source: https://github.com/rocm/aotriton/blob/main/PR.instructions.md

Example of an optional notes section for technical clarifications.

```markdown
## Notes

* The code uses `hdim_vo` instead of `hdim_v` for better alignment with
  `hdim_qk`, and emphasizes the output tensor (`o`) should follow Tensor `V`'s
  head dimension.
```

--------------------------------

### Build and Profile AOTriton

Source: https://github.com/rocm/aotriton/blob/main/docs/How To Generate Tuning Database.md

Commands to configure, build, and run profiling for AOTriton tuning. Ensure the install prefix is set to avoid conflicts.

```bash
mkdir cpptune_build
cd cpptune_build
# -DCMAKE_INSTALL_PREFIX is mandatory to avoid conflicts with aotriton bundled by AOTriton
cmake .. -DCMAKE_INSTALL_PREFIX=./install_dir -DCMAKE_BUILD_TYPE=Release -DAOTRITON_BUILD_FOR_TUNING=ON -DAOTRITON_NAME_SUFFIX=123 -G Ninja
# Optionally only build for one arch
# cmake .. -DCMAKE_INSTALL_PREFIX=./install_dir -DCMAKE_BUILD_TYPE=Release -DAOTRITON_BUILD_FOR_TUNING=ON -DAOTRITON_OVERRIDE_TARGET_GPUS=gfx942_mod0 -DAOTRITON_NAME_SUFFIX=123 -G Ninja
ninja install
cd ..
# Run profiling on Target GPU
# We do not recommend updating the tuning_database.sqlite3 directly
PYTHONPATH=cpptune_build/bindings/ python test/tune_flash.py --json_file ~/navi32-aotriton_0.8.json --use_multigpu -1
# Update tuning database from experiment data stored in JSON file
python v2python/table_tool.py --action rawjson -k FLASH -f v2python/rules/tuning_database.sqlite3 -i ~/navi32-aotriton_0.8.json
```

--------------------------------

### Configure Python Virtual Environment and Dependencies

Source: https://github.com/rocm/aotriton/blob/main/CMakeLists.txt

Initializes the virtual environment path and installs required Python packages from requirements.txt.

```cmake
execute_process(COMMAND ${CMAKE_COMMAND} -E env VIRTUAL_ENV=${VENV_DIR} "${VENV_BIN_PYTHON}" -c "import site; print(site.getsitepackages()[0], end='')" OUTPUT_VARIABLE VENV_SITE)
message("VENV_SITE ${VENV_SITE}")

execute_process(COMMAND ${CMAKE_COMMAND} -E env VIRTUAL_ENV=${VENV_DIR} "${VENV_BIN_PYTHON}" -m pip install -r "${CMAKE_CURRENT_LIST_DIR}/requirements.txt")
```

--------------------------------

### TemplateParameter Initialization Example

Source: https://github.com/rocm/aotriton/blob/main/v3python/base/README.md

Demonstrates the initialization process for TemplateParameter objects, including setting choices for tensor types, conditional values, and integer parameters. Shows how TypedChoice and TypedChoice.constexpr are used.

```Python
literal in TYPE_CHOICES
-> (parse_choices)
-> Stored in TemplateParameter.choices.
   Let tp1 = TypeParameter(('Q', 'K'), ['*fp16', '*bf16'])
       tp2 = TypeParameter(('V',), [CDETensor('USE_V', False, 0, 'Q')])
       tp3 = TypeParameter(('HDIM',), [16, 32, 48, 64])
       TC  = typed_choice
       TCC = typed_choice.constexpr
   tp1.choices = [ TC.tensor(elem_ty='fp16', rank=any), TC.tensor(elem_ty='bf16', rank=any) ]
   tp2.choices = [ TC.CDETensor('USE_V', False, 0, 'Q') ] // CDETensor unchanged
   tp3.choices = TypeParameter(('HDIM',), [TCC.int16(16), TCC.int16(32), TCC.int16(48), TCC.int16(64)])
   // Note tp3 (HDIM) is fully resolved (called "settled") now
-> (late_init() -> link_deferral_target())
-> tp2  : [ CDETensor(fp_x, False, 0, tp1) ] // Link CDETensor to tp1 and fp_x (USE_V)
-> (late_init() -> resolve_rank())
## Design I
-> tp1.type_matrix = { // argname -> list of TC
     'Q' : [ TC.tensor(elem_ty='fp16', rank=2), TC.tensor(elem_ty='fp16', rank=2) ]
     'K' : [ TC.tensor(elem_ty='fp16', rank=3), TC.tensor(elem_ty='fp16', rank=3) ] // Note rank=3
   }
   // tp1 is settled now
## Design II
-> tp1.choices = [
        TC.tensor(elem_ty='fp16', rank=any).specialized = {
            'Q' : TC.tensor(elem_ty='fp16', rank=2),
            'K' : TC.tensor(elem_ty='fp16', rank=3),
        },
        TC.tensor(elem_ty='bf16', rank=any).specialized = {
            'Q' : TC.tensor(elem_ty='bf16', rank=2),
            'K' : TC.tensor(elem_ty='bf16', rank=3),
        },
   ]

```

--------------------------------

### Configure AOTriton Build and Installation

Source: https://github.com/rocm/aotriton/blob/main/v3src/CMakeLists.txt

This CMake script manages the compilation of the AOTriton library, handles kernel tuning regeneration, sets target properties, and defines installation paths for headers and configuration files.

```cmake
if(AOTRITON_BUILD_FOR_TUNING)
  # Regenerate with --ignore_missing_kernels
  add_custom_target(aotriton_v2_regen_shim
    COMMAND ${CMAKE_COMMAND}
    -E env VIRTUAL_ENV=${VENV_DIR}
    AOTRITON_ENABLE_FP32=${AOTRITON_ENABLE_FP32}
    "${VENV_BIN_PYTHON}" -m v3python.generate
    --target_gpus ${EFFECTIVE_TARGET_GPUS}
    --build_dir "${AOTRITON_V2_BUILD_DIR}"
    ${GENERATE_OPTION}
    --build_for_tuning_second_pass
    WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_PARENT_DIR}")
  if(NOT AOTRITON_NOIMAGE_MODE)
    add_dependencies(aotriton_v2_regen_shim aotriton_v2_compile)
  endif(NOT AOTRITON_NOIMAGE_MODE)
endif(AOTRITON_BUILD_FOR_TUNING)

set(AOTRITON_SHIM_FLAGS "")
if(AOTRITON_NAME_SUFFIX)
  list(APPEND AOTRITON_SHIM_FLAGS "--library_suffix" "${AOTRITON_NAME_SUFFIX}")
endif()
message(STATUS "AOTRITON_SHIM_FLAGS ${AOTRITON_SHIM_FLAGS}")

file(STRINGS "${AOTRITON_V2_BUILD_DIR}/Bare.shim" SHIM_CC_FILES ENCODING UTF-8)

aux_source_directory(. CC_FILES)
aux_source_directory(flash/ FLASH_CC_FILES)
aux_source_directory(flash/aiter FLASH_AITER_CC_FILES)
add_library(aotriton_v2 SHARED ${SHIM_CC_FILES} ${CC_FILES} ${FLASH_CC_FILES} ${FLASH_AITER_CC_FILES})
set_target_properties(aotriton_v2 PROPERTIES LINKER_LANGUAGE CXX)
target_link_libraries(aotriton_v2 PRIVATE ${CMAKE_DL_LIBS})
if(AOTRITON_ENABLE_SUFFIX)
  set_target_properties(aotriton_v2 PROPERTIES OUTPUT_NAME "aotriton${AOTRITON_NAME_SUFFIX}_v2")
endif()

target_include_directories(aotriton_v2 PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_LIST_DIR}/../include>)
# for <build>/include/aotriton/config.h. Code should use <aotriton/config.h>
target_include_directories(aotriton_v2 PUBLIC $<BUILD_INTERFACE:${CMAKE_BINARY_DIR}/include>)
target_include_directories(aotriton_v2 INTERFACE $<INSTALL_INTERFACE:include>)
target_include_directories(aotriton_v2 PRIVATE ${CMAKE_CURRENT_BINARY_DIR})
set_target_properties(aotriton_v2 PROPERTIES
  POSITION_INDEPENDENT_CODE ON
)
if(AOTRITON_BUILD_FOR_TUNING)
  target_compile_definitions(aotriton_v2 PRIVATE -DAOTRITON_BUILD_FOR_TUNING=1)
  add_dependencies(aotriton_v2 aotriton_v2_regen_shim)
else(AOTRITON_BUILD_FOR_TUNING)
  target_compile_definitions(aotriton_v2 PRIVATE -DAOTRITON_BUILD_FOR_TUNING=0)
endif(AOTRITON_BUILD_FOR_TUNING)
target_link_libraries(aotriton_v2 PRIVATE lzma_interface)
target_link_libraries(aotriton_v2 PUBLIC hip::host hip::amdhip64)
set_target_properties(aotriton_v2 PROPERTIES VERSION
  ${AOTRITON_VERSION_MAJOR_INT}.${AOTRITON_VERSION_MINOR_INT}.${AOTRITON_VERSION_PATCH_INT})

if(NOT WIN32)
  execute_process(
    COMMAND ${CMAKE_COMMAND}
    -E env VIRTUAL_ENV=${VENV_DIR}
    "${VENV_BIN_PYTHON}" -m v3python.comment_only_asm
    -o "${AOTRITON_V2_BUILD_DIR}/set_aotriton_version.s"
    ${AOTRITON_VERSION_MAJOR_INT}
    ${AOTRITON_VERSION_MINOR_INT}
    ${AOTRITON_VERSION_PATCH_INT}
    WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_PARENT_DIR}"
    COMMAND_ERROR_IS_FATAL ANY
  )
  target_sources(aotriton_v2 PRIVATE "${AOTRITON_V2_BUILD_DIR}/set_aotriton_version.s")
else()
  # FIXME: Add version string to Windows binaries.
endif()
# Otherwise the binary size blows up
# FIXME: Properly export symbols
set_target_properties(aotriton_v2 PROPERTIES CXX_VISIBILITY_PRESET hidden)

include(GNUInstallDirs)
message("CMAKE_INSTALL_INCLUDEDIR ${CMAKE_INSTALL_INCLUDEDIR}")
install(DIRECTORY "${CMAKE_CURRENT_SOURCE_PARENT_DIR}/include/aotriton"
        DESTINATION ${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_INCLUDEDIR}
        FILES_MATCHING PATTERN "*.h")
install(TARGETS aotriton_v2
        EXPORT aotriton-targets
        DESTINATION ${CMAKE_INSTALL_PREFIX}/lib)
install(EXPORT aotriton-targets
  DESTINATION ${CMAKE_INSTALL_PREFIX}/lib/cmake/aotriton)

# Packages
include(CMakePackageConfigHelpers)
configure_package_config_file(${CMAKE_CURRENT_SOURCE_DIR}/config.cmake.in
  "${CMAKE_CURRENT_BINARY_DIR}/aotriton-config.cmake"
  INSTALL_DESTINATION "lib/cmake/aotriton"
  NO_SET_AND_CHECK_MACRO
  NO_CHECK_REQUIRED_COMPONENTS_MACRO)
write_basic_package_version_file(
  "${CMAKE_CURRENT_BINARY_DIR}/aotriton-config-version.cmake"
  VERSION "${AOTRITON_VERSION_MAJOR_INT}.${AOTRITON_VERSION_MINOR_INT}.${AOTRITON_VERSION_PATCH_INT}"
  COMPATIBILITY AnyNewerVersion)
install(FILES
  ${CMAKE_CURRENT_BINARY_DIR}/aotriton-config.cmake
  ${CMAKE_CURRENT_BINARY_DIR}/aotriton-config-version.cmake
  DESTINATION lib/cmake/aotriton)

if(NOT AOTRITON_NOIMAGE_MODE)
  install(DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/aotriton.images"
          DESTINATION ${CMAKE_INSTALL_PREFIX}/lib)
endif(NOT AOTRITON_NOIMAGE_MODE)

# Python binding only available for AOTriton V2 API
add_library(aotriton INTERFACE)
target_link_libraries(aotriton INTERFACE aotriton_v2)
```

--------------------------------

### Minor Changes Section Example

Source: https://github.com/rocm/aotriton/blob/main/PR.instructions.md

Example of a minor changes list for refactoring, documentation, and small fixes.

```markdown
## Minor Changes

* [test] Adjust unit tests for varlen compact LSE tensor
* [docs] Update comments about input dimensions in V2 header file
* [build] Replace `git log -1 --format=%H` with `git rev-parse HEAD`
```

--------------------------------

### Build aotriton

Source: https://github.com/rocm/aotriton/blob/main/README.md

Use these commands to build the aotriton library. Ensure all prerequisites are met before starting. The `export PKG_CONFIG_PATH` line is only needed when building with conda.

```bash
pip install -r requirements.txt
mkdir build
cd build
export PKG_CONFIG_PATH="${PKG_CONFIG_PATH}:${CONDA_PREFIX}/lib/pkgconfig"
cmake .. -DCMAKE_INSTALL_PREFIX=./install_dir -DCMAKE_BUILD_TYPE=Release -DAOTRITON_GPU_BUILD_TIMEOUT=0 -G Ninja
# Use ccmake to tweak options
ninja install/strip  # Use `ninja install` to keep symbols
```

--------------------------------

### Process Alternative Triton Wheel Configurations

Source: https://github.com/rocm/aotriton/blob/main/CMakeLists.txt

Parses a YAML configuration file to set up additional virtual environments and install specific Triton wheels.

```cmake
if(NOT AOTRITON_NOIMAGE_MODE AND AOTRITON_ALT_TRITON_WHEEL_CONFIG_FILE)
  execute_process(
    COMMAND ${CMAKE_COMMAND} -E env VIRTUAL_ENV=${VENV_DIR} "${VENV_BIN_PYTHON}" -c
    "import yaml,sys; f = open(sys.argv[1]); d = yaml.load(f, Loader=yaml.Loader); print(';'.join([f'{k};{v}' for k, v in d['venvs'].items()]))"
    "${AOTRITON_ALT_TRITON_WHEEL_CONFIG_FILE}"
    OUTPUT_VARIABLE ALT_VENVS
    OUTPUT_STRIP_TRAILING_WHITESPACE
  )
  while(ALT_VENVS)
    list(POP_FRONT ALT_VENVS ALT_VENV_NAME)
    list(POP_FRONT ALT_VENVS WHEEL)
    set(ALT_VENV_DIR "altvenvs/${ALT_VENV_NAME}")
    message(STATUS "ALT_VENV_DIR ${ALT_VENV_DIR}")
    message(STATUS "WHEEL ${WHEEL}")
    execute_process(COMMAND "${Python3_EXECUTABLE}" -m venv "${ALT_VENV_DIR}")
    set(ALT_VENV_BIN_PYTHON "${ALT_VENV_DIR}/${VENV_REL_PYTHON}")
    execute_process(COMMAND "${ALT_VENV_BIN_PYTHON}" -m pip install -r "${CMAKE_CURRENT_LIST_DIR}/requirements.txt")
    set(AOTRITON_TRITON_SO "${ALT_VENV_DIR}/${VENV_REL_TRITON}")
    add_custom_command(OUTPUT "${AOTRITON_TRITON_SO}"
      COMMAND "${ALT_VENV_BIN_PYTHON}" -m pip install ${WHEEL}
    )
    LIST(APPEND ALL_TRITON_SOS "${AOTRITON_TRITON_SO}")
  endwhile(ALT_VENVS)
endif(NOT AOTRITON_NOIMAGE_MODE AND AOTRITON_ALT_TRITON_WHEEL_CONFIG_FILE)
```

--------------------------------

### Major Changes Section Example

Source: https://github.com/rocm/aotriton/blob/main/PR.instructions.md

Example of a major changes list using categorized bullet points and sub-bullets.

```markdown
## Major Changes

* [api] **BREAKING** Use compact LSE for varlen inputs
* [kernel] Use compact LSE and Delta Tensor
* [kernel] `head_dim` argument in all SDPA Triton kernels is replaced by
  `hdim_qk` and `hdim_vo` arguments to support this feature.
* [shim] The dispatcher will use the last dimensions of `Q` and `V` tensor as
  `hdim_qk` and `hdim_vo`
* [test] Add test cases to test `hdim_qk != hdim_vo`. New cases are added to
  + `test_fast` case (`FOR_RELEASE=0`)
  + `test_hdim_qk_ne_vo` (`FOR_RELEASE=3`)
```

--------------------------------

### Run Tests

Source: https://context7.com/rocm/aotriton/llms.txt

Executes tests using pytest. The PYTHONPATH is set to include the installation directory, and tests are run verbosely.

```bash
PYTHONPATH=install_dir/lib/ pytest ../test/test_forward.py -v
```

--------------------------------

### Configure CMake Build

Source: https://context7.com/rocm/aotriton/llms.txt

Configures the project using CMake with specific build options. This command sets the installation prefix, build type, GPU build timeout, target architectures, and uses the Ninja generator.

```bash
mkdir build && cd build
cmake .. \
    -DCMAKE_INSTALL_PREFIX=./install_dir \
    -DCMAKE_BUILD_TYPE=Release \
    -DAOTRITON_GPU_BUILD_TIMEOUT=0 \
    -DAOTRITON_TARGET_ARCH="gfx90a;gfx942;gfx1100" \
    -G Ninja
```

--------------------------------

### Verify or Build Triton Dependency

Source: https://github.com/rocm/aotriton/blob/main/CMakeLists.txt

Checks for existing Triton installation or triggers a build process based on project configuration flags.

```cmake
if(AOTRITON_INHERIT_SYSTEM_SITE_TRITON)
  execute_process(
    COMMAND ${CMAKE_COMMAND} -E env VIRTUAL_ENV=${VENV_DIR} "${VENV_BIN_PYTHON}" -c "import triton; import triton.language; print('Triton found: ', triton.__version__)"
    RESULT_VARIABLE TRITON_CHECK_RESULT
    OUTPUT_VARIABLE TRITON_OUTPUT
    ERROR_VARIABLE TRITON_OUTPUT
    OUTPUT_STRIP_TRAILING_WHITESPACE
  )

  if(TRITON_CHECK_RESULT EQUAL 0)
    message(STATUS "[AOTriton] ${TRITON_OUTPUT}")
    add_custom_target(aotriton_venv_triton ALL)
  else()
    message(FATAL_ERROR "[AOTriton] AOTRITON_INHERIT_SYSTEM_SITE_TRITON is enabled but triton is not available in the virtual environment. Please install triton or disable this option. ${TRITON_OUTPUT}")
  endif()
else()
  set(TRITON_BUILD_DIR "${CMAKE_CURRENT_BINARY_DIR}/triton_build")
  execute_process(COMMAND ${CMAKE_COMMAND} -E make_directory "${TRITON_BUILD_DIR}")

  set(AOTRITON_TRITON_SO "${VENV_SITE}/${VENV_REL_TRITON}")

  if(AOTRITON_USE_LOCAL_TRITON_WHEEL)
    add_custom_command(OUTPUT "${AOTRITON_TRITON_SO}"
      COMMAND ${CMAKE_COMMAND} -E env VIRTUAL_ENV=${VENV_DIR} "${VENV_BIN_PYTHON}" -m pip install ${AOTRITON_USE_LOCAL_TRITON_WHEEL}
    )
  else()
    execute_process(
      COMMAND "${VENV_BIN_PYTHON}" -m pip --version
      RESULT_VARIABLE _PIP_RESULT
      OUTPUT_QUIET ERROR_QUIET
    )
    if(_PIP_RESULT EQUAL 0)
      add_custom_command(OUTPUT "${AOTRITON_TRITON_SO}"
        COMMAND ${CMAKE_COMMAND} -E env TRITON_BUILD_PROTON=OFF VIRTUAL_ENV=${VENV_DIR} TRITON_BUILD_DIR=${TRITON_BUILD_DIR} "${VENV_BIN_PYTHON}" -m pip install .
        WORKING_DIRECTORY "${CMAKE_CURRENT_LIST_DIR}/third_party/triton/"
      )
    else()
      message(STATUS "[AOTriton] pip not available, falling back to setup.py develop for triton")
      add_custom_command(OUTPUT "${AOTRITON_TRITON_SO}"
        COMMAND ${CMAKE_COMMAND} -E env TRITON_BUILD_PROTON=OFF VIRTUAL_ENV=${VENV_DIR} TRITON_BUILD_DIR=${TRITON_BUILD_DIR} "${VENV_BIN_PYTHON}" setup.py develop
        WORKING_DIRECTORY "${CMAKE_CURRENT_LIST_DIR}/third_party/triton/"
      )
    endif()
  endif()
  LIST(APPEND ALL_TRITON_SOS "${AOTRITON_TRITON_SO}")
endif()
```

--------------------------------

### C++ Tensor Type Example

Source: https://github.com/rocm/aotriton/blob/main/v3python/base/README.md

Illustrates how C++ templates can define functions accepting tensors of varying ranks but the same elemental type.

```C++
typename<typename T>
void func(Eigen::Tensor<T, 2>& Q, Eigen::Tensor<T, 3>& K);
```

--------------------------------

### Run AOTriton Tests

Source: https://github.com/rocm/aotriton/blob/main/docs/How To Run Tests.md

Execute AOTriton tests using pytest. Set the PYTHONPATH to include the installation directory to ensure the correct AOTriton library is used. The FOR_RELEASE=1 environment variable may be used for release builds.

```bash
FOR_RELEASE=1 PYTHONPATH=install_dir/lib/ pytest ../test/test_backward.py -v
```

--------------------------------

### AOTriton Flash Attention Forward Pass with PyTorch

Source: https://context7.com/rocm/aotriton/llms.txt

This snippet shows how to use the `attn_fwd` function from `pyaotriton.v3.flash` to perform a forward pass of flash attention. It includes tensor creation, parameter setup, and execution, demonstrating the integration with PyTorch tensors on CUDA.

```python
import torch
from pyaotriton import T0, T1, T2, T4, DType, Stream, hipError_t
from pyaotriton.v3.flash import attn_fwd, attn_fwd_params, attn_options

def cast_dtype(dtype):
    """Convert PyTorch dtype to AOTriton DType."""
    bits = dtype.itemsize * 8
    if dtype.is_floating_point:
        maintype = 'Float' if 'bfloat' not in str(dtype) else 'BFloat'
    else:
        maintype = 'Int' if 'uint' not in str(dtype) else 'UInt'
    return getattr(DType, f'k{maintype}{bits}')

def mk_aotensor(tensor):
    """Create AOTriton TensorView from PyTorch tensor."""
    if tensor is None:
        return None
    rank = len(tensor.shape)
    dtype = cast_dtype(tensor.dtype)
    if rank == 4:
        return T4(tensor.data_ptr(), tuple(tensor.size()), tensor.stride(), dtype)
    elif rank == 2:
        return T2(tensor.data_ptr(), tuple(tensor.size()), tensor.stride(), dtype)
    elif rank == 1:
        return T1(tensor.data_ptr(), tuple(tensor.size()), tensor.stride(), dtype)
    elif rank == 0 or tensor.numel() <= 1:
        return T0(tensor.data_ptr(), dtype)

# Create tensors on GPU
batch, heads, seq_len, head_dim = 2, 8, 512, 64
q = torch.randn(batch, heads, seq_len, head_dim, device='cuda', dtype=torch.float16)
k = torch.randn(batch, heads, seq_len, head_dim, device='cuda', dtype=torch.float16)
v = torch.randn(batch, heads, seq_len, head_dim, device='cuda', dtype=torch.float16)
out = torch.empty_like(q)
lse = torch.empty(batch * heads, seq_len, device='cuda', dtype=torch.float32)

# Set up parameters
params = attn_fwd_params()
params.Q = mk_aotensor(q)
params.K = mk_aotensor(k)
params.V = mk_aotensor(v)
params.B = T4(0, [0]*4, [0]*4, DType.kFloat16)  # Null tensor
params.Sm_scale = 1.0 / (head_dim ** 0.5)
params.L = mk_aotensor(lse)
params.Out = mk_aotensor(out)
params.dropout_p = 0.0
params.philox_seed_ptr = T0(0, DType.kUInt64)
params.philox_offset1 = T0(0, DType.kUInt64)
params.philox_offset2 = 0
params.philox_seed_output = T0(0, DType.kUInt64)
params.philox_offset_output = T0(0, DType.kUInt64)
params.encoded_softmax = T4(0, [0]*4, [0]*4, DType.kFloat16)
params.persistent_atomic_counter = T0(0, DType.kInt32)
params.causal_type = 0  # Non-causal
params.varlen_type = 0
params.window_left = 0
params.window_right = 0

# Execute
err = attn_fwd(params, attn_fwd_params.kVersion, Stream(), None)
assert err == hipError_t.hipSuccess
torch.cuda.synchronize()
```

--------------------------------

### Initialize Build Directories

Source: https://github.com/rocm/aotriton/blob/main/v3src/CMakeLists.txt

Sets up the primary build and kernel storage directories using CMake commands.

```cmake
message("CMAKE_SOURCE_DIR ${CMAKE_SOURCE_DIR}")
message("CMAKE_CURRENT_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}")
cmake_path(GET CMAKE_CURRENT_SOURCE_DIR PARENT_PATH CMAKE_CURRENT_SOURCE_PARENT_DIR)
message("CMAKE_CURRENT_SOURCE_PARENT_DIR ${CMAKE_CURRENT_SOURCE_PARENT_DIR}")
message("CMAKE_CURRENT_LIST_DIR ${CMAKE_CURRENT_LIST_DIR}")
message("CMAKE_CURRENT_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}")
set(AOTRITON_V2_BUILD_DIR "${CMAKE_CURRENT_BINARY_DIR}")
set(AOTRITON_KERNEL_STORAGE_V2_DIR "${AOTRITON_V2_BUILD_DIR}/aotriton.images")
execute_process(COMMAND ${CMAKE_COMMAND} -E make_directory "${AOTRITON_V2_BUILD_DIR}")
execute_process(COMMAND ${CMAKE_COMMAND} -E make_directory "${AOTRITON_KERNEL_STORAGE_V2_DIR}")
```

--------------------------------

### Create New Database File

Source: https://github.com/rocm/aotriton/blob/main/docs/How To Update Constraints of Tuning Database.md

Initializes a new tuning database file with the updated schema.

```bash
python tritonsrc/tune_flash.py \
		--db_file v2python/rules/new_tuning_database.sqlite3 \
		--stop_at 0 --create_table_only
```

--------------------------------

### Build AOTriton 0.10b

Source: https://github.com/rocm/aotriton/blob/main/dockerfile/README.md

Execute the build script for version 0.10b with specified GPU architectures.

```bash
cd aotriton/dockerfile
bash build.sh input tmpfs output 0.10b "gfx90a;gfx942;gfx950;gfx1201"
```

--------------------------------

### Build AOTriton 0.7.1b

Source: https://github.com/rocm/aotriton/blob/main/dockerfile/README.md

Execute the build script for versions prior to 0.10b using specific hardware targets.

```bash
cd aotriton/dockerfile
bash build.sh input tmpfs output 0.7.1b "MI300X;MI200;Navi31"
```

--------------------------------

### Download Cached Triton LLVM Package

Source: https://github.com/rocm/aotriton/blob/main/dockerfile/README.md

URL for the required LLVM/MLIR package to speed up the build process.

```text
https://oaitriton.blob.core.windows.net/public/llvm-builds/llvm-657ec732-almalinux-x64.tar.gz
```

--------------------------------

### Python Tensor Rank Definitions

Source: https://github.com/rocm/aotriton/blob/main/v3python/base/README.md

Defines tensor names and their corresponding ranks using Python dictionaries. This setup is used to manage tensor dimensions within the parameter system.

```Python
TYPE_CHOICES = {
    ('Q', 'K') : ['*fp16', '*bf16'],
}
TENSOR_RANKS = {
    'Q' : 2,
    'K' : 3,
}
```

--------------------------------

### Generate Bare and shim files using Python

Source: https://github.com/rocm/aotriton/blob/main/v3src/CMakeLists.txt

Invokes a Python module to generate build files with specific environment variables and target GPU configurations. Requires the virtual environment path and build directory to be defined.

```cmake
execute_process(
COMMAND ${CMAKE_COMMAND}
-E env VIRTUAL_ENV=${VENV_DIR}
AOTRITON_ENABLE_FP32=${AOTRITON_ENABLE_FP32}
"${VENV_BIN_PYTHON}" -X utf8 -m v3python.generate
--target_gpus ${EFFECTIVE_TARGET_GPUS}
--build_dir "${AOTRITON_V2_BUILD_DIR}"
${GENERATE_OPTION}
COMMAND_ECHO STDOUT
WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_PARENT_DIR}"
COMMAND_ERROR_IS_FATAL ANY
)
```

--------------------------------

### Initialize TensorView

Source: https://context7.com/rocm/aotriton/llms.txt

Creates non-owning views into GPU memory. The last stride must always be 1 for contiguous memory access.

```cpp
#include <aotriton/util.h>

// Create a 4D tensor view for attention tensors
// Shape: [batch, heads, seq_len, head_dim]
std::array<uint64_t, 4> sizes = {2, 8, 512, 64};
std::array<uint64_t, 4> strides = {8*512*64, 512*64, 64, 1};
intptr_t gpu_ptr = reinterpret_cast<intptr_t>(device_memory);

aotriton::TensorView<4> tensor(gpu_ptr, sizes, strides, aotriton::kFloat16);

// Access tensor properties
uint64_t batch_size = tensor.size(0);      // 2
uint64_t head_dim = tensor.size(3);        // 64
uint64_t seq_stride = tensor.stride(2);    // 64
void* data = tensor.data_ptr();
aotriton::DType dtype = tensor.dtype();    // kFloat16

// Create null tensor for optional parameters
auto null_tensor = aotriton::TensorView<4>::get_null_tensor(aotriton::kFloat16);

// Scalar tensor (rank 0) for philox seed/offset
aotriton::TensorView<0> scalar(seed_ptr, aotriton::kUInt64);
```

--------------------------------

### Enable Kernel Compression for AOTriton < 0.8

Source: https://github.com/rocm/aotriton/wiki/Common-Build-Problems

For AOTriton versions 0.5 through 0.7, set -DAOTRITON_COMPRESS_KERNEL=ON to prevent linker errors caused by GPU kernels being too large to embed.

```bash
-DAOTRITON_COMPRESS_KERNEL=ON
```

--------------------------------

### Configure Build Options

Source: https://github.com/rocm/aotriton/blob/main/v3src/CMakeLists.txt

Sets various compiler and generation flags based on project configuration variables.

```cmake
get_filename_component(AOTRITON_COMPILER "${CMAKE_CURRENT_LIST_DIR}/../v3python/compile.py" ABSOLUTE)
message("AOTRITON_COMPILER ${AOTRITON_COMPILER}")

set(HSACO_COMPILE_OPTION "")
if(AOTRITON_TERMINATE_WHEN_GPU_BUILD_TIMEOUT)
  list(APPEND HSACO_COMPILE_OPTION "--error_when_timeout")
endif()

if(AOTRITON_BUILD_FOR_TUNING)
  set(GENERATE_OPTION "--build_for_tuning")
else(AOTRITON_BUILD_FOR_TUNING)
  set(GENERATE_OPTION "")
endif()

if(AOTRITON_BUILD_FOR_TUNING_BUT_SKIP_KERNEL)
  if (NOT AOTRITON_BUILD_FOR_TUNING)
    message(FATAL_ERROR "AOTRITON_BUILD_FOR_TUNING_BUT_SKIP_KERNEL is only relevant with AOTRITON_BUILD_FOR_TUNING=ON")
  endif()
  list(APPEND GENERATE_OPTION "--build_for_tuning_but_skip_kernel")
  list(APPEND GENERATE_OPTION ${AOTRITON_BUILD_FOR_TUNING_BUT_SKIP_KERNEL})
endif()

if(AOTRITON_NOIMAGE_MODE)
  list(APPEND GENERATE_OPTION "--noimage_mode")
endif()

if(AOTRITON_ALT_TRITON_WHEEL_CONFIG_FILE)
  list(APPEND GENERATE_OPTION "--alt_venv_config")
  list(APPEND GENERATE_OPTION "${AOTRITON_ALT_TRITON_WHEEL_CONFIG_FILE}")
endif()

if(WIN32)
  find_package(dlfcn-win32 REQUIRED)
  set(CMAKE_DL_LIBS dlfcn-win32::dl)
endif()

if(AOTRITON_ENABLE_FP32_INPUTS)
  set(AOTRITON_ENABLE_FP32 1)
else()
  set(AOTRITON_ENABLE_FP32 0)
endif()
```

--------------------------------

### Clone AOTriton Repository Recursively

Source: https://github.com/rocm/aotriton/wiki/Common-Build-Problems

Use this command to ensure all submodules, including pybind11, are downloaded when cloning the AOTriton repository.

```bash
git clone --recursive
```

--------------------------------

### Replace Old Database

Source: https://github.com/rocm/aotriton/blob/main/docs/How To Update Constraints of Tuning Database.md

Overwrites the existing database file with the updated version.

```bash
cp v2python/rules/new_tuning_database.sqlite3 v2python/rules/tuning_database.sqlite3
```

--------------------------------

### Dump Database Table to CSV

Source: https://github.com/rocm/aotriton/blob/main/docs/How To Update Constraints of Tuning Database.md

Exports a specific table from the tuning database to a CSV file for migration.

```bash
python -m v2python.table_tool --action dumpcsv \
		-f v2python/rules/tuning_database.sqlite3 \
		--table_name 'FLASH$attn_fwd' --table_file attn_fwd.csv
```

--------------------------------

### GPU Detection Utilities

Source: https://context7.com/rocm/aotriton/llms.txt

Provides functions for querying GPU architecture and capabilities at runtime, enabling dispatch to architecture-specific kernels.

```APIDOC
## GPU Detection Utilities

### Description
Provides functions for querying GPU architecture and capabilities at runtime, enabling dispatch to architecture-specific kernels. These utilities detect the current GPU from a HIP stream.

### Functions
- `aotriton::getGpuFromStream(hipStream_t stream)`: Detects the GPU architecture from a HIP stream.
- `aotriton::isArchExperimentallySupported(hipStream_t stream)`: Checks if the current GPU architecture is experimentally supported.
- `aotriton::getMultiProcessorCount(hipStream_t stream)`: Returns the number of Streaming Multiprocessors (SMs) for occupancy calculations.
- `aotriton::Gpu2VendorArch(aotriton::Gpu gpu)`: Extracts the vendor and architecture from a `Gpu` enum value.

### Enums
- `aotriton::Gpu`: Represents GPU architectures (e.g., `GPU_AMD_ARCH_GFX942_MOD0` for MI300X).

### Example Usage
```cpp
#include <aotriton/util.h>

// ...

hipStream_t stream;
hipStreamCreate(&stream);

// Detect GPU architecture
aotriton::Gpu gpu = aotriton::getGpuFromStream(stream);

// Check for specific architectures
if (gpu == aotriton::GPU_AMD_ARCH_GFX942_MOD0) {
    // MI300X detected
}

// Check experimental support
bool experimental = aotriton::isArchExperimentallySupported(stream);

// Get SM count
int sm_count = aotriton::getMultiProcessorCount(stream);

// Get vendor architecture
uint32_t vendor_arch = aotriton::Gpu2VendorArch(gpu);
```
```

--------------------------------

### Create Virtual Environment

Source: https://github.com/rocm/aotriton/blob/main/CMakeLists.txt

Creates a Python virtual environment using the specified Python executable. It can optionally inherit system site packages.

```cmake
set(Python_ARTIFACTS_INTERACTIVE TRUE)

# Not a target, we need to override Python3_EXECUTABLE later
if(AOTRITON_INHERIT_SYSTEM_SITE_TRITON)
  execute_process(COMMAND "${Python3_EXECUTABLE}" -m venv --system-site-packages "${VENV_DIR}")
else()
  execute_process(COMMAND "${Python3_EXECUTable}" -m venv "${VENV_DIR}")
endif()

set(ENV{VIRTUAL_ENV} "${VENV_DIR}")
message("VENV_DIR ${VENV_DIR}")
if(WIN32)
  set(VENV_REL_PYTHON "Scripts/python")
  set(VENV_REL_TRITON "triton/_C/libtriton.pyd")
else()
  set(VENV_REL_PYTHON "bin/python")
  set(VENV_REL_TRITON "triton/_C/libtriton.so")
endif()
set(VENV_BIN_PYTHON "${VENV_DIR}/${VENV_REL_PYTHON}")
```

--------------------------------

### Manage GPU Streams

Source: https://context7.com/rocm/aotriton/llms.txt

Wraps HIP stream handles for asynchronous execution. The default constructor initializes a null stream.

```cpp
#include <aotriton/runtime.h>

// Use default stream
aotriton::Stream default_stream;

// Wrap existing HIP stream
hipStream_t hip_stream;
hipStreamCreate(&hip_stream);
aotriton::Stream stream(hip_stream);

// Get native HIP stream for direct HIP API calls
hipStream_t native = stream.native();
```

--------------------------------

### Checkout Git tag using execute_process

Source: https://github.com/rocm/aotriton/blob/main/v3src/CMakeLists.txt

Executes a Git checkout command within a specified working directory. Fails the build if the command returns an error.

```cmake
execute_process(
        COMMAND ${GIT_EXECUTABLE} checkout ${AITER_GIT_TAG}
        WORKING_DIRECTORY "${AITER_DIR}"
        COMMAND_ERROR_IS_FATAL ANY
      )
```

--------------------------------

### Link liblzma for Kernel Storage V2

Source: https://github.com/rocm/aotriton/blob/main/CMakeLists.txt

Configures the build to use liblzma for compression in Kernel Storage V2. It handles finding and linking the library differently for Unix-like systems and others.

```cmake
if(UNIX)
  include(FindPkgConfig)
  pkg_search_module(LZMA REQUIRED liblzma)
  add_library(lzma_interface INTERFACE)
  target_link_libraries(lzma_interface INTERFACE ${LZMA_LIBRARIES})
  target_link_directories(lzma_interface INTERFACE ${LZMA_LIBRARY_DIRS})
  target_include_directories(lzma_interface INTERFACE ${LZMA_INCLUDE_DIRS})
else()
  find_package(liblzma REQUIRED CONFIG)
  if(NOT liblzma_FOUND)
    message(FATAL_ERROR "liblzma package not found")
  else()
    message(STATUS "liblzma found: ${liblzma_VERSION}")
  endif()
  add_library(lzma_interface INTERFACE)
  target_link_libraries(lzma_interface INTERFACE liblzma::liblzma)
endif()
```

--------------------------------

### Run Parallel AOTriton Tests on Multi-GPUs

Source: https://github.com/rocm/aotriton/blob/main/docs/How To Run Tests.md

For AOTriton versions 0.10.0 and later, tests can be run in parallel across multiple GPUs. This requires the 'jq' utility and the 'amd-smi' Python CLI. The number of GPUs is determined dynamically.

```bash
# Optionally starting from AOTriton >= 0.10.0
# it is possible to run tests in parallel on multi-GPUs
# NGPUS=$(amd-smi list --json|jq length)
# FOR_RELEASE=1 PYTHONPATH=install_dir/lib/ pytest -n $NGPUS ../test/test_backward.py -v
```

--------------------------------

### Update Python Version on RHEL7

Source: https://github.com/rocm/aotriton/wiki/Common-Build-Problems

If AOTriton requires a newer Python version than available on RHEL7, use 'scl enable' to switch to a supported version like Python 3.8.

```bash
scl enable rh-python38
```

--------------------------------

### Set Runtime Attention Options in C++

Source: https://context7.com/rocm/aotriton/llms.txt

Configures kernel execution behavior such as backend selection and determinism using the attn_options struct.

```cpp
#include <aotriton/flash.h>

aotriton::v3::flash::attn_options options;

// Force specific backend implementation (useful for debugging/benchmarking)
options.force_backend_index = -1;  // -1 = auto-select best kernel
options.force_backend_index = 0;   // Force backend 0

// Enable deterministic mode (may be slower but reproducible)
options.deterministic = true;

// Pass options to attention kernel
hipError_t err = aotriton::v3::flash::attn_fwd(
    params,
    aotriton::v3::flash::attn_fwd_params::kVersion,
    stream,
    &options
);
```

--------------------------------

### Load CSV Data into New Database

Source: https://github.com/rocm/aotriton/blob/main/docs/How To Update Constraints of Tuning Database.md

Imports the previously dumped CSV data into the newly created database table.

```bash
python -m v2python.table_tool --action loadcsv \
	-f v2python/rules/new_tuning_database.sqlite3 \
	--table_name 'FLASH$attn_fwd' --table_file attn_fwd.csv -k ''
```

--------------------------------

### Query GPU Capabilities

Source: https://context7.com/rocm/aotriton/llms.txt

Detects GPU architecture and capabilities at runtime from a HIP stream.

```cpp
#include <aotriton/util.h>

hipStream_t stream;
hipStreamCreate(&stream);

// Detect GPU architecture
aotriton::Gpu gpu = aotriton::getGpuFromStream(stream);

// Check for specific architectures
if (gpu == aotriton::GPU_AMD_ARCH_GFX942_MOD0) {
    // MI300X detected
}

// Check if architecture is experimentally supported
bool experimental = aotriton::isArchExperimentallySupported(stream);

// Get SM count for occupancy calculations
int sm_count = aotriton::getMultiProcessorCount(stream);

// Extract vendor and architecture from Gpu enum
uint32_t vendor_arch = aotriton::Gpu2VendorArch(gpu);
```

--------------------------------

### Configure CMake for Testing

Source: https://context7.com/rocm/aotriton/llms.txt

Configures the project for testing with CMake, setting a specific name suffix to avoid conflicts with PyTorch. Uses the Ninja generator.

```bash
cmake .. \
    -DCMAKE_INSTALL_PREFIX=./install_dir \
    -DCMAKE_BUILD_TYPE=Release \
    -DAOTRITON_NAME_SUFFIX=test123 \
    -G Ninja
```

--------------------------------

### Populate Default Values

Source: https://github.com/rocm/aotriton/blob/main/docs/How To Update Constraints of Tuning Database.md

Updates the new table columns with default values using standard sqlite3 commands.

```bash
sqlite3 v2python/rules/new_tuning_database.sqlite3 'update FLASH$attn_fwd set inputs$BIAS_TYPE = 0;'
```