ROCm · bnemanich · May 2, 2026 · Jan 26, 2026 · Jan 27, 2026 · Jan 27, 2026
@@ -88,6 +88,10 @@ if(HIPBLASLT_ENABLE_HOST)
     option(HIPBLASLT_ENABLE_HIPBLAS_DIRECT "Use the hipblas header directly." OFF)
 endif()
 
+# mxDataGenerator is used by hipblaslt clients and tensilelite client; default OFF on Windows
+cmake_dependent_option(HIPBLASLT_ENABLE_MXDATAGENERATOR "Use mxDataGenerator for MX format data generation (clients/tests)." ON "NOT WIN32" OFF)
+message(STATUS "Enable mxDataGenerator for MX format data generation: ${HIPBLASLT_ENABLE_MXDATAGENERATOR}")
+
 set(CMAKE_SKIP_BUILD_RPATH FALSE CACHE BOOL "Skip build RPATH")
 set(CMAKE_BUILD_WITH_INSTALL_RPATH FALSE CACHE BOOL "Build with install RPATH")
 set(CMAKE_INSTALL_RPATH "$ORIGIN/../lib:$ORIGIN/../llvm/lib" CACHE STRING "Install RPATH")
@@ -255,6 +259,48 @@ if(NOT ROCM_LIBS_SUPERBUILD)
     endif()
 endif()
 
+# INTERFACE library that owns the public hipblaslt API headers (the in-tree
+# `library/include` subtree plus its build-tree counterpart, where
+# `hipblaslt-export.h` and `hipblaslt-version.h` are generated). Routing the
+# include directories through a target lets consumers pick them up via the
+# build graph rather than via `target_include_directories(<tgt> BEFORE
+# PRIVATE .../library/include ...)`. `hip::host` is exposed as an INTERFACE
+# link because the in-tree headers `#include <hip/...>`.
+add_library(hipblaslt-headers INTERFACE)
+add_library(hipblaslt::headers ALIAS hipblaslt-headers)
+target_include_directories(hipblaslt-headers
+    INTERFACE
+        $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/library/include>
+        $<BUILD_INTERFACE:${CMAKE_CURRENT_BINARY_DIR}/library/include>
+)
+target_link_libraries(hipblaslt-headers INTERFACE hip::host)
+
+# `hipblaslt-version.h` is consumed by `library/include/hipblaslt/hipblaslt.h`
+# (only included when HIPBLASLT_ENABLE_HOST is ON). When HOST is ON, the
+# canonical `configure_file` call in `library/include/CMakeLists.txt`
+# regenerates the same file from the same template; this generation is
+# unconditional only because the file is cheap to produce.
+configure_file(
+    "${CMAKE_CURRENT_SOURCE_DIR}/library/include/hipblaslt-version.h.in"
+    "${CMAKE_CURRENT_BINARY_DIR}/library/include/hipblaslt/hipblaslt-version.h"
+)
+
+# `hipblaslt::mxdatagen` lives in clients/common (next to its source), and is
+# only built when at least one consumer needs it. tensilelite/client and
+# tensilelite/tests reach for it on the same condition, so we drive the
+# subdir from here to keep the dependency wiring close to the gate that
+# decides whether to build it at all.
+set(_hipblaslt_mxdatagen_enabled FALSE)
+if(HIPBLASLT_ENABLE_MXDATAGENERATOR
+        AND (HIPBLASLT_ENABLE_CLIENT
+             OR TENSILELITE_ENABLE_CLIENT
+             OR TENSILELITE_BUILD_TESTING))
+    set(_hipblaslt_mxdatagen_enabled TRUE)
+endif()
+if(_hipblaslt_mxdatagen_enabled)
+    add_subdirectory(clients/common)
+endif()
+
 add_subdirectory(tensilelite)
 
 if(HIPBLASLT_ENABLE_HOST)

@@ -29,19 +29,19 @@ target_link_libraries(hipblaslt-clients-common
         tensilelite::tensilelite-host
 )
 
