ROCm · Micky774 · Apr 28, 2026 · Apr 28, 2026 · Apr 28, 2026 · Apr 28, 2026
@@ -8,48 +8,58 @@ project(ck_fused_attn LANGUAGES HIP CXX)
 
 set(AITER_MHA_INSTALL_PREFIX "transformer_engine" CACHE STRING "aiter mha shared lib install prefix in TE")
 
-#Corresponding runtime check is in nvte_get_fused_attn_backend()
-list(FIND CMAKE_HIP_ARCHITECTURES "gfx1250" _gfx1250_idx)
-if(NOT _gfx1250_idx EQUAL -1)
-  message(WARNING
-    "Removing unsupported gfx1250 from CMAKE_HIP_ARCHITECTURES for ck_fused_attn build.")
-  list(REMOVE_ITEM CMAKE_HIP_ARCHITECTURES "gfx1250")
-  list(LENGTH CMAKE_HIP_ARCHITECTURES _hip_arch_count)
-  if(_hip_arch_count EQUAL 0)
-    message(FATAL_ERROR
-            "No supported architectures remain for the ck_fused_attn build. "
-            "Re-run the build with FUSED_ATTN_CK backend disabled.")
-  endif()
-  set(GPU_TARGETS ${CMAKE_HIP_ARCHITECTURES})
-endif()
-set(__AITER_SOURCE_DIR "${CMAKE_CURRENT_LIST_DIR}/../../../3rdparty/QoLA/3rdparty/aiter")
+# gfx1250 carries AITER V3 bwd kernels only (hd128, bf16, batch mode). The
+# runtime envelope is enforced in nvte_get_fused_attn_backend().
+set(__QOLA_DIR "${CMAKE_CURRENT_LIST_DIR}/../../../3rdparty/QoLA")
+set(__AITER_SOURCE_DIR "${__QOLA_DIR}/build/third_party/aiter")
 set(__CK_SOURCE_DIR "${__AITER_SOURCE_DIR}/3rdparty/composable_kernel")
-
 set(CK_INCLUDE_DIR "${__CK_SOURCE_DIR}/include")
-message(STATUS "ck_include_dir: ${CK_INCLUDE_DIR}")
-if(NOT EXISTS "${CK_INCLUDE_DIR}")
-  message(FATAL_ERROR
-          "Could not find CK API. "
-          "Try running 'git submodule update --init --recursive' "
-          "within the Transformer Engine source.")
-endif()
-
 set(AITER_INCLUDE_DIR "${__AITER_SOURCE_DIR}/csrc/include")
-message(STATUS "aiter_include_dir: ${AITER_INCLUDE_DIR}")
-if(NOT EXISTS "${AITER_INCLUDE_DIR}")
-  message(FATAL_ERROR
-          "Could not find AITER API. "
-          "Try running 'git submodule update --init --recursive' "
-          "within the Transformer Engine source.")
-endif()
 
 if(NOT Python_EXECUTABLE)
   find_package(Python COMPONENTS Interpreter QUIET)
 endif()
 
+# Resolve the manifest-pinned AITER commit (defines AITER_SHA) and bring the
+# QoLA-managed AITER source tree to that commit before any consumer reads it
+# (header validation below, header includes for the .cpp build later, and
+# QoLA's own kernel build if the prebuilt cache misses).
+include("${CMAKE_CURRENT_LIST_DIR}/aiter_prebuilt.cmake")
+
 if(Python_EXECUTABLE)
