ck_tile: fix 2:4-sparse SWMMAC correctness on gfx1201/RDNA4 (3 bugs) + fail→pass repro

.gitignore

@@ -109,11 +109,6 @@ test_data/*
 # The experimental/builder directory should be tracked despite matching build*
 !experimental/builder
 !experimental/builder/**
-experimental/grouped_convolution_tile_instances/instances/*
-!experimental/grouped_convolution_tile_instances/instances/*.in
-!experimental/grouped_convolution_tile_instances/instances/*.inc
-!experimental/grouped_convolution_tile_instances/instances/*.hpp
-experimental/grouped_convolution_tile_instances/*.inc
 # Heuristics: benchmark data (never in git)
 dispatcher/heuristics/data/
 

.pre-commit-config.yaml

@@ -26,6 +26,16 @@ repos:
         verbose: false
         language: script
         types_or: [c++, python, shell, cmake]
+    -   id: ascii-only-checker
+        name: Check for non-ASCII characters in C/C++ sources
+        entry: projects/composablekernel/script/check_ascii_only.sh
+        language: script
+        types_or: [c++, inc]
+    -   id: crlf-checker
+        name: Check for CRLF line endings in C/C++ sources
+        entry: projects/composablekernel/script/check_no_crlf.sh
+        language: script
+        types_or: [c++, inc]
     -   id: remove-exec-bit
         name: Remove executable bit from non-executable files
         entry: projects/composablekernel/script/remove_exec_bit.sh

CHANGELOG.md

@@ -2,9 +2,10 @@
 
 Documentation for Composable Kernel available at [https://rocm.docs.amd.com/projects/composable_kernel/en/latest/](https://rocm.docs.amd.com/projects/composable_kernel/en/latest/).
 
-## (Unreleased) Composable Kernel 1.3.0
+## Composable Kernel 1.2.0 for ROCm 7.13
 
 ### Added
+
 * Added overload of load_tile_transpose that takes reference to output tensor as output parameter
 * Use data type from LDS tensor view when determining tile distribution for transpose in the GEMM pipeline
 * Added eightwarps support for abquant mode in blockscale GEMM.

CMakeLists.txt

@@ -136,6 +136,87 @@ else()
     set(CK_ENABLE_BF8 "ON")
 endif()
 
+# ---------------------------------------------------------------------------
+# Map GPU target strings to hex amdgcn_target_id values (arch.hpp).
+# Builds a -DCK_CMAKE_GPU_TARGET_IDS=0xHHHH,... definition that host-side
+# test code can consume without launching a device kernel.
+# ---------------------------------------------------------------------------
+function(_ck_gpu_target_string_to_id TARGET_STR OUT_VAR)
+    string(TOLOWER "${TARGET_STR}" _tgt)
+    string(REGEX REPLACE ":.*" "" _tgt "${_tgt}")
+    # GFX9
+    if(_tgt STREQUAL "gfx908")
+        set(${OUT_VAR} "0x0908" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx90a")
+        set(${OUT_VAR} "0x090A" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx942")
+        set(${OUT_VAR} "0x0942" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx950")
+        set(${OUT_VAR} "0x0950" PARENT_SCOPE)
+    # GFX10.3
+    elseif(_tgt STREQUAL "gfx1030")
+        set(${OUT_VAR} "0x1030" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1031")
+        set(${OUT_VAR} "0x1031" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1032")
+        set(${OUT_VAR} "0x1032" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1033")
+        set(${OUT_VAR} "0x1033" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1034")
+        set(${OUT_VAR} "0x1034" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1035")
+        set(${OUT_VAR} "0x1035" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1036")
+        set(${OUT_VAR} "0x1036" PARENT_SCOPE)
+    elseif(_tgt MATCHES "^gfx10-3-generic$")
+        set(${OUT_VAR} "0x103F" PARENT_SCOPE)
+    # GFX11
+    elseif(_tgt STREQUAL "gfx1100")
+        set(${OUT_VAR} "0x1100" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1101")
+        set(${OUT_VAR} "0x1101" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1102")
+        set(${OUT_VAR} "0x1102" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1103")
+        set(${OUT_VAR} "0x1103" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1150")
+        set(${OUT_VAR} "0x1150" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1151")
+        set(${OUT_VAR} "0x1151" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1152")
+        set(${OUT_VAR} "0x1152" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1153")
+        set(${OUT_VAR} "0x1153" PARENT_SCOPE)
+    elseif(_tgt MATCHES "^gfx11-generic$")
+        set(${OUT_VAR} "0x11FF" PARENT_SCOPE)
+    # GFX12
+    elseif(_tgt STREQUAL "gfx1200")
+        set(${OUT_VAR} "0x1200" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1201")
+        set(${OUT_VAR} "0x1201" PARENT_SCOPE)
+    elseif(_tgt MATCHES "^gfx12-generic$")
+        set(${OUT_VAR} "0x12FF" PARENT_SCOPE)
+    elseif(_tgt STREQUAL "gfx1250")
+        set(${OUT_VAR} "0x1250" PARENT_SCOPE)
+    else()
+        message(WARNING "_ck_gpu_target_string_to_id: unknown GPU target '${TARGET_STR}', skipping")
+        set(${OUT_VAR} "" PARENT_SCOPE)
+    endif()
+endfunction()
+
+# Inject a macro CK_CMAKE_GPU_TARGET_IDS containing all configure-level GPU targets into all source files.
+set(_ck_hex_ids)
+foreach(_tgt IN LISTS GPU_TARGETS)
+    _ck_gpu_target_string_to_id("${_tgt}" _ck_hex)
+    if(_ck_hex AND NOT _ck_hex STREQUAL "0x0000")
+        list(APPEND _ck_hex_ids "${_ck_hex}")
+    endif()
+endforeach()
+list(JOIN _ck_hex_ids "," _ck_hex_str)
+if(_ck_hex_str)
+    add_compile_definitions(CK_CMAKE_GPU_TARGET_IDS=${_ck_hex_str})
+endif()
+
 #for f8/bf8_t type
 add_compile_options(-Wno-bit-int-extension)
 add_compile_options(-Wno-pass-failed)
@@ -164,6 +245,8 @@ if(NOT DISABLE_DL_KERNELS AND GPU_TARGETS MATCHES "gfx101|gfx103|gfx10-1|gfx10-3
     set(DL_KERNELS "ON")
     set(CK_ENABLE_DL_KERNELS "ON")
 endif()
+message(STATUS "Disabling DPP kernels by default")
+option(DISABLE_DPP_KERNELS "Disable DPP kernels by default" ON)
 if(NOT DISABLE_DPP_KERNELS)
     add_definitions(-DDPP_KERNELS)
     set(DPP_KERNELS "ON")
@@ -741,6 +824,10 @@ SET(BUILD_DEV ON CACHE BOOL "BUILD_DEV")
 if(BUILD_DEV)
     add_compile_options(-Werror)
     add_compile_options(-Weverything)
+    add_compile_options(-Wno-lifetime-safety-intra-tu-suggestions)
+    add_compile_options(-Wno-lifetime-safety-cross-tu-suggestions)
+    add_compile_options(-Wno-lifetime-safety-lifetimebound-violation)
+    add_compile_options(-Wno-unknown-warning-option)
 endif()
 message(STATUS "CMAKE_CXX_FLAGS: ${CMAKE_CXX_FLAGS}")
 
@@ -761,8 +848,6 @@ if(NOT MIOPEN_REQ_LIBS_ONLY AND NOT HIPTENSOR_REQ_LIBS_ONLY)
 endif()
 
 
-option(MIOPEN_REQ_LIBS_ONLY "Build only the MIOpen required libraries" OFF)
-option(HIPTENSOR_REQ_LIBS_ONLY "Build only the HipTensor required libraries" OFF)
 option(DISABLE_OFFLOAD_COMPRESS "Disable offload compress compiler flag when building instances" OFF)
 option(BUILD_MHA_LIB "Build the static library for flash attention" OFF)
 option(BUILD_CK_DEVICE_INSTANCES "Build device operation instances in library/" ON)
@@ -822,7 +907,6 @@ endif()
 
 if (CK_EXPERIMENTAL_BUILDER)
   add_subdirectory(experimental/builder)
-  add_subdirectory(experimental/grouped_convolution_tile_instances)
 endif()
 
 if(NOT GPU_ARCHS AND USER_GPU_TARGETS AND NOT MIOPEN_REQ_LIBS_ONLY AND NOT HIPTENSOR_REQ_LIBS_ONLY)

Dockerfile

@@ -25,7 +25,7 @@ ENV HIP_PLATFORM=amd
 
 # Add rocm repository
 RUN set -xe && \
-    apt-get update && apt-get install -y --allow-unauthenticated apt-utils wget gnupg2 curl cmake git vim nano zip
+    apt-get update && apt-get install -y apt-utils wget gnupg2 curl cmake git vim nano zip
 
 RUN if [ "$compiler_version" = "therock" ]; then \
         rm -rf /opt/rocm && mkdir /opt/rocm && \
@@ -60,7 +60,7 @@ RUN set -x && \
     tar -xzf sccache.tar.gz --strip-components=1 -C ${SCCACHE_INSTALL_LOCATION} && \
     chmod +x ${SCCACHE_INSTALL_LOCATION}/sccache && \
 # Install dependencies
-    DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \
+    DEBIAN_FRONTEND=noninteractive apt-get install -y \
     build-essential \
     jq \
     libelf-dev \