+if(HIPBLASLT_ENABLE_MXDATAGENERATOR)
+    # PUBLIC because public headers such as testing_matmul.hpp guard
+    # `<mxDataGen.hpp>` on `HIPBLASLT_ENABLE_MXDATAGENERATOR`, so downstream
+    # targets (hipblaslt-test, hipblaslt-bench, ...) need to see the macro,
+    # the include path, and the C++20 floor all the way through.
+    target_link_libraries(hipblaslt-clients-common PUBLIC hipblaslt::mxdatagen)
+endif()
+
 if(HIPBLASLT_ENABLE_ROCROLLER)
-    if(NOT ROCM_LIBS_SUPERBUILD)
-        if(HIPBLASLT_ENABLE_THEROCK)
-            find_package(mxDataGenerator REQUIRED)
-        else()
-            add_subdirectory("${CMAKE_CURRENT_SOURCE_DIR}/../../../shared/mxdatagenerator" "${CMAKE_CURRENT_BINARY_DIR}/mxdatagenerator")
-        endif()
-    endif()
     target_compile_definitions(hipblaslt-clients-common PRIVATE HIPBLASLT_USE_ROCROLLER)
-    target_link_libraries(hipblaslt-clients-common PRIVATE roc::mxDataGenerator)
-    target_compile_features(hipblaslt-clients-common PRIVATE cxx_std_20)
 endif()
 
+
 if(HIPBLASLT_ENABLE_ASAN)
     hipblaslt_target_configure_sanitizers(hipblaslt-clients-common PUBLIC)
 endif()
@@ -59,7 +59,13 @@ if(HIPBLASLT_ENABLE_OPENMP)
     target_link_libraries(hipblaslt-bench PRIVATE OpenMP::OpenMP_CXX)
 endif()
 
-add_subdirectory(common)
+# `clients/common/CMakeLists.txt` itself only declares the mxdatagen helper
+# (which the project root already added when needed); the `include` and
+# `src` subdirectories populate hipblaslt-clients-common's sources, so they
+# are added directly here, avoiding a second `add_subdirectory(common)` that
+# would clash with the root-level one.
+add_subdirectory(common/include)
+add_subdirectory(common/src)
 add_subdirectory(bench)
 
 add_executable(hipblaslt-api-overhead "${CMAKE_CURRENT_SOURCE_DIR}/bench/src/client_api_overhead.cpp")

@@ -2,7 +2,7 @@
  *
  * MIT License
  *
- * Copyright (C) 2022-2025 Advanced Micro Devices, Inc.
+ * Copyright (C) 2022-2026 Advanced Micro Devices, Inc.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -969,16 +969,16 @@ try
     if(isBlockScaling(arg.scaleA))
     {
         if(arg.a_type != HIP_R_8F_E4M3 && arg.a_type != HIP_R_8F_E5M2
-           && arg.a_type != HIP_R_4F_E2M1_EXT && arg.a_type != HIP_R_6F_E2M3_EXT
-           && arg.a_type != HIP_R_6F_E3M2_EXT)
+           && arg.a_type != HIP_R_4F_E2M1 && arg.a_type != HIP_R_6F_E2M3
+           && arg.a_type != HIP_R_6F_E3M2)
             throw std::invalid_argument("Invalid a_type for block scaling format: "s
                                         + hip_datatype_to_string(arg.a_type));
     }
     if(isBlockScaling(arg.scaleB))
     {
         if(arg.b_type != HIP_R_8F_E4M3 && arg.b_type != HIP_R_8F_E5M2
-           && arg.b_type != HIP_R_4F_E2M1_EXT && arg.b_type != HIP_R_6F_E2M3_EXT
-           && arg.b_type != HIP_R_6F_E3M2_EXT)
+           && arg.b_type != HIP_R_4F_E2M1 && arg.b_type != HIP_R_6F_E2M3
+           && arg.b_type != HIP_R_6F_E3M2)
             throw std::invalid_argument("Invalid b_type for block scaling format: "s
                                         + hip_datatype_to_string(arg.b_type));
     }