+  set(__QOLA_MANIFEST "${CMAKE_CURRENT_LIST_DIR}/qola_manifest.toml")
+  # Redirect GIT_CONFIG_GLOBAL to a tempfile carrying `safe.directory = *` so
+  # git operations inside the QoLA-managed AITER tree (and its recursive
+  # submodules) work in containerized builds where the bind-mounted .git is
+  # owned by a different UID than the build process. Mirrors the pattern in
+  # transformer_engine/common/CMakeLists.txt:get_git_commit().
   execute_process(
-    COMMAND ${Python_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/check_aiter_mha_args.py --mode both --te-dir "${CMAKE_CURRENT_LIST_DIR}/../../.."
+    COMMAND sh -c
+      "tmp=$(mktemp /tmp/gitconfig.XXXXXX) || exit 1; \
+GIT_CONFIG_GLOBAL=$tmp git config --global --add safe.directory '*' >/dev/null 2>&1; \
+GIT_CONFIG_GLOBAL=$tmp PYTHONPATH=\"${__QOLA_DIR}:$PYTHONPATH\" '${Python_EXECUTABLE}' -m qola.cli checkout \
+--manifest '${__QOLA_MANIFEST}' \
+--aiter-root '${__AITER_SOURCE_DIR}'; \
+rc=$?; rm -f \"$tmp\"; exit $rc"
+    RESULT_VARIABLE AITER_CHECKOUT_RESULT
+    OUTPUT_VARIABLE AITER_CHECKOUT_OUTPUT
+    ERROR_VARIABLE AITER_CHECKOUT_ERROR
+    OUTPUT_STRIP_TRAILING_WHITESPACE
+    ERROR_STRIP_TRAILING_WHITESPACE
+  )
+  if(NOT AITER_CHECKOUT_RESULT EQUAL 0)
+    message(FATAL_ERROR
+            "Failed to sync AITER source tree at ${__AITER_SOURCE_DIR} to "
+            "manifest-pinned commit ${AITER_SHA}.\n"
+            "${AITER_CHECKOUT_OUTPUT}\n${AITER_CHECKOUT_ERROR}")
+  endif()
+  message(STATUS "[AITER] Synced ${__AITER_SOURCE_DIR} to ${AITER_SHA}")
+
+  execute_process(
+    COMMAND ${Python_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/check_aiter_mha_args.py
+            --mode both
+            --te-dir "${CMAKE_CURRENT_LIST_DIR}/../../.."
+            --aiter-root "${__AITER_SOURCE_DIR}"
     RESULT_VARIABLE AITER_ARG_CHECK_RESULT
     OUTPUT_VARIABLE AITER_ARG_CHECK_OUTPUT
     ERROR_VARIABLE AITER_ARG_CHECK_ERROR
@@ -64,50 +74,125 @@ if(Python_EXECUTABLE)
   endif()
   message(STATUS "AITER API validation passed via check_aiter_mha_args.py")
 else()
-  message(WARNING "Python interpreter not found; skipping AITER API validation.")
+  message(WARNING "Python interpreter not found; skipping AITER source-tree sync and API validation.")
 endif()
 
-if(DEFINED AITER_MHA_PATH)
-  message(STATUS "[AITER-BUILD] Using AITER_MHA_PATH=${AITER_MHA_PATH}")
-  # use pre-built te_libmha_fwd.so te_libmha_bwd.so
-  set(__AITER_MHA_PATH ${AITER_MHA_PATH})
-else()
-  set(__AITER_MHA_PATH "")
-  include("${CMAKE_CURRENT_LIST_DIR}/aiter_prebuilt.cmake")
-  get_prebuilt_aiter(__AITER_MHA_PATH)
+# Sanity-check the resolved include directories now that `qola checkout` has
+# materialized the AITER tree.
+message(STATUS "ck_include_dir: ${CK_INCLUDE_DIR}")
+if(NOT EXISTS "${CK_INCLUDE_DIR}")
+  message(FATAL_ERROR
+          "Could not find CK API at ${CK_INCLUDE_DIR}. "
+          "Re-run the build to let `qola checkout` clone AITER and its "
+          "composable_kernel submodule.")
+endif()
 
-  if(__AITER_MHA_PATH STREQUAL "")
-    # If not available, fallback: Build from source via QoLA
-    list(JOIN CMAKE_HIP_ARCHITECTURES ";" GPU_ARCHS_STR)
-    message(STATUS "[AITER-BUILD] Building AITER kernels for ${GPU_ARCHS_STR} via QoLA.")
-    set(__QOLA_DIR "${CMAKE_CURRENT_LIST_DIR}/../../../3rdparty/QoLA")
-    set(__QOLA_BUILD_DIR "${__QOLA_DIR}/build")
-    set(__QOLA_MANIFEST "${CMAKE_CURRENT_LIST_DIR}/qola_manifest.toml")
-    execute_process(
-      COMMAND ${CMAKE_COMMAND} -E env "PYTHONPATH=${__QOLA_DIR}:$ENV{PYTHONPATH}"
-              ${Python_EXECUTABLE} -m qola.cli build
-              --manifest ${__QOLA_MANIFEST}
-              --aiter-root ${__AITER_SOURCE_DIR}
-              --output-dir ${__QOLA_BUILD_DIR}
-              --arch "${GPU_ARCHS_STR}"
-      RESULT_VARIABLE QOLA_BUILD_RESULT
-    )
-    if(NOT QOLA_BUILD_RESULT EQUAL 0)
-      message(FATAL_ERROR "[AITER-BUILD] QoLA build failed.")
+message(STATUS "aiter_include_dir: ${AITER_INCLUDE_DIR}")
+if(NOT EXISTS "${AITER_INCLUDE_DIR}")
+  message(FATAL_ERROR
+          "Could not find AITER API at ${AITER_INCLUDE_DIR}. "
+          "Re-run the build to let `qola checkout` clone AITER.")
+endif()
+
+# Partition the requested HIP architectures into the CK-full set (CDNA, where
+# the AITER CK FMHA template headers compile) and the V3-asm-only set. gfx1250
+# (RDNA4) has AITER V3 *backward* asm kernels but no CK FMHA support and no
+# forward kernels, so it is built as a separate CK-free library (namespace
+# te_v3, manifest qola_manifest_gfx1250.toml) and dispatched at runtime in
+# ck_attn_bwd. The two tiers coexist via distinct QoLA namespaces.
+set(__CK_FULL_ARCHS ${CMAKE_HIP_ARCHITECTURES})
+set(__HAS_GFX1250 FALSE)
+list(FIND __CK_FULL_ARCHS "gfx1250" __GFX1250_IDX)
+if(NOT __GFX1250_IDX EQUAL -1)
+  set(__HAS_GFX1250 TRUE)
+  list(REMOVE_ITEM __CK_FULL_ARCHS "gfx1250")
+endif()
+list(LENGTH __CK_FULL_ARCHS __CK_FULL_ARCH_COUNT)
+if(__CK_FULL_ARCH_COUNT EQUAL 0 AND NOT __HAS_GFX1250)
+  message(FATAL_ERROR "ck_fused_attn: no target architectures requested.")
+endif()
+
+set(__AITER_MHA_PATH "")
+set(__HAVE_CK_FULL FALSE)
+
+# --- CK-full libraries (CDNA): te_libmha_fwd.so / te_libmha_bwd.so ---
+if(__CK_FULL_ARCH_COUNT GREATER 0)
+  set(__HAVE_CK_FULL TRUE)
+  if(DEFINED AITER_MHA_PATH)
+    message(STATUS "[AITER-BUILD] Using AITER_MHA_PATH=${AITER_MHA_PATH}")
+    # use pre-built te_libmha_fwd.so te_libmha_bwd.so
+    set(__AITER_MHA_PATH ${AITER_MHA_PATH})
+  else()
+    get_prebuilt_aiter(__AITER_MHA_PATH)
+
+    if(__AITER_MHA_PATH STREQUAL "")
+      # If not available, fallback: Build from source via QoLA
+      list(JOIN __CK_FULL_ARCHS ";" GPU_ARCHS_STR)
+      message(STATUS "[AITER-BUILD] Building CK-full AITER kernels for ${GPU_ARCHS_STR} via QoLA.")
+      set(__QOLA_BUILD_DIR "${__QOLA_DIR}/build")
+      execute_process(
+        COMMAND ${CMAKE_COMMAND} -E env "PYTHONPATH=${__QOLA_DIR}:$ENV{PYTHONPATH}"
+                ${Python_EXECUTABLE} -m qola.cli build
+                --manifest ${__QOLA_MANIFEST}
+                --aiter-root ${__AITER_SOURCE_DIR}
+                --output-dir ${__QOLA_BUILD_DIR}
+                --arch "${GPU_ARCHS_STR}"
+        RESULT_VARIABLE QOLA_BUILD_RESULT
+      )
+      if(NOT QOLA_BUILD_RESULT EQUAL 0)
+        message(FATAL_ERROR "[AITER-BUILD] QoLA build failed.")
+      endif()
+
+      # Copy the final .so libs and exported public headers into the aiter
+      # prebuilt cache so downstream consumers see a self-contained tree.
+      get_default_aiter_cache_dir(__QOLA_CACHE_DIR)
+      set(__QOLA_CACHE_LIB "${__QOLA_CACHE_DIR}/lib")
+      file(MAKE_DIRECTORY ${__QOLA_CACHE_LIB})
+      file(GLOB __QOLA_BUILT_LIBS "${__QOLA_BUILD_DIR}/lib/*.so")
+      file(COPY ${__QOLA_BUILT_LIBS} DESTINATION ${__QOLA_CACHE_LIB})
+      file(COPY "${__QOLA_BUILD_DIR}/include" DESTINATION "${__QOLA_CACHE_DIR}")
+      set(__AITER_MHA_PATH "${__QOLA_CACHE_LIB}")
+    else()
+      message(STATUS "[AITER-BUILD] Using pre-built AITER from ${__AITER_MHA_PATH}")
     endif()
+  endif()
+endif()
+
+# --- V3-asm-only backward library (gfx1250): te_v3_libmha_bwd.so ---
+# There is no prebuilt cache path for gfx1250 (no public prebuilt, and a CK-free
+# asm build is cheap), so always build it from source via QoLA. Both manifests
+# pin the same AITER commit and share the already-checked-out source tree.
+if(__HAS_GFX1250)
+  set(__QOLA_MANIFEST_V3 "${CMAKE_CURRENT_LIST_DIR}/qola_manifest_gfx1250.toml")
+  set(__QOLA_BUILD_DIR_V3 "${__QOLA_DIR}/build_gfx1250")
+  message(STATUS "[AITER-BUILD] Building CK-free V3 backward (gfx1250) via QoLA.")
+  # The asm-only / CK-free flags (ONLY_FAV3=1, ENABLE_CK=0) are carried by the
+  # gfx1250 manifest's libmha_bwd module, so no special env is needed here.
+  execute_process(
+    COMMAND ${CMAKE_COMMAND} -E env "PYTHONPATH=${__QOLA_DIR}:$ENV{PYTHONPATH}"
+            ${Python_EXECUTABLE} -m qola.cli build
+            --manifest ${__QOLA_MANIFEST_V3}
+            --aiter-root ${__AITER_SOURCE_DIR}
+            --output-dir ${__QOLA_BUILD_DIR_V3}
+            --arch "gfx1250"
+    RESULT_VARIABLE QOLA_V3_BUILD_RESULT
+  )
+  if(NOT QOLA_V3_BUILD_RESULT EQUAL 0)
+    message(FATAL_ERROR "[AITER-BUILD] QoLA gfx1250 V3 build failed.")
+  endif()
 
-    # Copy the final .so libs and exported public headers into the aiter
-    # prebuilt cache so downstream consumers see a self-contained tree.
+  # Stage the v3 lib next to the CK-full libs so a single link/-L/install path
+  # covers both. For a gfx1250-only build there are no CK-full libs, so set up
+  # the cache lib dir here and stage the v3 public headers too.
+  if(__AITER_MHA_PATH STREQUAL "")
     get_default_aiter_cache_dir(__QOLA_CACHE_DIR)
     set(__QOLA_CACHE_LIB "${__QOLA_CACHE_DIR}/lib")
     file(MAKE_DIRECTORY ${__QOLA_CACHE_LIB})
-    file(GLOB __QOLA_BUILT_LIBS "${__QOLA_BUILD_DIR}/lib/*.so")
-    file(COPY ${__QOLA_BUILT_LIBS} DESTINATION ${__QOLA_CACHE_LIB})
-    file(COPY "${__QOLA_BUILD_DIR}/include" DESTINATION "${__QOLA_CACHE_DIR}")
+    file(COPY "${__QOLA_BUILD_DIR_V3}/include" DESTINATION "${__QOLA_CACHE_DIR}")
     set(__AITER_MHA_PATH "${__QOLA_CACHE_LIB}")
-  else()
-    message(STATUS "[AITER-BUILD] Using pre-built AITER from ${__AITER_MHA_PATH}")
   endif()
+  file(GLOB __QOLA_V3_LIBS "${__QOLA_BUILD_DIR_V3}/lib/te_v3_*.so")
+  file(COPY ${__QOLA_V3_LIBS} DESTINATION ${__AITER_MHA_PATH})
 endif()
 
 set(ck_fused_attn_SOURCES)
@@ -124,7 +209,18 @@ endforeach()
 add_library(ck_fused_attn SHARED ${ck_fused_attn_SOURCES})
 set(CK_FUSED_ATTN_COMPILE_OPTIONS)
 list(APPEND CK_FUSED_ATTN_COMPILE_OPTIONS
-    -DCK_TILE_FLOAT_TO_BFLOAT16_DEFAULT=${CK_FUSED_ATTN_FLOAT_TO_BFLOAT16_DEFAULT})
+    -DCK_TILE_FLOAT_TO_BFLOAT16_DEFAULT=${CK_FUSED_ATTN_FLOAT_TO_BFLOAT16_DEFAULT}
+    -DENABLE_CK=1)
+# Tier guards consumed by src/ck_fused_attn_{fwd,bwd}.cpp:
+#   NVTE_AITER_CK_FULL        -> qola::te::{mha_fwd,mha_bwd} (CDNA) are linked
+#   NVTE_AITER_V3_BWD_GFX1250 -> qola::te_v3::mha_bwd (gfx1250) is linked +
+#                                runtime-dispatched in ck_attn_bwd
+if(__HAVE_CK_FULL)
+  list(APPEND CK_FUSED_ATTN_COMPILE_OPTIONS -DNVTE_AITER_CK_FULL)
+endif()
+if(__HAS_GFX1250)
+  list(APPEND CK_FUSED_ATTN_COMPILE_OPTIONS -DNVTE_AITER_V3_BWD_GFX1250)
+endif()
 
 # Public QoLA headers ship alongside the .so libs in ${__AITER_MHA_PATH}/../include
 # (emitted by qola.cli build, or copied from the QoLA build dir above for the
@@ -141,10 +237,22 @@ target_include_directories(ck_fused_attn PRIVATE ${__QOLA_INCLUDE_DIR})
 
 find_package(hip)
 target_link_directories(ck_fused_attn PUBLIC ${__AITER_MHA_PATH})
-list(APPEND ck_fused_attn_LINKER_LIBS hip::host hip::device roctx64 -l:te_libmha_fwd.so -l:te_libmha_bwd.so)
+list(APPEND ck_fused_attn_LINKER_LIBS hip::host hip::device roctx64)
+set(__INSTALL_AITER_LIBS)
+if(__HAVE_CK_FULL)
+  list(APPEND ck_fused_attn_LINKER_LIBS -l:te_libmha_fwd.so -l:te_libmha_bwd.so)
+  list(APPEND __INSTALL_AITER_LIBS
+       ${__AITER_MHA_PATH}/te_libmha_fwd.so
+       ${__AITER_MHA_PATH}/te_libmha_bwd.so)
+endif()
+if(__HAS_GFX1250)
+  list(APPEND ck_fused_attn_LINKER_LIBS -l:te_v3_libmha_bwd.so)
+  list(APPEND __INSTALL_AITER_LIBS
+       ${__AITER_MHA_PATH}/te_v3_libmha_bwd.so)
+endif()
 target_link_libraries(ck_fused_attn PUBLIC ${ck_fused_attn_LINKER_LIBS})
 target_compile_options(ck_fused_attn PRIVATE ${CK_FUSED_ATTN_COMPILE_OPTIONS})
 set_target_properties(ck_fused_attn PROPERTIES INSTALL_RPATH "$ORIGIN")
 
-install(FILES ${__AITER_MHA_PATH}/te_libmha_fwd.so ${__AITER_MHA_PATH}/te_libmha_bwd.so DESTINATION ${CMAKE_INSTALL_PREFIX}/${AITER_MHA_INSTALL_PREFIX}/lib)
+install(FILES ${__INSTALL_AITER_LIBS} DESTINATION ${CMAKE_INSTALL_PREFIX}/${AITER_MHA_INSTALL_PREFIX}/lib)
 install(TARGETS ck_fused_attn DESTINATION ${CMAKE_INSTALL_PREFIX}/${AITER_MHA_INSTALL_PREFIX}/lib)
@@ -18,8 +18,22 @@ string(STRIP "${ROCM_VER_CONTENT}" ROCM_VER_CONTENT)
 string(REGEX MATCH "^[0-9]+\\.[0-9]+" ROCM_VER "${ROCM_VER_CONTENT}")
 string(REGEX MATCH "^[0-9]+" ROCM_VER_MAJOR "${ROCM_VER}")
 
-# AITER commit
-get_git_commit("${__AITER_SOURCE_DIR}" AITER_SHA)
+# AITER commit — read from the QoLA manifest so the cache key tracks the
+# commit QoLA will actually check out and build, not whatever happens to be
+# the submodule's current HEAD at configure time.
+set(__QOLA_MANIFEST "${CMAKE_CURRENT_LIST_DIR}/qola_manifest.toml")
+set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS "${__QOLA_MANIFEST}")
+file(STRINGS "${__QOLA_MANIFEST}" __AITER_COMMIT_LINES
+     REGEX "^[ \t]*aiter_commit[ \t]*=[ \t]*\"[^\"]+\"")
+list(LENGTH __AITER_COMMIT_LINES __AITER_COMMIT_COUNT)
+if(NOT __AITER_COMMIT_COUNT EQUAL 1)
+  message(FATAL_ERROR
+          "Expected exactly one 'aiter_commit = \"...\"' line in "
+          "${__QOLA_MANIFEST}, found ${__AITER_COMMIT_COUNT}.")
+endif()
+list(GET __AITER_COMMIT_LINES 0 __AITER_COMMIT_LINE)
+string(REGEX MATCH "\"([^\"]+)\"" _UNUSED "${__AITER_COMMIT_LINE}")
+set(AITER_SHA "${CMAKE_MATCH_1}")
 
 # Cache key & local paths
 set(AITER_CACHE_ROOT "${CMAKE_CURRENT_LIST_DIR}/../../../build/aiter-prebuilts")

@@ -31,7 +31,7 @@ def parse_with_skip_comments(buffer, line, regex, outputs):
 
 
 def extract_fields_from_header(text: str, struct_name: str) -> List[str]:
-    struct_field_re = re.compile(r"([A-Za-z_][A-Za-z0-9_]*)\s*(?:=[^;]*)?;\s*$")
+    struct_field_re = re.compile(r"([A-Za-z_][A-Za-z0-9_]*)\s*(?:=[^;]*|\{[^;]*\})?;\s*$")
     struct_end_re = re.compile(r"^\s*};\s*$")
 
     struct_start_re = re.compile(rf"\bstruct\s+{re.escape(struct_name)}\b")
@@ -64,11 +64,14 @@ def main() -> int:
     parser = argparse.ArgumentParser(description="Check aiter args usage vs header definition")
     parser.add_argument("--mode", choices=["fwd", "bwd", "both"], default="both", help="Mode: fwd, bwd, or both")
     parser.add_argument("--te-dir", type=Path, default=Path(__file__).parent.parent.parent.parent, help="Root directory of TransformerEngine")
+    parser.add_argument("--aiter-root", type=Path, default=None,
+                        help="AITER source tree root. Defaults to <te-dir>/3rdparty/aiter.")
     args = parser.parse_args()
+    aiter_root = args.aiter_root if args.aiter_root else args.te_dir / "3rdparty/aiter"
     modes = ["fwd", "bwd"] if args.mode == "both" else [args.mode]
     mismatch = 0
     for mode in modes:
-        header_path = args.te_dir / f"3rdparty/aiter/csrc/include/mha_{mode}.h"
+        header_path = aiter_root / f"csrc/include/mha_{mode}.h"
         source_path = args.te_dir / f"transformer_engine/common/ck_fused_attn/src/ck_fused_attn_{mode}.cpp"
         header_text = header_path.read_text(encoding="utf-8")
         source_text = source_path.read_text(encoding="utf-8")

@@ -113,7 +113,6 @@ struct CkAttnBwdArgs : CKAttnCommonArgs {
   // dQ
   void* dq_ptr = nullptr;
   uint64_t stride_b_dq = 0, stride_h_dq = 0, stride_s_dq = 0;
-  void* dq_acc_ptr = nullptr;
 
   // dK / dV expanded (MQA/GQA reduction inputs; null when h==hg)
   void* dk_expanded_ptr = nullptr;

@@ -1,5 +1,5 @@
 [qola]
-aiter_commit = "33f2e6af5f39379c739720080ed0033d533f5cb2"    # pinned AITER submodule commit
+aiter_commit = "f03a4ec572bb3d9e15da3b346763c8f126feec0d"    # pinned AITER submodule commit
 namespace = "te"
 rocm_versions = ["7.2"]
 
@@ -9,9 +9,11 @@ architectures = ["gfx950", "gfx942"]
 [[modules]]
 name = "libmha_fwd"
 mode = "cpp_itfs"
+receipt = 700
 drop_srcs = ["mha_fwd_split.cu", "mha_fwd_batch_prefill.cu"]
 drop_directions = ["fwd_splitkv", "batch_prefill"]
 
 [[modules]]
 name = "libmha_bwd"
 mode = "cpp_itfs"
+receipt = 700
+1 −0		.github/CODEOWNERS
+11 −0		.gitignore
+0 −3		.gitmodules
+0 −1		3rdparty/aiter
+27 −6		CLAUDE.md
+196 −0		LICENSE
+56 −8		README.md
+6 −1		example/te-manifest.toml
+41 −0		patches/aiter/0001-jit-gpu-targets.patch
+101 −0		patches/aiter/0002-hsa-only-aiter-hip-common.patch
+114 −0		patches/aiter/0003-ck-te-receipt-700.patch
+42 −0		patches/aiter/README.md
+2 −1		qola/build_tools/__init__.py
+29 −3		qola/build_tools/builder.py
+60 −3		qola/build_tools/config.py
+244 −0		qola/build_tools/submodule.py
+70 −2		qola/cli.py
+28 −4		qola/cpp_itfs/registry.toml