@@ -1,2 +1,78 @@
-add_subdirectory(include)
-add_subdirectory(src)
+# Copyright Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier:  MIT
+
+# Owns the `hipblaslt::mxdatagen` STATIC helper. The source and header live
+# under this directory, so the target is defined here rather than at the root.
+#
+# This file is added by the project root via `add_subdirectory(clients/common)`
+# only when:
+#   * HIPBLASLT_ENABLE_MXDATAGENERATOR is ON, and
+#   * at least one consumer subtree is enabled (HIPBLASLT_ENABLE_CLIENT,
+#     TENSILELITE_ENABLE_CLIENT, or TENSILELITE_BUILD_TESTING).
+#
+# Consumers should `target_link_libraries(<tgt> PUBLIC hipblaslt::mxdatagen)`
+# (PRIVATE if no further consumer needs the macro). Linking propagates:
+#   * the include path for `<mxDataGen.hpp>`
+#   * the link to `roc::mxDataGenerator`
+#   * the C++20 requirement
+#   * the `HIPBLASLT_ENABLE_MXDATAGENERATOR` macro
+#   * `hipblaslt::headers`, which carries the in-tree `<hipblaslt/...>` API
+#     headers (the source pulls in `<hipblaslt/hipblaslt-export.h>` and
+#     `<hipblaslt/hipblaslt-types.h>`).
+
+if(NOT ROCM_LIBS_SUPERBUILD)
+    if(HIPBLASLT_ENABLE_THEROCK)
+        find_package(mxDataGenerator REQUIRED)
+    else()
+        # `${PROJECT_SOURCE_DIR}/../../shared/mxdatagenerator` resolves to the
+        # sibling `shared/mxdatagenerator` subtree of the project. The binary
+        # dir is anchored under the hipblaslt project's build tree to match
+        # the upstream subdirectory's existing layout assumptions.
+        add_subdirectory(
+            "${PROJECT_SOURCE_DIR}/../../shared/mxdatagenerator"
+            "${PROJECT_BINARY_DIR}/mxdatagenerator"
+        )
+    endif()
+endif()
+
+add_library(hipblaslt-mxdatagen STATIC
+    "${CMAKE_CURRENT_SOURCE_DIR}/src/mxDataGen.cpp"
+)
+add_library(hipblaslt::mxdatagen ALIAS hipblaslt-mxdatagen)
+
+target_include_directories(hipblaslt-mxdatagen
+    PUBLIC
+        $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>
+)
+
+target_link_libraries(hipblaslt-mxdatagen
+    PUBLIC
+        hipblaslt::headers
+        roc::mxDataGenerator
+    PRIVATE
+        # `-x hip` (via hip::device's INTERFACE_COMPILE_OPTIONS) makes the
+        # translation unit compile through hipcc/clang. mxDataGen.cpp itself
+        # does not launch kernels, but the modern ROCm bf16 header it pulls
+        # in transitively uses clang-only builtins (e.g. __builtin_elementwise_rint).
+        hip::device
+)
+
+target_compile_features(hipblaslt-mxdatagen PUBLIC cxx_std_20)
+target_compile_definitions(hipblaslt-mxdatagen PUBLIC HIPBLASLT_ENABLE_MXDATAGENERATOR)
+set_target_properties(hipblaslt-mxdatagen
+    PROPERTIES POSITION_INDEPENDENT_CODE ON
+)
+
+# `<hipblaslt/hipblaslt-export.h>` is generated by
+# `library/include/CMakeLists.txt`'s `generate_export_header(hipblaslt ...)`
+# call when HIPBLASLT_ENABLE_HOST=ON. In a tensilelite-only build that target
+# does not exist, so produce the same file from this helper instead. Both
+# code paths emit the file to the same location so `hipblaslt::headers`'s
+# `${CMAKE_BINARY_DIR}/library/include` interface dir resolves it either way.
+if(NOT HIPBLASLT_ENABLE_HOST)
+    include(GenerateExportHeader)
+    generate_export_header(hipblaslt-mxdatagen
+        BASE_NAME hipblaslt
+        EXPORT_FILE_NAME "${PROJECT_BINARY_DIR}/library/include/hipblaslt/hipblaslt-export.h"
+    )
+endif()
@@ -2,7 +2,7 @@
  *
  * MIT License
  *
- * Copyright (C) 2022-2025 Advanced Micro Devices, Inc.
+ * Copyright (C) 2022-2026 Advanced Micro Devices, Inc.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -86,8 +86,8 @@ inline std::size_t realDataTypeSize(hipDataType dtype)
 {
     // These types were not defined in older versions of ROCm, so need to be handled specially here.
     auto const dtype_int = static_cast<int>(dtype);
-    if(dtype_int == HIP_R_4F_E2M1_EXT || dtype_int == HIP_R_6F_E2M3_EXT
-       || dtype_int == HIP_R_6F_E3M2_EXT)
+    if(dtype_int == HIP_R_4F_E2M1 || dtype_int == HIP_R_6F_E2M3
+       || dtype_int == HIP_R_6F_E3M2)
     {
         return 1;
     }

@@ -2,7 +2,7 @@
  *
  * MIT License
  *
- * Copyright (C) 2022-2025 Advanced Micro Devices, Inc.
+ * Copyright (C) 2022-2026 Advanced Micro Devices, Inc.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -153,13 +153,13 @@ inline void hipblaslt_init(void*       A,
         hipblaslt_init<hipblasLtInt8>(
             static_cast<hipblasLtInt8*>(A), M, N, lda, stride, batch_count);
         break;
-    case HIP_R_6F_E2M3_EXT:
+    case HIP_R_6F_E2M3:
         hipblaslt_cerr << "hipblaslt_init not supports FP6" << std::endl;
         break;
-    case HIP_R_6F_E3M2_EXT:
+    case HIP_R_6F_E3M2:
         hipblaslt_cerr << "hipblaslt_init not supports BF6" << std::endl;
         break;
-    case HIP_R_4F_E2M1_EXT:
+    case HIP_R_4F_E2M1:
         hipblaslt_cerr << "hipblaslt_init not supports FP4" << std::endl;
         break;
     default:
@@ -263,13 +263,13 @@ inline void hipblaslt_init_sin(void*       A,
         hipblaslt_init_sin<hipblasLtInt8>(
             static_cast<hipblasLtInt8*>(A), M, N, lda, stride, batch_count);
         break;
-    case HIP_R_6F_E2M3_EXT:
+    case HIP_R_6F_E2M3:
         hipblaslt_cerr << "hipblaslt_init_sin not supports FP6" << std::endl;
         break;
-    case HIP_R_6F_E3M2_EXT:
+    case HIP_R_6F_E3M2:
         hipblaslt_cerr << "hipblaslt_init_sin not supports BF6" << std::endl;
         break;
-    case HIP_R_4F_E2M1_EXT:
+    case HIP_R_4F_E2M1:
         hipblaslt_cerr << "hipblaslt_init_sin not supports FP4" << std::endl;
         break;
     default:
@@ -353,13 +353,13 @@ inline void hipblaslt_init_alternating_sign(void*       A,
         hipblaslt_init_alternating_sign<hipblasLtInt8>(
             static_cast<hipblasLtInt8*>(A), M, N, lda, stride, batch_count);
         break;
-    case HIP_R_6F_E2M3_EXT:
+    case HIP_R_6F_E2M3:
         hipblaslt_cerr << "hipblaslt_init_alternating_sign not supports FP6" << std::endl;
         break;
-    case HIP_R_6F_E3M2_EXT:
+    case HIP_R_6F_E3M2:
         hipblaslt_cerr << "hipblaslt_init_alternating_sign not supports BF6" << std::endl;
         break;
-    case HIP_R_4F_E2M1_EXT:
+    case HIP_R_4F_E2M1:
         hipblaslt_cerr << "hipblaslt_init_alternating_sign not supports FP4" << std::endl;
         break;
     default:
@@ -440,13 +440,13 @@ inline void hipblaslt_init_hpl_alternating_sign(void*       A,
         hipblaslt_init_hpl_alternating_sign<hipblasLtInt8>(
             static_cast<hipblasLtInt8*>(A), M, N, lda, stride, batch_count);
         break;
-    case HIP_R_6F_E2M3_EXT:
+    case HIP_R_6F_E2M3:
         hipblaslt_cerr << "hipblaslt_init_hpl_alternating_sign not supports FP6" << std::endl;
         break;
-    case HIP_R_6F_E3M2_EXT:
+    case HIP_R_6F_E3M2:
         hipblaslt_cerr << "hipblaslt_init_hpl_alternating_sign not supports BF6" << std::endl;
         break;
-    case HIP_R_4F_E2M1_EXT:
+    case HIP_R_4F_E2M1:
         hipblaslt_cerr << "hipblaslt_init_hpl_alternating_sign not supports FP4" << std::endl;
         break;
     default:
@@ -521,13 +521,13 @@ inline void hipblaslt_init_cos(void*       A,
         hipblaslt_init_cos<hipblasLtInt8>(
             static_cast<hipblasLtInt8*>(A), M, N, lda, stride, batch_count);
         break;
-    case HIP_R_6F_E2M3_EXT:
+    case HIP_R_6F_E2M3:
         hipblaslt_cerr << "hipblaslt_init_cos not supports FP6" << std::endl;
         break;
-    case HIP_R_6F_E3M2_EXT:
+    case HIP_R_6F_E3M2:
         hipblaslt_cerr << "hipblaslt_init_cos not supports BF6" << std::endl;
         break;
-    case HIP_R_4F_E2M1_EXT:
+    case HIP_R_4F_E2M1:
         hipblaslt_cerr << "hipblaslt_init_cos not supports FP4" << std::endl;
         break;
     default:
@@ -608,13 +608,13 @@ inline void hipblaslt_init_hpl(void*       A,
         hipblaslt_init_hpl<hipblasLtInt8>(
             static_cast<hipblasLtInt8*>(A), M, N, lda, stride, batch_count);
         break;
-    case HIP_R_6F_E2M3_EXT:
+    case HIP_R_6F_E2M3:
         hipblaslt_cerr << "hipblaslt_init_hpl not supports FP6" << std::endl;
         break;
-    case HIP_R_6F_E3M2_EXT:
+    case HIP_R_6F_E3M2:
         hipblaslt_cerr << "hipblaslt_init_hpl not supports BF6" << std::endl;
         break;
-    case HIP_R_4F_E2M1_EXT:
+    case HIP_R_4F_E2M1:
         hipblaslt_cerr << "hipblaslt_init_hpl not supports FP4" << std::endl;
         break;
     default:
@@ -678,13 +678,13 @@ inline void hipblaslt_init_nan(void* A, size_t N, hipDataType type)
     case HIP_R_8I:
         hipblaslt_init_nan<hipblasLtInt8>(static_cast<hipblasLtInt8*>(A), N);
         break;
-    case HIP_R_6F_E2M3_EXT:
+    case HIP_R_6F_E2M3:
         hipblaslt_cerr << "hipblaslt_init_nan not supports FP6" << std::endl;
         break;
-    case HIP_R_6F_E3M2_EXT:
+    case HIP_R_6F_E3M2:
         hipblaslt_cerr << "hipblaslt_init_nan not supports BF6" << std::endl;
         break;
-    case HIP_R_4F_E2M1_EXT:
+    case HIP_R_4F_E2M1:
         hipblaslt_cerr << "hipblaslt_init_nan not supports FP4" << std::endl;
         break;
     default:
@@ -733,13 +733,13 @@ inline void hipblaslt_init_nan(void* A, size_t start_offset, size_t end_offset,
     case HIP_R_8I:
         hipblaslt_init_nan<hipblasLtInt8>(static_cast<hipblasLtInt8*>(A), start_offset, end_offset);
         break;
-    case HIP_R_6F_E2M3_EXT:
+    case HIP_R_6F_E2M3:
         hipblaslt_cerr << "hipblaslt_init_nan not supports FP6" << std::endl;
         break;
-    case HIP_R_6F_E3M2_EXT:
+    case HIP_R_6F_E3M2:
         hipblaslt_cerr << "hipblaslt_init_nan not supports BF6" << std::endl;
         break;
-    case HIP_R_4F_E2M1_EXT:
+    case HIP_R_4F_E2M1:
         hipblaslt_cerr << "hipblaslt_init_nan not supports FP4" << std::endl;
         break;
     default: