From 4bd45072b9444902c9ec57ce4783a6a73cac33c7 Mon Sep 17 00:00:00 2001
From: Estevan Vedovelli <Estevan.Vedovelli@amd.com>
Date: Thu, 4 Jun 2026 10:51:59 -0400
Subject: [PATCH 01/10] [rocblas] Kernel Selection trace logging option

---
 projects/rocblas/CHANGELOG.md                 |   4 +
 projects/rocblas/clients/gtest/CMakeLists.txt |   1 +
 .../clients/gtest/kernel_select_gtest.cpp     | 341 ++++++++++++++++++
 .../clients/include/testing_logging.hpp       |  38 ++
 .../docs/data/reference/logging-env.rst       |   1 +
 .../docs/how-to/logging-in-rocblas.rst        |  20 +
 .../library/include/internal/rocblas-types.h  |   5 +
 .../src/blas3/Tensile/gemm_templates.cpp      |  79 ++++
 .../src/blas3/rocblas_gemm_batched_imp.hpp    |   3 +
 .../library/src/blas3/rocblas_gemm_imp.hpp    |   2 +
 .../rocblas_gemm_strided_batched_imp.hpp      |   3 +
 .../src/blas3/rocblas_hemm_batched_imp.hpp    |   2 +
 .../library/src/blas3/rocblas_hemm_imp.hpp    |   2 +
 .../rocblas_hemm_strided_batched_imp.hpp      |   2 +
 .../src/blas3/rocblas_her2k_batched_imp.hpp   |   2 +
 .../library/src/blas3/rocblas_her2k_imp.hpp   |   2 +
 .../rocblas_her2k_strided_batched_imp.hpp     |   2 +
 .../src/blas3/rocblas_herk_batched_imp.hpp    |   2 +
 .../library/src/blas3/rocblas_herk_imp.hpp    |   2 +
 .../rocblas_herk_strided_batched_imp.hpp      |   2 +
 .../src/blas3/rocblas_herkx_batched_imp.hpp   |   2 +
 .../library/src/blas3/rocblas_herkx_imp.hpp   |   2 +
 .../rocblas_herkx_strided_batched_imp.hpp     |   2 +
 .../src/blas3/rocblas_symm_batched_imp.hpp    |   2 +
 .../library/src/blas3/rocblas_symm_imp.hpp    |   2 +
 .../rocblas_symm_strided_batched_imp.hpp      |   2 +
 .../src/blas3/rocblas_syr2k_batched_imp.hpp   |   2 +
 .../library/src/blas3/rocblas_syr2k_imp.hpp   |   2 +
 .../rocblas_syr2k_strided_batched_imp.hpp     |   2 +
 .../src/blas3/rocblas_syrk_batched_imp.hpp    |   2 +
 .../library/src/blas3/rocblas_syrk_imp.hpp    |   2 +
 .../rocblas_syrk_strided_batched_imp.hpp      |   2 +
 .../src/blas3/rocblas_syrkx_batched_imp.hpp   |   2 +
 .../library/src/blas3/rocblas_syrkx_imp.hpp   |   2 +
 .../rocblas_syrkx_strided_batched_imp.hpp     |   2 +
 .../src/blas3/rocblas_trmm_batched_imp.hpp    |   2 +
 .../library/src/blas3/rocblas_trmm_imp.hpp    |   2 +
 .../rocblas_trmm_strided_batched_imp.hpp      |   2 +
 .../src/blas3/rocblas_trsm_batched_imp.hpp    |   2 +
 .../library/src/blas3/rocblas_trsm_imp.hpp    |   2 +
 .../rocblas_trsm_strided_batched_imp.hpp      |   2 +
 .../library/src/blas3/rocblas_trtri.cpp       |   2 +
 .../src/blas3/rocblas_trtri_batched.cpp       |   2 +
 .../blas3/rocblas_trtri_strided_batched.cpp   |   2 +
 .../blas_ex/rocblas_gemm_batched_ex_imp.hpp   |   2 +
 .../src/blas_ex/rocblas_gemm_ex_imp.hpp       |   2 +
 .../rocblas_gemm_strided_batched_ex_imp.hpp   |   3 +
 .../src/blas_ex/rocblas_gemmt_batched_imp.hpp |   2 +
 .../library/src/blas_ex/rocblas_gemmt_imp.hpp |   2 +
 .../rocblas_gemmt_strided_batched_imp.hpp     |   2 +
 projects/rocblas/library/src/handle.cpp       |   4 +-
 .../rocblas/library/src/hipblaslt_host.cpp    |  43 ++-
 .../rocblas/library/src/include/handle.hpp    |  32 ++
 .../library/src/include/hipblaslt_host.hpp    |   3 +-
 .../rocblas/library/src/include/logging.hpp   |  12 +
 projects/rocblas/library/src/tensile_host.cpp | 114 +++++-
 56 files changed, 768 insertions(+), 16 deletions(-)
 create mode 100644 projects/rocblas/clients/gtest/kernel_select_gtest.cpp

diff --git a/projects/rocblas/CHANGELOG.md b/projects/rocblas/CHANGELOG.md
index 09521afe4e53..53ec69ea1179 100644
--- a/projects/rocblas/CHANGELOG.md
+++ b/projects/rocblas/CHANGELOG.md
@@ -5,6 +5,10 @@ rocBLAS documentation is available at
 
 ## rocBLAS 5.5.0
 
+### Added
+
+* `ROCBLAS_LAYER=0x10` (`rocblas_layer_mode_log_kernel_select`): emits a `rocblas-bench`-replayable line per internal GEMM sub-problem, annotated with the selected kernel name/index and the source backend (`hipblaslt`, `tensile`, or `source`). Useful for decomposing higher-level BLAS3 calls (`trsm`, `trmm`, `syrk`, ...) into individually benchmarkable kernels without parsing `TENSILE_DB`. Lines also carry `fallback_from=` and `parent_api=` for context.
+
 ## rocBLAS 5.4.0
 
 ### Added
diff --git a/projects/rocblas/clients/gtest/CMakeLists.txt b/projects/rocblas/clients/gtest/CMakeLists.txt
index 6195ecff4b28..92c690111298 100644
--- a/projects/rocblas/clients/gtest/CMakeLists.txt
+++ b/projects/rocblas/clients/gtest/CMakeLists.txt
@@ -45,6 +45,7 @@ set(rocblas_no_tensile_test_source
     set_get_pointer_mode_gtest.cpp
     set_get_atomics_mode_gtest.cpp
     logging_mode_gtest.cpp
+    kernel_select_gtest.cpp
     ostream_threadsafety_gtest.cpp
     set_get_vector_gtest.cpp
     set_get_matrix_gtest.cpp
diff --git a/projects/rocblas/clients/gtest/kernel_select_gtest.cpp b/projects/rocblas/clients/gtest/kernel_select_gtest.cpp
new file mode 100644
index 000000000000..271325de23e2
--- /dev/null
+++ b/projects/rocblas/clients/gtest/kernel_select_gtest.cpp
@@ -0,0 +1,341 @@
+/* ************************************************************************
+ * Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell cop-
+ * ies of the Software, and to permit persons to whom the Software is furnished
+ * to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IM-
+ * PLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
+ * FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
+ * COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER
+ * IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNE-
+ * CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+ *
+ * ************************************************************************ */
+
+// Targeted tests for ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select).
+// Self-contained: drives the C API directly with hipMalloc/hipMemcpy, sets env vars
+// per-test, and parses the trace file to verify schema-level invariants. The selected
+// kernel name depends on GPU arch and backend build, so we never golden-file the line.
+
+#include "client_utility.hpp"
+#include "device_vector.hpp"
+#include "rocblas_test.hpp"
+#include <cstdio>
+#include <cstdlib>
+#include <fstream>
+#include <gtest/gtest.h>
+#include <hip/hip_runtime.h>
+#include <rocblas/rocblas.h>
+#include <sstream>
+#include <string>
+#ifndef WIN32
+#include <unistd.h>
+#endif
+
+#ifdef WIN32
+#define setenv(A, B, C) _putenv_s(A, B)
+#define unsetenv(A) _putenv_s(A, "")
+#endif
+
+namespace
+{
+    // Reads the entire trace file produced by a test into a single string for substring matching.
+    std::string read_trace(const std::string& path)
+    {
+        std::ifstream     ifs(path);
+        std::stringstream ss;
+        ss << ifs.rdbuf();
+        return ss.str();
+    }
+
+    // Counts occurrences of `needle` in `haystack`. Plain substring scan, no regex.
+    size_t count_occurrences(const std::string& haystack, const std::string& needle)
+    {
+        if(needle.empty())
+            return 0;
+        size_t n = 0;
+        for(size_t pos = 0; (pos = haystack.find(needle, pos)) != std::string::npos;
+            pos += needle.size())
+            ++n;
+        return n;
+    }
+
+    // Generates a unique trace-file path under /tmp for this test invocation.
+    std::string make_trace_path(const char* test_name)
+    {
+        std::string p = "/tmp/rocblas_kernel_select_";
+        p += test_name;
+        p += "_";
+        p += std::to_string(getpid());
+        p += ".log";
+        std::remove(p.c_str());
+        return p;
+    }
+
+    // RAII: snapshot the env vars we mutate, restore on scope exit so test ordering doesn't matter.
+    struct env_guard
+    {
+        env_guard()
+            : m_layer(getenv("ROCBLAS_LAYER") ? getenv("ROCBLAS_LAYER") : "")
+            , m_path(getenv("ROCBLAS_LOG_TRACE_PATH") ? getenv("ROCBLAS_LOG_TRACE_PATH") : "")
+            , m_uselt(getenv("ROCBLAS_USE_HIPBLASLT") ? getenv("ROCBLAS_USE_HIPBLASLT") : "")
+            , m_layer_set(getenv("ROCBLAS_LAYER") != nullptr)
+            , m_path_set(getenv("ROCBLAS_LOG_TRACE_PATH") != nullptr)
+            , m_uselt_set(getenv("ROCBLAS_USE_HIPBLASLT") != nullptr)
+        {
+        }
+        ~env_guard()
+        {
+            auto restore = [](const char* name, const std::string& val, bool was_set) {
+                if(was_set)
+                    setenv(name, val.c_str(), 1);
+                else
+                    unsetenv(name);
+            };
+            restore("ROCBLAS_LAYER", m_layer, m_layer_set);
+            restore("ROCBLAS_LOG_TRACE_PATH", m_path, m_path_set);
+            restore("ROCBLAS_USE_HIPBLASLT", m_uselt, m_uselt_set);
+        }
+
+    private:
+        std::string m_layer, m_path, m_uselt;
+        bool        m_layer_set, m_path_set, m_uselt_set;
+    };
+
+} // namespace
+
+// (1) trsm decomposition produces many sub-problem lines, all attributed to rocblas_dtrsm.
+// Covers: the rocblas_internal_api_scope at trsm's _impl, multi-line decomposition output,
+// and demonstrates the primary use-case the feature was designed for.
+TEST(KernelSelect, TrsmDecompositionParentApi)
+{
+    env_guard         guard;
+    const std::string trace_path = make_trace_path("trsm");
+    constexpr int     m          = 512;
+    constexpr int     n          = 512;
+    constexpr double  alpha      = 1.0;
+
+    setenv("ROCBLAS_LAYER", "0x10", 1);
+    setenv("ROCBLAS_LOG_TRACE_PATH", trace_path.c_str(), 1);
+
+    {
+        rocblas_handle handle;
+        ASSERT_EQ(rocblas_create_handle(&handle), rocblas_status_success);
+
+        device_vector<double> dA(size_t(m) * m);
+        device_vector<double> dB(size_t(m) * n);
+
+        ASSERT_EQ(rocblas_dtrsm(handle,
+                                rocblas_side_left,
+                                rocblas_fill_upper,
+                                rocblas_operation_transpose,
+                                rocblas_diagonal_non_unit,
+                                m,
+                                n,
+                                &alpha,
+                                dA,
+                                m,
+                                dB,
+                                m),
+                  rocblas_status_success);
+
+        ASSERT_EQ(rocblas_destroy_handle(handle), rocblas_status_success);
+    }
+
+    const std::string trace = read_trace(trace_path);
+    // trsm at this size decomposes into >>1 sub-problems; require multiple lines all carrying
+    // the trsm parent. (Exact count depends on internal block size, so we only lower-bound.)
+    const size_t lines = count_occurrences(trace, "parent_api=rocblas_dtrsm");
+    EXPECT_GT(lines, size_t(1)) << "expected multiple sub-problems attributed to rocblas_dtrsm, "
+                                << "got " << lines << " in " << trace_path;
+    EXPECT_NE(trace.find("-f gemm_strided_batched_ex"), std::string::npos);
+    EXPECT_NE(trace.find("# source="), std::string::npos);
+    EXPECT_NE(trace.find("kernel="), std::string::npos);
+
+    std::remove(trace_path.c_str());
+}
+
+// (2) Forcing ROCBLAS_USE_HIPBLASLT=0 should route every GEMM through Tensile.
+// Covers: source=tensile branch, capture of solution->name() / solution->index, and the
+// Tensile kernel-name format (Cijk_* is Tensile's canonical naming convention).
+TEST(KernelSelect, TensileBackendForced)
+{
+#ifndef BUILD_WITH_TENSILE
+    GTEST_SKIP() << "Tensile not built in; cannot test tensile backend path";
+#else
+    env_guard         guard;
+    const std::string trace_path = make_trace_path("tensile");
+    constexpr int     m          = 256;
+    constexpr int     n          = 256;
+    constexpr int     k          = 256;
+    constexpr float   alpha      = 1.0f;
+    constexpr float   beta       = 0.0f;
+
+    setenv("ROCBLAS_LAYER", "0x10", 1);
+    setenv("ROCBLAS_LOG_TRACE_PATH", trace_path.c_str(), 1);
+    setenv("ROCBLAS_USE_HIPBLASLT", "0", 1);
+
+    {
+        rocblas_handle handle;
+        ASSERT_EQ(rocblas_create_handle(&handle), rocblas_status_success);
+
+        device_vector<float> dA(size_t(m) * k);
+        device_vector<float> dB(size_t(k) * n);
+        device_vector<float> dC(size_t(m) * n);
+
+        ASSERT_EQ(rocblas_sgemm(handle,
+                                rocblas_operation_none,
+                                rocblas_operation_none,
+                                m,
+                                n,
+                                k,
+                                &alpha,
+                                dA,
+                                m,
+                                dB,
+                                k,
+                                &beta,
+                                dC,
+                                m),
+                  rocblas_status_success);
+
+        ASSERT_EQ(rocblas_destroy_handle(handle), rocblas_status_success);
+    }
+
+    const std::string trace = read_trace(trace_path);
+    EXPECT_NE(trace.find("source=tensile"), std::string::npos)
+        << "with ROCBLAS_USE_HIPBLASLT=0 expected source=tensile, got:\n"
+        << trace;
+    // Tensile solution names start with "Cijk_" by convention; if absent the capture is broken.
+    EXPECT_NE(trace.find("kernel=Cijk_"), std::string::npos)
+        << "expected Tensile-style kernel name, got:\n"
+        << trace;
+    EXPECT_EQ(trace.find("source=hipblaslt"), std::string::npos)
+        << "ROCBLAS_USE_HIPBLASLT=0 should suppress hipblaslt source, got:\n"
+        << trace;
+
+    std::remove(trace_path.c_str());
+#endif
+}
+
+// (3) Every emitted line must include the fallback_from= field, regardless of backend.
+// Covers: a weak schema check the existing logging test omits — guards against accidental
+// removal of the field or a code path that forgets to format it.
+TEST(KernelSelect, FallbackFromFieldAlwaysPresent)
+{
+    env_guard         guard;
+    const std::string trace_path = make_trace_path("fallback_field");
+    constexpr int     m          = 256;
+    constexpr int     n          = 256;
+    constexpr int     k          = 256;
+    constexpr float   alpha      = 1.0f;
+    constexpr float   beta       = 0.0f;
+
+    setenv("ROCBLAS_LAYER", "0x10", 1);
+    setenv("ROCBLAS_LOG_TRACE_PATH", trace_path.c_str(), 1);
+
+    {
+        rocblas_handle handle;
+        ASSERT_EQ(rocblas_create_handle(&handle), rocblas_status_success);
+
+        device_vector<float> dA(size_t(m) * k);
+        device_vector<float> dB(size_t(k) * n);
+        device_vector<float> dC(size_t(m) * n);
+
+        ASSERT_EQ(rocblas_sgemm(handle,
+                                rocblas_operation_none,
+                                rocblas_operation_none,
+                                m,
+                                n,
+                                k,
+                                &alpha,
+                                dA,
+                                m,
+                                dB,
+                                k,
+                                &beta,
+                                dC,
+                                m),
+                  rocblas_status_success);
+
+        ASSERT_EQ(rocblas_destroy_handle(handle), rocblas_status_success);
+    }
+
+    const std::string trace = read_trace(trace_path);
+    ASSERT_NE(trace.find("# source="), std::string::npos)
+        << "no kernel-select line emitted, trace was:\n"
+        << trace;
+    // Every kernel-select line should have the field — counting source= vs fallback_from= gives
+    // a uniform per-line check without per-line parsing.
+    const size_t source_lines   = count_occurrences(trace, "# source=");
+    const size_t fallback_lines = count_occurrences(trace, "fallback_from=");
+    EXPECT_EQ(source_lines, fallback_lines)
+        << "every kernel-select line must carry fallback_from=, got " << source_lines
+        << " source= but " << fallback_lines << " fallback_from=";
+
+    std::remove(trace_path.c_str());
+}
+
+// (4) Bits 0x8 and 0x10 are independent and can be combined; the same trace stream carries both.
+// Covers: env-parse OR of multiple bits, and the back-compat invariant that 0x8 still works.
+TEST(KernelSelect, CombinedWithInternalLayer)
+{
+    env_guard         guard;
+    const std::string trace_path = make_trace_path("combined");
+    constexpr int     m          = 256;
+    constexpr int     n          = 256;
+    constexpr int     k          = 256;
+    constexpr float   alpha      = 1.0f;
+    constexpr float   beta       = 0.0f;
+
+    setenv("ROCBLAS_LAYER", "0x18", 1); // 0x8 internal + 0x10 kernel-select
+    setenv("ROCBLAS_LOG_TRACE_PATH", trace_path.c_str(), 1);
+
+    {
+        rocblas_handle handle;
+        ASSERT_EQ(rocblas_create_handle(&handle), rocblas_status_success);
+
+        device_vector<float> dA(size_t(m) * k);
+        device_vector<float> dB(size_t(k) * n);
+        device_vector<float> dC(size_t(m) * n);
+
+        ASSERT_EQ(rocblas_sgemm(handle,
+                                rocblas_operation_none,
+                                rocblas_operation_none,
+                                m,
+                                n,
+                                k,
+                                &alpha,
+                                dA,
+                                m,
+                                dB,
+                                k,
+                                &beta,
+                                dC,
+                                m),
+                  rocblas_status_success);
+
+        ASSERT_EQ(rocblas_destroy_handle(handle), rocblas_status_success);
+    }
+
+    const std::string trace = read_trace(trace_path);
+    // 0x8 emits one of these backend tags per GEMM.
+    const bool has_0x8 = trace.find("rocblas_gemm_hipblaslt_backend") != std::string::npos
+                         || trace.find("rocblas_gemm_tensile_backend") != std::string::npos
+                         || trace.find("rocblas_gemm_source_backend") != std::string::npos;
+    EXPECT_TRUE(has_0x8) << "0x8 internal backend line missing from combined trace:\n" << trace;
+    // 0x10 emits the bench-style line.
+    EXPECT_NE(trace.find("# source="), std::string::npos)
+        << "0x10 kernel-select line missing from combined trace:\n"
+        << trace;
+
+    std::remove(trace_path.c_str());
+}
diff --git a/projects/rocblas/clients/include/testing_logging.hpp b/projects/rocblas/clients/include/testing_logging.hpp
index 2b258e066adf..6a97afc7ed29 100644
--- a/projects/rocblas/clients/include/testing_logging.hpp
+++ b/projects/rocblas/clients/include/testing_logging.hpp
@@ -651,6 +651,44 @@ void testing_logging(const Arguments& arg)
 
 #endif // BUILD_WITH_HIPBLASLT
 
+    // Smoke check for ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select). We don't
+    // golden-file the line because the selected kernel name varies by GPU arch and backend
+    // build — just assert at least one bench-style record with the expected schema was emitted.
+    if constexpr(!rocblas_is_complex<T>)
+    {
+        const fs::path kernel_select_fspath
+            = tmp_dir + std::string("kernel_select_") + test_run + std::string(".txt");
+        const std::string kernel_select_path = kernel_select_fspath.generic_string();
+
+        setenv_status = setenv("ROCBLAS_LAYER", "0x10", true);
+        ASSERT_EQ(setenv_status, 0);
+        setenv_status = setenv("ROCBLAS_LOG_TRACE_PATH", kernel_select_path.c_str(), true);
+        ASSERT_EQ(setenv_status, 0);
+
+        {
+            rocblas_local_handle handle_ks;
+            rocblas_set_pointer_mode(handle_ks, test_pointer_mode);
+            CHECK_ROCBLAS_ERROR(rocblas_gemm<T>(
+                handle_ks, transA, transB, m, n, k, alpha, da, lda, db, ldb, beta, dc, ldc));
+        } // handle dtor flushes the stream
+
+        std::ifstream ks_ifs(kernel_select_path);
+        ASSERT_TRUE(ks_ifs.good()) << "could not open " << kernel_select_path;
+        bool found = false;
+        for(std::string line; std::getline(ks_ifs, line);)
+        {
+            if(line.find("-f gemm_strided_batched_ex") != std::string::npos
+               && line.find("# source=") != std::string::npos
+               && line.find("kernel=") != std::string::npos
+               && line.find("parent_api=") != std::string::npos)
+            {
+                found = true;
+                break;
+            }
+        }
+        EXPECT_TRUE(found) << "ROCBLAS_LAYER=0x10 did not emit a kernel-select line";
+    }
+
     setenv_status = setenv("ROCBLAS_LAYER", "0", true);
     ASSERT_EQ(setenv_status, 0);
 
diff --git a/projects/rocblas/docs/data/reference/logging-env.rst b/projects/rocblas/docs/data/reference/logging-env.rst
index 55f25ad89d02..d0549366f82a 100644
--- a/projects/rocblas/docs/data/reference/logging-env.rst
+++ b/projects/rocblas/docs/data/reference/logging-env.rst
@@ -12,6 +12,7 @@
         | :code:`ROCBLAS_LAYER & 2 == 1`: Bench logging is enabled.
         | :code:`ROCBLAS_LAYER & 4 == 1`: Profile logging is enabled.
         | :code:`ROCBLAS_LAYER & 8 == 1`: Internal API logging is enabled.
+        | :code:`ROCBLAS_LAYER & 0x10 == 1`: Kernel-selection logging is enabled.
 
     * - | ``ROCBLAS_LOG_PATH``
         | Sets the full path for logging.
diff --git a/projects/rocblas/docs/how-to/logging-in-rocblas.rst b/projects/rocblas/docs/how-to/logging-in-rocblas.rst
index e060014265c1..881a87ff1063 100644
--- a/projects/rocblas/docs/how-to/logging-in-rocblas.rst
+++ b/projects/rocblas/docs/how-to/logging-in-rocblas.rst
@@ -41,6 +41,26 @@ of the function.
 Internal API logging outputs information like the GEMM backend used for a particular GEMM call.
 Not all internal APIs are logged. The log output goes to the same stream as trace logging.
 
+Kernel-selection logging (``ROCBLAS_LAYER=0x10``) emits one ``rocblas-bench``-replayable line per
+internal GEMM sub-problem dispatched, including those produced by decomposed BLAS3 ops such as
+``trsm``, ``trmm``, ``syrk``, and friends. Each line ends in a ``#`` comment that names the source
+backend (``hipblaslt``, ``tensile``, or ``source``), the selected kernel identifier, whether the
+call fell back from another backend, and the user-facing API that triggered it. The line itself is
+copy-pasteable into ``rocblas-bench`` to reproduce that single sub-problem in isolation. Output
+goes to the same stream as trace logging. Example:
+
+.. code-block:: text
+
+   ./rocblas-bench -f gemm_strided_batched_ex --transposeA N --transposeB N -m 4096 -n 4096 -k 4096 \
+     --alpha 1 --a_type f32_r --lda 4096 --stride_a 0 --b_type f32_r --ldb 4096 --stride_b 0 \
+     --beta 0 --c_type f32_r --ldc 4096 --stride_c 0 --d_type f32_r --ldd 4096 --stride_d 0 \
+     --batch_count 1 --compute_type f32_r --algo 0 --solution_index 0 --flags 0 \
+     # source=tensile kernel=Cijk_Ailk_Bljk_SB_MT128x128x8 tensile_index=42 \
+     fallback_from=hipblaslt parent_api=rocblas_strsm
+
+Bits can be combined: ``ROCBLAS_LAYER=0x19`` enables trace (``0x1``), internal (``0x8``), and
+kernel-selection (``0x10``) simultaneously.
+
 The default stream for logging output is standard error. :ref:`Four
 environment variables <rocblas_logging_env>` can set the full path name for a
 log file.
diff --git a/projects/rocblas/library/include/internal/rocblas-types.h b/projects/rocblas/library/include/internal/rocblas-types.h
index ec3a828ec26d..f3168d901682 100644
--- a/projects/rocblas/library/include/internal/rocblas-types.h
+++ b/projects/rocblas/library/include/internal/rocblas-types.h
@@ -256,6 +256,11 @@ typedef enum rocblas_layer_mode_
     rocblas_layer_mode_log_profile = 0x4,
     /*! \brief Outputs to the same stream as trace logging with limited internal API details, such as the GEMM backend that was used. */
     rocblas_layer_mode_log_internal = 0x8,
+    /*! \brief Outputs to the same stream as trace logging a rocblas-bench-replayable line for each
+     * internal GEMM sub-problem dispatched (including those produced by decomposed BLAS3 ops such as
+     * trsm, trmm, syrk, ...), annotated with the selected kernel name/index and the source backend
+     * (hipblaslt, tensile, or source). */
+    rocblas_layer_mode_log_kernel_select = 0x10,
 } rocblas_layer_mode;
 
 /*! \brief Indicates whether the layer is active by means of a bitmask.*/
diff --git a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
index 6c2439a8476a..3ca66b710977 100644
--- a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
+++ b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
@@ -31,7 +31,9 @@
 
 #include "check_numerics_matrix.hpp"
 #include "handle.hpp"
+#include "int64_helpers.hpp"
 #include "logging.hpp"
+#include "utility.hpp"
 
 /*
  * ===========================================================================
@@ -164,6 +166,83 @@ rocblas_status rocblas_internal_gemm(rocblas_handle    handle,
                          ldc);
     }
 
+    // ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select): bench-style line for the
+    // source-GEMM fallback path. Reached when Tensile returned not_implemented (or BUILD_WITH_TENSILE
+    // is off), so fallback_from is "tensile" when Tensile was built in, else "none".
+    if(handle->layer_mode & rocblas_layer_mode_log_kernel_select)
+    {
+#ifdef BUILD_WITH_TENSILE
+        const char* fallback = "tensile";
+#else
+        const char* fallback = "none";
+#endif
+        const char* parent_api
+            = handle->current_api_name ? handle->current_api_name : "unknown";
+
+        rocblas_internal_ostream kernel_field;
+        kernel_field << "kernel=rocblas_gemm_source_solution_64";
+        rocblas_internal_ostream source_field;
+        source_field << "# source=source";
+        rocblas_internal_ostream fallback_field;
+        fallback_field << "fallback_from=" << fallback;
+        rocblas_internal_ostream parent_field;
+        parent_field << "parent_api=" << parent_api;
+
+        rocblas_internal_logger logger;
+        logger.log_kernel_select(handle,
+                                 ROCBLAS_API_BENCH " -f gemm_strided_batched_ex",
+                                 "--transposeA",
+                                 rocblas_transpose_letter(trans_a),
+                                 "--transposeB",
+                                 rocblas_transpose_letter(trans_b),
+                                 "-m",
+                                 m,
+                                 "-n",
+                                 n,
+                                 "-k",
+                                 k,
+                                 LOG_BENCH_SCALAR_VALUE(handle, alpha),
+                                 "--a_type",
+                                 rocblas_precision_string<TScal>,
+                                 "--lda",
+                                 lda,
+                                 "--stride_a",
+                                 stride_a,
+                                 "--b_type",
+                                 rocblas_precision_string<TScal>,
+                                 "--ldb",
+                                 ldb,
+                                 "--stride_b",
+                                 stride_b,
+                                 LOG_BENCH_SCALAR_VALUE(handle, beta),
+                                 "--c_type",
+                                 rocblas_precision_string<TScal>,
+                                 "--ldc",
+                                 ldc,
+                                 "--stride_c",
+                                 stride_c,
+                                 "--d_type",
+                                 rocblas_precision_string<TScal>,
+                                 "--ldd",
+                                 ldc,
+                                 "--stride_d",
+                                 stride_c,
+                                 "--batch_count",
+                                 batch_count,
+                                 "--compute_type",
+                                 rocblas_precision_string<TScal>,
+                                 "--algo",
+                                 0,
+                                 "--solution_index",
+                                 0,
+                                 "--flags",
+                                 0,
+                                 source_field.str(),
+                                 kernel_field.str(),
+                                 fallback_field.str(),
+                                 parent_field.str());
+    }
+
     if(k == 0 || (alpha && *alpha == 0))
     {
         return rocblas_gemm_scale_launcher_64(
diff --git a/projects/rocblas/library/src/blas3/rocblas_gemm_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_gemm_batched_imp.hpp
index f0ee90d71eff..73737cc210f7 100644
--- a/projects/rocblas/library/src/blas3/rocblas_gemm_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_gemm_batched_imp.hpp
@@ -65,6 +65,9 @@ namespace
     {
         if(!handle)
             return rocblas_status_invalid_handle;
+
+        rocblas_internal_api_scope api_scope(handle, rocblas_gemm_batched_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device
diff --git a/projects/rocblas/library/src/blas3/rocblas_gemm_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_gemm_imp.hpp
index 62f582c48bfd..45fad00aeafc 100644
--- a/projects/rocblas/library/src/blas3/rocblas_gemm_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_gemm_imp.hpp
@@ -62,6 +62,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_gemm_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device
diff --git a/projects/rocblas/library/src/blas3/rocblas_gemm_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_gemm_strided_batched_imp.hpp
index 98b3265196a2..3da749112860 100644
--- a/projects/rocblas/library/src/blas3/rocblas_gemm_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_gemm_strided_batched_imp.hpp
@@ -76,6 +76,9 @@ namespace
     {
         if(!handle)
             return rocblas_status_invalid_handle;
+
+        rocblas_internal_api_scope api_scope(handle, rocblas_gemm_strided_batched_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device
diff --git a/projects/rocblas/library/src/blas3/rocblas_hemm_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_hemm_batched_imp.hpp
index 002ed9994fc7..73d65166cf80 100644
--- a/projects/rocblas/library/src/blas3/rocblas_hemm_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_hemm_batched_imp.hpp
@@ -55,6 +55,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_hemm_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         auto                    layer_mode     = handle->layer_mode;
diff --git a/projects/rocblas/library/src/blas3/rocblas_hemm_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_hemm_imp.hpp
index aa6ae6f5b0df..25128ce939de 100644
--- a/projects/rocblas/library/src/blas3/rocblas_hemm_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_hemm_imp.hpp
@@ -53,6 +53,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_hemm_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         auto                    layer_mode     = handle->layer_mode;
diff --git a/projects/rocblas/library/src/blas3/rocblas_hemm_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_hemm_strided_batched_imp.hpp
index 24a87ae20cb5..1359540a4ad2 100644
--- a/projects/rocblas/library/src/blas3/rocblas_hemm_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_hemm_strided_batched_imp.hpp
@@ -58,6 +58,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_hemm_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         auto                    layer_mode     = handle->layer_mode;
diff --git a/projects/rocblas/library/src/blas3/rocblas_her2k_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_her2k_batched_imp.hpp
index 48860a274661..3afef70c61dd 100644
--- a/projects/rocblas/library/src/blas3/rocblas_her2k_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_her2k_batched_imp.hpp
@@ -55,6 +55,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_her2k_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_her2k_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_her2k_imp.hpp
index 64c9d04ad221..9884a85d4420 100644
--- a/projects/rocblas/library/src/blas3/rocblas_her2k_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_her2k_imp.hpp
@@ -52,6 +52,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_her2k_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_her2k_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_her2k_strided_batched_imp.hpp
index 7fd998413930..4ea3e8d63637 100644
--- a/projects/rocblas/library/src/blas3/rocblas_her2k_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_her2k_strided_batched_imp.hpp
@@ -58,6 +58,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_her2k_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_herk_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_herk_batched_imp.hpp
index d10908955264..9e43832cd7ea 100644
--- a/projects/rocblas/library/src/blas3/rocblas_herk_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_herk_batched_imp.hpp
@@ -56,6 +56,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_herk_name<T>);
+
         //Check if the handle is in the device memory size query, as there are two algorithms one which requires extra workspace memory and one which doesn't
         if(handle->is_device_memory_size_query())
         {
diff --git a/projects/rocblas/library/src/blas3/rocblas_herk_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_herk_imp.hpp
index 734ee3a4a714..29c009aec660 100644
--- a/projects/rocblas/library/src/blas3/rocblas_herk_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_herk_imp.hpp
@@ -52,6 +52,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_herk_name<T>);
+
         //Check if the handle is in the device memory size query, as there are two algorithms one which requires extra workspace memory and one which doesn't
         if(handle->is_device_memory_size_query())
         {
diff --git a/projects/rocblas/library/src/blas3/rocblas_herk_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_herk_strided_batched_imp.hpp
index b08d21ee309a..8df70f657bdb 100644
--- a/projects/rocblas/library/src/blas3/rocblas_herk_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_herk_strided_batched_imp.hpp
@@ -57,6 +57,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_herk_name<T>);
+
         //Check if the handle is in the device memory size query, as there are two algorithms one which requires extra workspace memory and one which doesn't
         if(handle->is_device_memory_size_query())
         {
diff --git a/projects/rocblas/library/src/blas3/rocblas_herkx_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_herkx_batched_imp.hpp
index 2644cb6c1a7b..2a34e8c0edff 100644
--- a/projects/rocblas/library/src/blas3/rocblas_herkx_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_herkx_batched_imp.hpp
@@ -55,6 +55,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_herkx_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_herkx_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_herkx_imp.hpp
index 5d1d71db69f0..320520a6d0f8 100644
--- a/projects/rocblas/library/src/blas3/rocblas_herkx_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_herkx_imp.hpp
@@ -52,6 +52,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_herkx_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_herkx_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_herkx_strided_batched_imp.hpp
index 302700d540aa..7cb9e81d6a01 100644
--- a/projects/rocblas/library/src/blas3/rocblas_herkx_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_herkx_strided_batched_imp.hpp
@@ -58,6 +58,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_herkx_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_symm_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_symm_batched_imp.hpp
index 1587a2fcf70e..be2587b9cfc9 100644
--- a/projects/rocblas/library/src/blas3/rocblas_symm_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_symm_batched_imp.hpp
@@ -59,6 +59,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_symm_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         auto                    layer_mode     = handle->layer_mode;
diff --git a/projects/rocblas/library/src/blas3/rocblas_symm_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_symm_imp.hpp
index 278ca9c5e92b..180288bdd366 100644
--- a/projects/rocblas/library/src/blas3/rocblas_symm_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_symm_imp.hpp
@@ -56,6 +56,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_symm_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         auto                    layer_mode     = handle->layer_mode;
diff --git a/projects/rocblas/library/src/blas3/rocblas_symm_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_symm_strided_batched_imp.hpp
index 2ae6dc5b30e9..08941b5910d9 100644
--- a/projects/rocblas/library/src/blas3/rocblas_symm_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_symm_strided_batched_imp.hpp
@@ -62,6 +62,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_symm_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         auto                    layer_mode     = handle->layer_mode;
diff --git a/projects/rocblas/library/src/blas3/rocblas_syr2k_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_syr2k_batched_imp.hpp
index 6fb3e6989f8c..b5ce79fc4d54 100644
--- a/projects/rocblas/library/src/blas3/rocblas_syr2k_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_syr2k_batched_imp.hpp
@@ -59,6 +59,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_syr2k_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_syr2k_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_syr2k_imp.hpp
index c25502cffb46..31c597835091 100644
--- a/projects/rocblas/library/src/blas3/rocblas_syr2k_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_syr2k_imp.hpp
@@ -56,6 +56,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_syr2k_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_syr2k_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_syr2k_strided_batched_imp.hpp
index 12c3e0b7c0be..7c2b1df6c461 100644
--- a/projects/rocblas/library/src/blas3/rocblas_syr2k_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_syr2k_strided_batched_imp.hpp
@@ -62,6 +62,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_syr2k_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_syrk_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_syrk_batched_imp.hpp
index 92e9caad919f..ec533a1fd400 100644
--- a/projects/rocblas/library/src/blas3/rocblas_syrk_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_syrk_batched_imp.hpp
@@ -57,6 +57,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_syrk_name<T>);
+
         //Check if the handle is in the device memory size query, as there are two algorithms one which requires extra workspace memory and one which doesn't
         if(handle->is_device_memory_size_query())
         {
diff --git a/projects/rocblas/library/src/blas3/rocblas_syrk_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_syrk_imp.hpp
index 7ef8f855cd5b..0b318600bcf8 100644
--- a/projects/rocblas/library/src/blas3/rocblas_syrk_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_syrk_imp.hpp
@@ -55,6 +55,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_syrk_name<T>);
+
         //Check if the handle is in the device memory size query, as there are two algorithms one which requires extra workspace memory and one which doesn't
         if(handle->is_device_memory_size_query())
         {
diff --git a/projects/rocblas/library/src/blas3/rocblas_syrk_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_syrk_strided_batched_imp.hpp
index 1cc6ba92c38b..ee67dc0397d7 100644
--- a/projects/rocblas/library/src/blas3/rocblas_syrk_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_syrk_strided_batched_imp.hpp
@@ -60,6 +60,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_syrk_name<T>);
+
         //Check if the handle is in the device memory size query, as there are two algorithms one which requires extra workspace memory and one which doesn't
         if(handle->is_device_memory_size_query())
         {
diff --git a/projects/rocblas/library/src/blas3/rocblas_syrkx_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_syrkx_batched_imp.hpp
index dc01f7b5b41f..a7c5a8a55cd9 100644
--- a/projects/rocblas/library/src/blas3/rocblas_syrkx_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_syrkx_batched_imp.hpp
@@ -59,6 +59,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_syrkx_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_syrkx_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_syrkx_imp.hpp
index 918c5c616d04..fd4c523a9371 100644
--- a/projects/rocblas/library/src/blas3/rocblas_syrkx_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_syrkx_imp.hpp
@@ -56,6 +56,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_syrkx_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_syrkx_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_syrkx_strided_batched_imp.hpp
index 15af2bd0cca7..1aa04523cb53 100644
--- a/projects/rocblas/library/src/blas3/rocblas_syrkx_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_syrkx_strided_batched_imp.hpp
@@ -62,6 +62,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_syrkx_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         // Copy alpha and beta to host if on device. This is because gemm is called and it
diff --git a/projects/rocblas/library/src/blas3/rocblas_trmm_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_trmm_batched_imp.hpp
index 5b55678916c2..37b605fa335a 100644
--- a/projects/rocblas/library/src/blas3/rocblas_trmm_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_trmm_batched_imp.hpp
@@ -60,6 +60,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_trmm_batched_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         T        alpha_h, beta_h;
diff --git a/projects/rocblas/library/src/blas3/rocblas_trmm_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_trmm_imp.hpp
index 36535e79894b..b0ee0ea05f23 100644
--- a/projects/rocblas/library/src/blas3/rocblas_trmm_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_trmm_imp.hpp
@@ -57,6 +57,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_trmm_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         T        alpha_h, beta_h;
diff --git a/projects/rocblas/library/src/blas3/rocblas_trmm_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_trmm_strided_batched_imp.hpp
index 6e51de189e84..c3fadf488377 100644
--- a/projects/rocblas/library/src/blas3/rocblas_trmm_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_trmm_strided_batched_imp.hpp
@@ -65,6 +65,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_trmm_strided_batched_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         T        alpha_h, beta_h;
diff --git a/projects/rocblas/library/src/blas3/rocblas_trsm_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_trsm_batched_imp.hpp
index b2bf8375a674..8ca9879a6377 100644
--- a/projects/rocblas/library/src/blas3/rocblas_trsm_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_trsm_batched_imp.hpp
@@ -73,6 +73,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_trsm_name<T>);
+
         auto                    check_numerics = handle->check_numerics;
         rocblas_internal_logger logger;
         /////////////
diff --git a/projects/rocblas/library/src/blas3/rocblas_trsm_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_trsm_imp.hpp
index 2a566969e428..01290a44d230 100644
--- a/projects/rocblas/library/src/blas3/rocblas_trsm_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_trsm_imp.hpp
@@ -70,6 +70,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_trsm_name<T>);
+
         auto                    check_numerics = handle->check_numerics;
         rocblas_internal_logger logger;
         /////////////
diff --git a/projects/rocblas/library/src/blas3/rocblas_trsm_strided_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_trsm_strided_batched_imp.hpp
index 3f91c509c1c6..924bc0ba08eb 100644
--- a/projects/rocblas/library/src/blas3/rocblas_trsm_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_trsm_strided_batched_imp.hpp
@@ -76,6 +76,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_trsm_name<T>);
+
         auto                    check_numerics = handle->check_numerics;
         rocblas_internal_logger logger;
         /////////////
diff --git a/projects/rocblas/library/src/blas3/rocblas_trtri.cpp b/projects/rocblas/library/src/blas3/rocblas_trtri.cpp
index c6d061890117..cde952658184 100644
--- a/projects/rocblas/library/src/blas3/rocblas_trtri.cpp
+++ b/projects/rocblas/library/src/blas3/rocblas_trtri.cpp
@@ -51,6 +51,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_trtri_name<T>);
+
         size_t size = rocblas_internal_trtri_temp_elements(n, 1) * sizeof(T);
         if(handle->is_device_memory_size_query())
         {
diff --git a/projects/rocblas/library/src/blas3/rocblas_trtri_batched.cpp b/projects/rocblas/library/src/blas3/rocblas_trtri_batched.cpp
index 198c71ce503d..f98cd20c345b 100644
--- a/projects/rocblas/library/src/blas3/rocblas_trtri_batched.cpp
+++ b/projects/rocblas/library/src/blas3/rocblas_trtri_batched.cpp
@@ -53,6 +53,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_trtri_name<T>);
+
         // Compute the optimal size for temporary device memory
         size_t els   = rocblas_internal_trtri_temp_elements(n, 1);
         size_t size  = els * batch_count * sizeof(T);
diff --git a/projects/rocblas/library/src/blas3/rocblas_trtri_strided_batched.cpp b/projects/rocblas/library/src/blas3/rocblas_trtri_strided_batched.cpp
index 199f77d62805..4a3b50d78bc6 100644
--- a/projects/rocblas/library/src/blas3/rocblas_trtri_strided_batched.cpp
+++ b/projects/rocblas/library/src/blas3/rocblas_trtri_strided_batched.cpp
@@ -55,6 +55,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_trtri_name<T>);
+
         // Compute the optimal size for temporary device memory
         size_t size = rocblas_internal_trtri_temp_elements(n, batch_count) * sizeof(T);
         if(handle->is_device_memory_size_query())
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_gemm_batched_ex_imp.hpp b/projects/rocblas/library/src/blas_ex/rocblas_gemm_batched_ex_imp.hpp
index e67b38c375e0..c09e0e34ca5e 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_gemm_batched_ex_imp.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_gemm_batched_ex_imp.hpp
@@ -60,6 +60,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, ROCBLAS_API_STR(rocblas_gemm_batched_ex));
+
         const bool HPA = compute_type == rocblas_datatype_f32_r
                          && (a_type == rocblas_datatype_f16_r || a_type == rocblas_datatype_bf16_r);
 
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex_imp.hpp b/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex_imp.hpp
index 1bd9b0bd65af..2fb4f7fb3fbf 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex_imp.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex_imp.hpp
@@ -59,6 +59,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, ROCBLAS_API_STR(rocblas_gemm_ex));
+
         const bool HPA = compute_type == rocblas_datatype_f32_r
                          && (a_type == rocblas_datatype_f16_r || a_type == rocblas_datatype_bf16_r);
 
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_gemm_strided_batched_ex_imp.hpp b/projects/rocblas/library/src/blas_ex/rocblas_gemm_strided_batched_ex_imp.hpp
index 6321eac3e62e..db9d01c9f139 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_gemm_strided_batched_ex_imp.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_gemm_strided_batched_ex_imp.hpp
@@ -63,6 +63,9 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle,
+                                             ROCBLAS_API_STR(rocblas_gemm_strided_batched_ex));
+
         const bool HPA = compute_type == rocblas_datatype_f32_r
                          && (a_type == rocblas_datatype_f16_r || a_type == rocblas_datatype_bf16_r);
 
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_gemmt_batched_imp.hpp b/projects/rocblas/library/src/blas_ex/rocblas_gemmt_batched_imp.hpp
index d721a4b0ab96..16ba465ad31e 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_gemmt_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_gemmt_batched_imp.hpp
@@ -62,6 +62,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_gemmt_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
         rocblas_internal_logger logger;
         auto                    layer_mode     = handle->layer_mode;
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_gemmt_imp.hpp b/projects/rocblas/library/src/blas_ex/rocblas_gemmt_imp.hpp
index 8f1cf1f56dd3..eb1793be7610 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_gemmt_imp.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_gemmt_imp.hpp
@@ -59,6 +59,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_gemmt_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         auto                    layer_mode     = handle->layer_mode;
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_gemmt_strided_batched_imp.hpp b/projects/rocblas/library/src/blas_ex/rocblas_gemmt_strided_batched_imp.hpp
index c84168b44cff..6d57f4ec3927 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_gemmt_strided_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_gemmt_strided_batched_imp.hpp
@@ -65,6 +65,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, rocblas_gemmt_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         auto                    layer_mode     = handle->layer_mode;
diff --git a/projects/rocblas/library/src/handle.cpp b/projects/rocblas/library/src/handle.cpp
index bb354b889a74..28b19686d30f 100644
--- a/projects/rocblas/library/src/handle.cpp
+++ b/projects/rocblas/library/src/handle.cpp
@@ -981,7 +981,9 @@ void _rocblas_handle::init_logging()
         layer_mode = static_cast<rocblas_layer_mode>(strtol(str_layer_mode, 0, 0));
 
         // open log_trace file
-        if(layer_mode & (rocblas_layer_mode_log_trace | rocblas_layer_mode_log_internal))
+        if(layer_mode
+           & (rocblas_layer_mode_log_trace | rocblas_layer_mode_log_internal
+              | rocblas_layer_mode_log_kernel_select))
             log_trace_os = open_log_stream("ROCBLAS_LOG_TRACE_PATH");
 
         // open log_bench file
diff --git a/projects/rocblas/library/src/hipblaslt_host.cpp b/projects/rocblas/library/src/hipblaslt_host.cpp
index 9459d531d6f7..13f28e364203 100644
--- a/projects/rocblas/library/src/hipblaslt_host.cpp
+++ b/projects/rocblas/library/src/hipblaslt_host.cpp
@@ -418,7 +418,8 @@ namespace
 template <typename Ti, typename To, typename Tc>
 rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<Ti, To, Tc>& prob,
                                               rocblas_gemm_algo                            algo,
-                                              int32_t solution_index)
+                                              int32_t  solution_index,
+                                              int32_t* selected_algo_index_out)
 {
     bool solution_query = algo == rocblas_gemm_algo_solution_index
                           && prob.flags & rocblas_gemm_flags_check_solution_index;
@@ -446,6 +447,9 @@ rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<Ti
             return init;
         }
 
+        if(selected_algo_index_out)
+            *selected_algo_index_out = hipblaslt_ext::getIndexFromAlgo(heuristicResult.algo);
+
         auto gsu_malloc = prob.handle->gsu_malloc_by_size(workspace_size);
 
         if(!gsu_malloc)
@@ -493,6 +497,9 @@ rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<Ti
             return init;
         }
 
+        if(selected_algo_index_out)
+            *selected_algo_index_out = hipblaslt_ext::getIndexFromAlgo(heuristicResult.algo);
+
         auto gsu_malloc = prob.handle->gsu_malloc_by_size(workspace_size);
 
         if(!gsu_malloc)
@@ -712,52 +719,64 @@ rocblas_status getAllSolutionsHipBlasLT(const RocblasContractionProblem<Ti, To,
  ******************************************************************************/
 
 // Non-HPA/GEMM types
-template rocblas_status runContractionProblemHipBlasLT(
-    const RocblasContractionProblem<rocblas_half>&, rocblas_gemm_algo algo, int32_t solution_index);
+template rocblas_status
+    runContractionProblemHipBlasLT(const RocblasContractionProblem<rocblas_half>&,
+                                   rocblas_gemm_algo algo,
+                                   int32_t           solution_index,
+                                   int32_t*          selected_algo_index_out);
 
 template rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<float>&,
                                                        rocblas_gemm_algo algo,
-                                                       int32_t           solution_index);
+                                                       int32_t           solution_index,
+                                                       int32_t*          selected_algo_index_out);
 
 template rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<double>&,
                                                        rocblas_gemm_algo algo,
-                                                       int32_t           solution_index);
+                                                       int32_t           solution_index,
+                                                       int32_t*          selected_algo_index_out);
 
 template rocblas_status
     runContractionProblemHipBlasLT(const RocblasContractionProblem<rocblas_float_complex>&,
                                    rocblas_gemm_algo algo,
-                                   int32_t           solution_index);
+                                   int32_t           solution_index,
+                                   int32_t*          selected_algo_index_out);
 
 template rocblas_status
     runContractionProblemHipBlasLT(const RocblasContractionProblem<rocblas_double_complex>&,
                                    rocblas_gemm_algo algo,
-                                   int32_t           solution_index);
+                                   int32_t           solution_index,
+                                   int32_t*          selected_algo_index_out);
 
 // HPA types
 template rocblas_status runContractionProblemHipBlasLT(
     const RocblasContractionProblem<rocblas_half, rocblas_half, float>&,
     rocblas_gemm_algo algo,
-    int32_t           solution_index);
+    int32_t           solution_index,
+    int32_t*          selected_algo_index_out);
 
 template rocblas_status
     runContractionProblemHipBlasLT(const RocblasContractionProblem<rocblas_half, float, float>&,
                                    rocblas_gemm_algo algo,
-                                   int32_t           solution_index);
+                                   int32_t           solution_index,
+                                   int32_t*          selected_algo_index_out);
 
 template rocblas_status runContractionProblemHipBlasLT(
     const RocblasContractionProblem<rocblas_bfloat16, rocblas_bfloat16, float>&,
     rocblas_gemm_algo algo,
-    int32_t           solution_index);
+    int32_t           solution_index,
+    int32_t*          selected_algo_index_out);
 
 template rocblas_status
     runContractionProblemHipBlasLT(const RocblasContractionProblem<rocblas_bfloat16, float, float>&,
                                    rocblas_gemm_algo algo,
-                                   int32_t           solution_index);
+                                   int32_t           solution_index,
+                                   int32_t*          selected_algo_index_out);
 
 template rocblas_status
     runContractionProblemHipBlasLT(const RocblasContractionProblem<int8_t, int32_t, int32_t>&,
                                    rocblas_gemm_algo algo,
-                                   int32_t           solution_index);
+                                   int32_t           solution_index,
+                                   int32_t*          selected_algo_index_out);
 
 // ********** get all solutions explicits ********
 // Non-HPA/GEMM types
diff --git a/projects/rocblas/library/src/include/handle.hpp b/projects/rocblas/library/src/include/handle.hpp
index 499b0d70eb41..96e09b19e66f 100644
--- a/projects/rocblas/library/src/include/handle.hpp
+++ b/projects/rocblas/library/src/include/handle.hpp
@@ -437,6 +437,11 @@ struct _rocblas_handle
     std::unique_ptr<rocblas_internal_ostream> log_bench_os;
     std::unique_ptr<rocblas_internal_ostream> log_profile_os;
     void                                      init_logging();
+
+    // Name of the user-facing C API call currently executing on this handle. Set by an RAII helper
+    // at the top of each *_impl and cleared on scope exit. Used by rocblas_layer_mode_log_kernel_select
+    // (0x10) to annotate internal-GEMM trace lines with their parent API. Null when no API is active.
+    const char* current_api_name = nullptr;
     void                                      init_check_numerics();
 
     // data pointer for rocSOLVER
@@ -922,6 +927,33 @@ struct _rocblas_handle
     };
 };
 
+// RAII helper that stamps the user-facing C API name onto the handle for the duration of a call.
+// Used by rocblas_layer_mode_log_kernel_select (0x10) to attribute each internal-GEMM trace line
+// to its parent API. Construct one line at the top of each *_impl after the handle null-check.
+// Save/restore protects against re-entry: if an internal path calls another rocBLAS API, the
+// outer name is restored on scope exit. Single-threaded per handle, matching the rest of rocBLAS.
+struct rocblas_internal_api_scope
+{
+    rocblas_internal_api_scope(rocblas_handle h, const char* name)
+        : m_handle(h)
+        , m_prev(h ? h->current_api_name : nullptr)
+    {
+        if(m_handle)
+            m_handle->current_api_name = name;
+    }
+    ~rocblas_internal_api_scope()
+    {
+        if(m_handle)
+            m_handle->current_api_name = m_prev;
+    }
+    rocblas_internal_api_scope(const rocblas_internal_api_scope&)            = delete;
+    rocblas_internal_api_scope& operator=(const rocblas_internal_api_scope&) = delete;
+
+private:
+    rocblas_handle m_handle;
+    const char*    m_prev;
+};
+
 // For functions which don't use temporary device memory, and won't be likely
 // to use them in the future, the RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle)
 // macro can be used to return from a rocblas function with a requested size of 0.
diff --git a/projects/rocblas/library/src/include/hipblaslt_host.hpp b/projects/rocblas/library/src/include/hipblaslt_host.hpp
index 7a9e0255a753..05f3d210085b 100644
--- a/projects/rocblas/library/src/include/hipblaslt_host.hpp
+++ b/projects/rocblas/library/src/include/hipblaslt_host.hpp
@@ -45,7 +45,8 @@
 template <typename Ti, typename To, typename Tc>
 rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<Ti, To, Tc>& problem,
                                               rocblas_gemm_algo algo = rocblas_gemm_algo_standard,
-                                              int32_t           solution_index = 0);
+                                              int32_t           solution_index          = 0,
+                                              int32_t*          selected_algo_index_out = nullptr);
 
 template <typename Ti, typename To, typename Tc>
 rocblas_status getAllSolutionsHipBlasLT(const RocblasContractionProblem<Ti, To, Tc>& prob,
diff --git a/projects/rocblas/library/src/include/logging.hpp b/projects/rocblas/library/src/include/logging.hpp
index b5df2254396b..97653d94e5dd 100644
--- a/projects/rocblas/library/src/include/logging.hpp
+++ b/projects/rocblas/library/src/include/logging.hpp
@@ -250,6 +250,18 @@ class rocblas_internal_logger
         log_arguments(*handle->log_bench_os, " ", std::forward<Ts>(xs)..., atomics_str);
     }
 
+    // if kernel-select logging is turned on with
+    // (handle->layer_mode & rocblas_layer_mode_log_kernel_select) != 0
+    // log_kernel_select emits one space-separated line per internal GEMM sub-problem to the
+    // trace stream. The line is formatted like a rocblas-bench invocation so it can be copied
+    // and replayed directly; trailing kernel-selection metadata is conventionally placed in a
+    // `# source=... kernel=... fallback_from=... parent_api=...` comment by the caller.
+    template <typename... Ts>
+    void log_kernel_select(rocblas_handle handle, Ts&&... xs)
+    {
+        log_arguments(*handle->log_trace_os, " ", std::forward<Ts>(xs)...);
+    }
+
     // if profile logging is turned on with
     // (handle->layer_mode & rocblas_layer_mode_log_profile) != 0
     // log_profile will call rocblas_internal_argument_profile to profile actual arguments,
diff --git a/projects/rocblas/library/src/tensile_host.cpp b/projects/rocblas/library/src/tensile_host.cpp
index d36f686cfeff..c2d1e140d445 100644
--- a/projects/rocblas/library/src/tensile_host.cpp
+++ b/projects/rocblas/library/src/tensile_host.cpp
@@ -33,6 +33,7 @@
 #include "hipblaslt_host.hpp"
 #endif
 
+#include "int64_helpers.hpp"
 #include "logging.hpp"
 #include "tensile_host.hpp"
 
@@ -709,7 +710,9 @@ namespace
                 rocblas_layer_mode layer_mode
                     = static_cast<rocblas_layer_mode>(strtol(str_layer_mode, 0, 0));
 
-                if(layer_mode & (rocblas_layer_mode_log_trace | rocblas_layer_mode_log_internal))
+                if(layer_mode
+                   & (rocblas_layer_mode_log_trace | rocblas_layer_mode_log_internal
+                      | rocblas_layer_mode_log_kernel_select))
                 {
                     char buf[256];
                     rocblas_get_version_string(buf, 256);
@@ -1234,6 +1237,15 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
     rocblas_status status            = rocblas_status_internal_error;
     bool           hipblaslt_backend = false;
 
+    // Per-sub-problem kernel-selection state, consumed by rocblas_layer_mode_log_kernel_select (0x10).
+    // Capture is gated by the layer bit to avoid the std::string copy of solution->name() when off.
+    const bool kernel_select_logging
+        = prob.handle->layer_mode & rocblas_layer_mode_log_kernel_select;
+    bool        hipblaslt_attempted      = false;
+    int32_t     selected_hipblaslt_index = -1;
+    std::string selected_tensile_name;
+    int64_t     selected_tensile_index = -1;
+
     bool solutionBased = algo == rocblas_gemm_algo_solution_index;
     bool hipBLASLtOnly = false, TensileOnly = false;
     bool useDefaultSolution = rocblas_default_solution_index(solution_index);
@@ -1249,12 +1261,17 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
 #ifdef BUILD_WITH_HIPBLASLT
         if(useHipBLASLt(prob))
         {
+            hipblaslt_attempted = true;
             try
             {
                 rocblas_int hipblaslt_idx = map_index_rocblas_to_hipblaslt(solution_index);
 
                 rocblas_internal_ostream msg;
-                auto hipblasltResult = runContractionProblemHipBlasLT(prob, algo, hipblaslt_idx);
+                auto                     hipblasltResult = runContractionProblemHipBlasLT(
+                    prob,
+                    algo,
+                    hipblaslt_idx,
+                    kernel_select_logging ? &selected_hipblaslt_index : nullptr);
 
                 if(hipblasltResult == rocblas_status_success || (solutionBased && hipBLASLtOnly))
                 {
@@ -1317,6 +1334,12 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
             if(!solution && fallbackTensileProblem(tensile_prob))
                 solution = library->findBestSolution(tensile_prob, *hardware, fitness_query);
 
+            if(solution && kernel_select_logging)
+            {
+                selected_tensile_name  = solution->name();
+                selected_tensile_index = solution->index;
+            }
+
             if(!solution)
             {
                 if(!useDefaultSolution)
@@ -1428,6 +1451,93 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
                          prob.col_stride_d);
     }
 
+    // ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select):
+    // one rocblas-bench-replayable line per internal GEMM sub-problem, with kernel-selection
+    // metadata in a trailing `#` comment so the line itself remains copy-pasteable.
+    if(kernel_select_logging && status == rocblas_status_success)
+    {
+        std::string alphas, betas;
+        (void)rocblas_internal_log_bench_alpha_beta_ex(
+            rocblas_datatype_from_type<Tc>, prob.alpha, prob.beta, alphas, betas);
+
+        const char* source   = hipblaslt_backend ? "hipblaslt" : "tensile";
+        const char* fallback = (hipblaslt_attempted && !hipblaslt_backend) ? "hipblaslt" : "none";
+        const char* parent_api
+            = prob.handle->current_api_name ? prob.handle->current_api_name : "unknown";
+
+        rocblas_internal_ostream kernel_field;
+        if(hipblaslt_backend)
+            kernel_field << "kernel=hipblaslt_algo_index=" << selected_hipblaslt_index;
+        else if(!selected_tensile_name.empty())
+            kernel_field << "kernel=" << selected_tensile_name
+                         << " tensile_index=" << selected_tensile_index;
+        else
+            kernel_field << "kernel=unknown";
+
+        rocblas_internal_ostream source_field;
+        source_field << "# source=" << source;
+
+        rocblas_internal_ostream fallback_field;
+        fallback_field << "fallback_from=" << fallback;
+
+        rocblas_internal_ostream parent_field;
+        parent_field << "parent_api=" << parent_api;
+
+        rocblas_internal_logger logger;
+        logger.log_kernel_select(prob.handle,
+                                 ROCBLAS_API_BENCH " -f gemm_strided_batched_ex",
+                                 "--transposeA",
+                                 rocblas_transpose_letter(prob.trans_a),
+                                 "--transposeB",
+                                 rocblas_transpose_letter(prob.trans_b),
+                                 "-m",
+                                 prob.m,
+                                 "-n",
+                                 prob.n,
+                                 "-k",
+                                 prob.k,
+                                 alphas,
+                                 "--a_type",
+                                 rocblas_datatype_string(rocblas_datatype_from_type<Ti>),
+                                 "--lda",
+                                 prob.col_stride_a,
+                                 "--stride_a",
+                                 prob.batch_stride_a,
+                                 "--b_type",
+                                 rocblas_datatype_string(rocblas_datatype_from_type<Ti>),
+                                 "--ldb",
+                                 prob.col_stride_b,
+                                 "--stride_b",
+                                 prob.batch_stride_b,
+                                 betas,
+                                 "--c_type",
+                                 rocblas_datatype_string(rocblas_datatype_from_type<To>),
+                                 "--ldc",
+                                 prob.col_stride_c,
+                                 "--stride_c",
+                                 prob.batch_stride_c,
+                                 "--d_type",
+                                 rocblas_datatype_string(rocblas_datatype_from_type<To>),
+                                 "--ldd",
+                                 prob.col_stride_d,
+                                 "--stride_d",
+                                 prob.batch_stride_d,
+                                 "--batch_count",
+                                 prob.batch_count,
+                                 "--compute_type",
+                                 rocblas_datatype_string(rocblas_datatype_from_type<Tc>),
+                                 "--algo",
+                                 (int)algo,
+                                 "--solution_index",
+                                 solution_index,
+                                 "--flags",
+                                 (uint32_t)prob.flags,
+                                 source_field.str(),
+                                 kernel_field.str(),
+                                 fallback_field.str(),
+                                 parent_field.str());
+    }
+
     return status;
 }
 

From ade55b04e633a65a355067d7641bfbcbd9c9796f Mon Sep 17 00:00:00 2001
From: Estevan Vedovelli <Estevan.Vedovelli@amd.com>
Date: Mon, 6 Jul 2026 16:57:35 -0400
Subject: [PATCH 02/10] Fix clang format

---
 .../src/blas3/Tensile/gemm_templates.cpp      |  3 +--
 .../rocblas/library/src/include/handle.hpp    | 23 +++++++++++--------
 2 files changed, 14 insertions(+), 12 deletions(-)

diff --git a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
index 3ca66b710977..71c3edf57d63 100644
--- a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
+++ b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
@@ -176,8 +176,7 @@ rocblas_status rocblas_internal_gemm(rocblas_handle    handle,
 #else
         const char* fallback = "none";
 #endif
-        const char* parent_api
-            = handle->current_api_name ? handle->current_api_name : "unknown";
+        const char* parent_api = handle->current_api_name ? handle->current_api_name : "unknown";
 
         rocblas_internal_ostream kernel_field;
         kernel_field << "kernel=rocblas_gemm_source_solution_64";
diff --git a/projects/rocblas/library/src/include/handle.hpp b/projects/rocblas/library/src/include/handle.hpp
index 5204e93164ed..ff264d8c8ebf 100644
--- a/projects/rocblas/library/src/include/handle.hpp
+++ b/projects/rocblas/library/src/include/handle.hpp
@@ -119,7 +119,7 @@ ROCBLAS_INTERNAL_EXPORT_NOINLINE rocblas_status
 
 // cached device properties for handle device
 ROCBLAS_INTERNAL_EXPORT_NOINLINE const hipDeviceProp_t*
-                                       rocblas_internal_get_device_prop(rocblas_handle handle);
+    rocblas_internal_get_device_prop(rocblas_handle handle);
 
 /*******************************************************************************
  * \brief rocblas_handle is a structure holding the rocblas library context.
@@ -231,7 +231,7 @@ struct _rocblas_handle
     _rocblas_handle();
     ~_rocblas_handle();
 
-    _rocblas_handle(const _rocblas_handle&) = delete;
+    _rocblas_handle(const _rocblas_handle&)            = delete;
     _rocblas_handle& operator=(const _rocblas_handle&) = delete;
 
     // Set the HIP default device ID to the handle's device ID, and restore on exit
@@ -446,7 +446,7 @@ struct _rocblas_handle
     // at the top of each *_impl and cleared on scope exit. Used by rocblas_layer_mode_log_kernel_select
     // (0x10) to annotate internal-GEMM trace lines with their parent API. Null when no API is active.
     const char* current_api_name = nullptr;
-    void                                      init_check_numerics();
+    void        init_check_numerics();
 
     // data pointer for rocSOLVER
     std::shared_ptr<void> data_ptr;
@@ -523,9 +523,10 @@ struct _rocblas_handle
     // Sets the optimal size(s) of device memory for a kernel call
     // Maximum size is accumulated in device_memory_query_size
     // Returns rocblas_status_size_increased or rocblas_status_size_unchanged
-    template <typename... Ss,
-              std::enable_if_t<sizeof...(Ss) && conjunction<std::is_convertible<Ss, size_t>...>{},
-                               int> = 0>
+    template <
+        typename... Ss,
+        std::enable_if_t<sizeof...(Ss) && conjunction<std::is_convertible<Ss, size_t>...>{}, int>
+        = 0>
     rocblas_status set_optimal_device_memory_size(Ss... sizes)
     {
         if(!device_memory_size_query)
@@ -905,9 +906,10 @@ struct _rocblas_handle
 
 public:
     // Allocate one or more sizes
-    template <typename... Ss,
-              std::enable_if_t<sizeof...(Ss) && conjunction<std::is_convertible<Ss, size_t>...>{},
-                               int> = 0>
+    template <
+        typename... Ss,
+        std::enable_if_t<sizeof...(Ss) && conjunction<std::is_convertible<Ss, size_t>...>{}, int>
+        = 0>
     auto device_malloc(Ss... sizes)
     {
         return _device_malloc(this, size_t(sizes)...);
@@ -915,7 +917,8 @@ struct _rocblas_handle
 
     template <typename... Ss,
               std::enable_if_t<sizeof...(Ss) && conjunction<std::is_convertible<Ss, size_t>...>{},
-                               int> = 0>
+                               int>
+              = 0>
     auto device_malloc_with_GSU(Ss... sizes) //assume last size is gsu size
     {
         //have to assume it can be called to resize in following iteration (resize if > cur?)

From 7cdc3829c1d58996fd5acb1c2a4e18717cc245b8 Mon Sep 17 00:00:00 2001
From: Estevan Vedovelli <Estevan.Vedovelli@amd.com>
Date: Mon, 6 Jul 2026 17:18:50 -0400
Subject: [PATCH 03/10] Rename unit test file

---
 projects/rocblas/clients/gtest/CMakeLists.txt                   | 2 +-
 .../gtest/{kernel_select_gtest.cpp => kernel_select_test.cpp}   | 0
 2 files changed, 1 insertion(+), 1 deletion(-)
 rename projects/rocblas/clients/gtest/{kernel_select_gtest.cpp => kernel_select_test.cpp} (100%)

diff --git a/projects/rocblas/clients/gtest/CMakeLists.txt b/projects/rocblas/clients/gtest/CMakeLists.txt
index 18acde8ebfe8..928d4b818c02 100644
--- a/projects/rocblas/clients/gtest/CMakeLists.txt
+++ b/projects/rocblas/clients/gtest/CMakeLists.txt
@@ -46,7 +46,7 @@ set(rocblas_no_tensile_test_source
     set_get_atomics_mode_gtest.cpp
     set_get_alpha_beta_stride_gtest.cpp
     logging_mode_gtest.cpp
-    kernel_select_gtest.cpp
+    kernel_select_test.cpp
     ostream_threadsafety_gtest.cpp
     set_get_vector_gtest.cpp
     set_get_matrix_gtest.cpp
diff --git a/projects/rocblas/clients/gtest/kernel_select_gtest.cpp b/projects/rocblas/clients/gtest/kernel_select_test.cpp
similarity index 100%
rename from projects/rocblas/clients/gtest/kernel_select_gtest.cpp
rename to projects/rocblas/clients/gtest/kernel_select_test.cpp

From 799caad7f0db71cdc0076fc1747d08a98946cc47 Mon Sep 17 00:00:00 2001
From: Estevan Vedovelli <Estevan.Vedovelli@amd.com>
Date: Mon, 10 Aug 2026 17:17:33 -0400
Subject: [PATCH 04/10] Remove Kernel idx from logs and reduces impacts of
 inits when logging option is disabled

---
 projects/rocblas/CHANGELOG.md                 |  5 +-
 .../clients/gtest/kernel_select_test.cpp      |  5 +-
 .../docs/how-to/logging-in-rocblas.rst        | 10 ++--
 .../library/include/internal/rocblas-types.h  |  4 +-
 .../rocblas/library/src/hipblaslt_host.cpp    | 47 +++++--------------
 .../rocblas/library/src/include/handle.hpp    |  6 ++-
 .../library/src/include/hipblaslt_host.hpp    |  3 +-
 projects/rocblas/library/src/tensile_host.cpp | 44 +++++------------
 8 files changed, 41 insertions(+), 83 deletions(-)

diff --git a/projects/rocblas/CHANGELOG.md b/projects/rocblas/CHANGELOG.md
index efececb880a4..f88b57fa2f06 100644
--- a/projects/rocblas/CHANGELOG.md
+++ b/projects/rocblas/CHANGELOG.md
@@ -5,6 +5,10 @@ rocBLAS documentation is available at
 
 ## rocBLAS 5.7.0
 
+### Added
+
+* `ROCBLAS_LAYER=0x10` (`rocblas_layer_mode_log_kernel_select`): emits a `rocblas-bench`-replayable line per internal GEMM sub-problem, annotated with the selected kernel name (Tensile backend only) and the source backend (`hipblaslt` or `tensile`). Useful for decomposing higher-level BLAS3 calls (`trsm`, `trmm`, `syrk`, ...) into individually benchmarkable kernels. Lines also carry `fallback_from=` and `parent_api=` for context.
+
 ## rocBLAS 5.6.0
 
 ### Added
@@ -25,7 +29,6 @@ rocBLAS documentation is available at
 
 ### Added
 
-* `ROCBLAS_LAYER=0x10` (`rocblas_layer_mode_log_kernel_select`): emits a `rocblas-bench`-replayable line per internal GEMM sub-problem, annotated with the selected kernel name/index and the source backend (`hipblaslt`, `tensile`, or `source`). Useful for decomposing higher-level BLAS3 calls (`trsm`, `trmm`, `syrk`, ...) into individually benchmarkable kernels without parsing `TENSILE_DB`. Lines also carry `fallback_from=` and `parent_api=` for context.
 * Per-batch `alpha`/`beta` support for Level 2 batched and strided-batched `gemv` via `rocblas_set_batch_alpha_stride` and `rocblas_set_batch_beta_stride` (device pointer mode).
 * Per-batch `alpha` support for Level 2 batched and strided-batched `ger`, `geru`, and `gerc` via `rocblas_set_batch_alpha_stride` (device pointer mode).
 * Per-batch `alpha` (scalar vector) API support for `axpy_batched`, `axpy_strided_batched`, and their `_ex` forms through `rocblas_set_batch_alpha_stride` when `rocblas_handle` is in `rocblas_pointer_mode_device`.
diff --git a/projects/rocblas/clients/gtest/kernel_select_test.cpp b/projects/rocblas/clients/gtest/kernel_select_test.cpp
index b8d01790806b..c09f1899e00e 100644
--- a/projects/rocblas/clients/gtest/kernel_select_test.cpp
+++ b/projects/rocblas/clients/gtest/kernel_select_test.cpp
@@ -158,14 +158,13 @@ TEST(KernelSelect, TrsmDecompositionParentApi)
                                 << "got " << lines << " in " << trace_path;
     EXPECT_NE(trace.find("-f gemm_strided_batched_ex"), std::string::npos);
     EXPECT_NE(trace.find("# source="), std::string::npos);
-    EXPECT_NE(trace.find("kernel="), std::string::npos);
 
     std::remove(trace_path.c_str());
 }
 
 // (2) Forcing ROCBLAS_USE_HIPBLASLT=0 should route every GEMM through Tensile.
-// Covers: source=tensile branch, capture of solution->name() / solution->index, and the
-// Tensile kernel-name format (Cijk_* is Tensile's canonical naming convention).
+// Covers: source=tensile branch, capture of solution->name(), and the Tensile kernel-name
+// format (Cijk_* is Tensile's canonical naming convention).
 TEST(KernelSelect, TensileBackendForced)
 {
 #ifndef BUILD_WITH_TENSILE
diff --git a/projects/rocblas/docs/how-to/logging-in-rocblas.rst b/projects/rocblas/docs/how-to/logging-in-rocblas.rst
index 881a87ff1063..677d485990b4 100644
--- a/projects/rocblas/docs/how-to/logging-in-rocblas.rst
+++ b/projects/rocblas/docs/how-to/logging-in-rocblas.rst
@@ -44,10 +44,10 @@ Not all internal APIs are logged. The log output goes to the same stream as trac
 Kernel-selection logging (``ROCBLAS_LAYER=0x10``) emits one ``rocblas-bench``-replayable line per
 internal GEMM sub-problem dispatched, including those produced by decomposed BLAS3 ops such as
 ``trsm``, ``trmm``, ``syrk``, and friends. Each line ends in a ``#`` comment that names the source
-backend (``hipblaslt``, ``tensile``, or ``source``), the selected kernel identifier, whether the
-call fell back from another backend, and the user-facing API that triggered it. The line itself is
-copy-pasteable into ``rocblas-bench`` to reproduce that single sub-problem in isolation. Output
-goes to the same stream as trace logging. Example:
+backend (``hipblaslt`` or ``tensile``), the selected kernel name (Tensile backend only), whether
+the call fell back from another backend, and the user-facing API that triggered it. The line
+itself is copy-pasteable into ``rocblas-bench`` to reproduce that single sub-problem in isolation.
+Output goes to the same stream as trace logging. Example:
 
 .. code-block:: text
 
@@ -55,7 +55,7 @@ goes to the same stream as trace logging. Example:
      --alpha 1 --a_type f32_r --lda 4096 --stride_a 0 --b_type f32_r --ldb 4096 --stride_b 0 \
      --beta 0 --c_type f32_r --ldc 4096 --stride_c 0 --d_type f32_r --ldd 4096 --stride_d 0 \
      --batch_count 1 --compute_type f32_r --algo 0 --solution_index 0 --flags 0 \
-     # source=tensile kernel=Cijk_Ailk_Bljk_SB_MT128x128x8 tensile_index=42 \
+     # source=tensile kernel=Cijk_Ailk_Bljk_SB_MT128x128x8 \
      fallback_from=hipblaslt parent_api=rocblas_strsm
 
 Bits can be combined: ``ROCBLAS_LAYER=0x19`` enables trace (``0x1``), internal (``0x8``), and
diff --git a/projects/rocblas/library/include/internal/rocblas-types.h b/projects/rocblas/library/include/internal/rocblas-types.h
index f3168d901682..6efcca439ccc 100644
--- a/projects/rocblas/library/include/internal/rocblas-types.h
+++ b/projects/rocblas/library/include/internal/rocblas-types.h
@@ -258,8 +258,8 @@ typedef enum rocblas_layer_mode_
     rocblas_layer_mode_log_internal = 0x8,
     /*! \brief Outputs to the same stream as trace logging a rocblas-bench-replayable line for each
      * internal GEMM sub-problem dispatched (including those produced by decomposed BLAS3 ops such as
-     * trsm, trmm, syrk, ...), annotated with the selected kernel name/index and the source backend
-     * (hipblaslt, tensile, or source). */
+     * trsm, trmm, syrk, ...), annotated with the selected kernel name (Tensile backend only) and the
+     * source backend (hipblaslt or tensile). */
     rocblas_layer_mode_log_kernel_select = 0x10,
 } rocblas_layer_mode;
 
diff --git a/projects/rocblas/library/src/hipblaslt_host.cpp b/projects/rocblas/library/src/hipblaslt_host.cpp
index 5ce3b1df4a98..bb3f4b2183ac 100644
--- a/projects/rocblas/library/src/hipblaslt_host.cpp
+++ b/projects/rocblas/library/src/hipblaslt_host.cpp
@@ -558,8 +558,7 @@ namespace
 template <typename Ti, typename To, typename Tc>
 rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<Ti, To, Tc>& prob,
                                               rocblas_gemm_algo                            algo,
-                                              int32_t  solution_index,
-                                              int32_t* selected_algo_index_out)
+                                              int32_t solution_index)
 {
 #if defined(HIPBLASLT_VERSION_MAJOR) && defined(HIPBLASLT_VERSION_MINOR) \
     && defined(HIPBLASLT_VERSION_PATCH)                                  \
@@ -734,10 +733,6 @@ rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<Ti
         {
             throw rocblas_status_memory_error;
         }
-        if(selected_algo_index_out)
-        {
-            *selected_algo_index_out = hipblaslt_ext::getIndexFromAlgo(heuristicResult.algo);
-        }
         workspace = prob.handle->gsu_workspace;
         hipblaslt_alpha_beta_type<Tc> alpha, beta;
         if(prob.alpha != nullptr)
@@ -917,9 +912,6 @@ rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<Ti
             return init;
         }
 
-        if(selected_algo_index_out)
-            *selected_algo_index_out = hipblaslt_ext::getIndexFromAlgo(heuristicResult.algo);
-
         auto gsu_malloc = prob.handle->gsu_malloc_by_size(workspace_size);
 
         if(!gsu_malloc)
@@ -967,9 +959,6 @@ rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<Ti
             return init;
         }
 
-        if(selected_algo_index_out)
-            *selected_algo_index_out = hipblaslt_ext::getIndexFromAlgo(heuristicResult.algo);
-
         auto gsu_malloc = prob.handle->gsu_malloc_by_size(workspace_size);
 
         if(!gsu_malloc)
@@ -1208,64 +1197,52 @@ rocblas_status getAllSolutionsHipBlasLT(const RocblasContractionProblem<Ti, To,
  ******************************************************************************/
 
 // Non-HPA/GEMM types
-template rocblas_status
-    runContractionProblemHipBlasLT(const RocblasContractionProblem<rocblas_half>&,
-                                   rocblas_gemm_algo algo,
-                                   int32_t           solution_index,
-                                   int32_t*          selected_algo_index_out);
+template rocblas_status runContractionProblemHipBlasLT(
+    const RocblasContractionProblem<rocblas_half>&, rocblas_gemm_algo algo, int32_t solution_index);
 
 template rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<float>&,
                                                        rocblas_gemm_algo algo,
-                                                       int32_t           solution_index,
-                                                       int32_t*          selected_algo_index_out);
+                                                       int32_t           solution_index);
 
 template rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<double>&,
                                                        rocblas_gemm_algo algo,
-                                                       int32_t           solution_index,
-                                                       int32_t*          selected_algo_index_out);
+                                                       int32_t           solution_index);
 
 template rocblas_status
     runContractionProblemHipBlasLT(const RocblasContractionProblem<rocblas_float_complex>&,
                                    rocblas_gemm_algo algo,
-                                   int32_t           solution_index,
-                                   int32_t*          selected_algo_index_out);
+                                   int32_t           solution_index);
 
 template rocblas_status
     runContractionProblemHipBlasLT(const RocblasContractionProblem<rocblas_double_complex>&,
                                    rocblas_gemm_algo algo,
-                                   int32_t           solution_index,
-                                   int32_t*          selected_algo_index_out);
+                                   int32_t           solution_index);
 
 // HPA types
 template rocblas_status runContractionProblemHipBlasLT(
     const RocblasContractionProblem<rocblas_half, rocblas_half, float>&,
     rocblas_gemm_algo algo,
-    int32_t           solution_index,
-    int32_t*          selected_algo_index_out);
+    int32_t           solution_index);
 
 template rocblas_status
     runContractionProblemHipBlasLT(const RocblasContractionProblem<rocblas_half, float, float>&,
                                    rocblas_gemm_algo algo,
-                                   int32_t           solution_index,
-                                   int32_t*          selected_algo_index_out);
+                                   int32_t           solution_index);
 
 template rocblas_status runContractionProblemHipBlasLT(
     const RocblasContractionProblem<rocblas_bfloat16, rocblas_bfloat16, float>&,
     rocblas_gemm_algo algo,
-    int32_t           solution_index,
-    int32_t*          selected_algo_index_out);
+    int32_t           solution_index);
 
 template rocblas_status
     runContractionProblemHipBlasLT(const RocblasContractionProblem<rocblas_bfloat16, float, float>&,
                                    rocblas_gemm_algo algo,
-                                   int32_t           solution_index,
-                                   int32_t*          selected_algo_index_out);
+                                   int32_t           solution_index);
 
 template rocblas_status
     runContractionProblemHipBlasLT(const RocblasContractionProblem<int8_t, int32_t, int32_t>&,
                                    rocblas_gemm_algo algo,
-                                   int32_t           solution_index,
-                                   int32_t*          selected_algo_index_out);
+                                   int32_t           solution_index);
 
 // ********** get all solutions explicits ********
 // Non-HPA/GEMM types
diff --git a/projects/rocblas/library/src/include/handle.hpp b/projects/rocblas/library/src/include/handle.hpp
index ff264d8c8ebf..6af8dd85aaa7 100644
--- a/projects/rocblas/library/src/include/handle.hpp
+++ b/projects/rocblas/library/src/include/handle.hpp
@@ -957,13 +957,15 @@ struct _rocblas_handle
 // RAII helper that stamps the user-facing C API name onto the handle for the duration of a call.
 // Used by rocblas_layer_mode_log_kernel_select (0x10) to attribute each internal-GEMM trace line
 // to its parent API. Construct one line at the top of each *_impl after the handle null-check.
+// When kernel-select logging is off (the common case) this is a no-op: the constructor only tests
+// the layer bit and touches the handle when the bit is set, so it costs nothing on the hot path.
 // Save/restore protects against re-entry: if an internal path calls another rocBLAS API, the
 // outer name is restored on scope exit. Single-threaded per handle, matching the rest of rocBLAS.
 struct rocblas_internal_api_scope
 {
     rocblas_internal_api_scope(rocblas_handle h, const char* name)
-        : m_handle(h)
-        , m_prev(h ? h->current_api_name : nullptr)
+        : m_handle((h && (h->layer_mode & rocblas_layer_mode_log_kernel_select)) ? h : nullptr)
+        , m_prev(m_handle ? m_handle->current_api_name : nullptr)
     {
         if(m_handle)
             m_handle->current_api_name = name;
diff --git a/projects/rocblas/library/src/include/hipblaslt_host.hpp b/projects/rocblas/library/src/include/hipblaslt_host.hpp
index 05f3d210085b..7a9e0255a753 100644
--- a/projects/rocblas/library/src/include/hipblaslt_host.hpp
+++ b/projects/rocblas/library/src/include/hipblaslt_host.hpp
@@ -45,8 +45,7 @@
 template <typename Ti, typename To, typename Tc>
 rocblas_status runContractionProblemHipBlasLT(const RocblasContractionProblem<Ti, To, Tc>& problem,
                                               rocblas_gemm_algo algo = rocblas_gemm_algo_standard,
-                                              int32_t           solution_index          = 0,
-                                              int32_t*          selected_algo_index_out = nullptr);
+                                              int32_t           solution_index = 0);
 
 template <typename Ti, typename To, typename Tc>
 rocblas_status getAllSolutionsHipBlasLT(const RocblasContractionProblem<Ti, To, Tc>& prob,
diff --git a/projects/rocblas/library/src/tensile_host.cpp b/projects/rocblas/library/src/tensile_host.cpp
index ef81d25b37f7..3e3c2194eda4 100644
--- a/projects/rocblas/library/src/tensile_host.cpp
+++ b/projects/rocblas/library/src/tensile_host.cpp
@@ -1270,10 +1270,8 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
     // Capture is gated by the layer bit to avoid the std::string copy of solution->name() when off.
     const bool kernel_select_logging
         = prob.handle->layer_mode & rocblas_layer_mode_log_kernel_select;
-    bool        hipblaslt_attempted      = false;
-    int32_t     selected_hipblaslt_index = -1;
+    bool        hipblaslt_attempted = false;
     std::string selected_tensile_name;
-    int64_t     selected_tensile_index = -1;
 
     bool solutionBased = algo == rocblas_gemm_algo_solution_index;
     bool hipBLASLtOnly = false, TensileOnly = false;
@@ -1296,11 +1294,7 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
                 rocblas_int hipblaslt_idx = map_index_rocblas_to_hipblaslt(solution_index);
 
                 rocblas_internal_ostream msg;
-                auto                     hipblasltResult = runContractionProblemHipBlasLT(
-                    prob,
-                    algo,
-                    hipblaslt_idx,
-                    kernel_select_logging ? &selected_hipblaslt_index : nullptr);
+                auto hipblasltResult = runContractionProblemHipBlasLT(prob, algo, hipblaslt_idx);
 
                 if(hipblasltResult == rocblas_status_success || (solutionBased && hipBLASLtOnly))
                 {
@@ -1366,10 +1360,7 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
                 solution = library->findBestSolution(tensile_prob, *hardware, fitness_query);
 
             if(solution && kernel_select_logging)
-            {
-                selected_tensile_name  = solution->name();
-                selected_tensile_index = solution->index;
-            }
+                selected_tensile_name = solution->name();
 
             if(!solution)
             {
@@ -1500,23 +1491,13 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
         const char* parent_api
             = prob.handle->current_api_name ? prob.handle->current_api_name : "unknown";
 
-        rocblas_internal_ostream kernel_field;
-        if(hipblaslt_backend)
-            kernel_field << "kernel=hipblaslt_algo_index=" << selected_hipblaslt_index;
-        else if(!selected_tensile_name.empty())
-            kernel_field << "kernel=" << selected_tensile_name
-                         << " tensile_index=" << selected_tensile_index;
-        else
-            kernel_field << "kernel=unknown";
-
-        rocblas_internal_ostream source_field;
-        source_field << "# source=" << source;
-
-        rocblas_internal_ostream fallback_field;
-        fallback_field << "fallback_from=" << fallback;
-
-        rocblas_internal_ostream parent_field;
-        parent_field << "parent_api=" << parent_api;
+        // Trailing kernel-selection metadata, kept in one `#` comment so the bench line stays
+        // copy-pasteable. For tensile we keep the kernel name only.
+        rocblas_internal_ostream metadata_field;
+        metadata_field << "# source=" << source;
+        if(!hipblaslt_backend && !selected_tensile_name.empty())
+            metadata_field << " kernel=" << selected_tensile_name;
+        metadata_field << " fallback_from=" << fallback << " parent_api=" << parent_api;
 
         rocblas_internal_logger logger;
         logger.log_kernel_select(prob.handle,
@@ -1567,10 +1548,7 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
                                  solution_index,
                                  "--flags",
                                  (uint32_t)prob.flags,
-                                 source_field.str(),
-                                 kernel_field.str(),
-                                 fallback_field.str(),
-                                 parent_field.str());
+                                 metadata_field.str());
     }
 
     return status;

From 027edb8ac5743afe4b277b482cf0e715ef74bb08 Mon Sep 17 00:00:00 2001
From: Estevan Vedovelli <Estevan.Vedovelli@amd.com>
Date: Tue, 11 Aug 2026 12:16:52 -0400
Subject: [PATCH 05/10] Fix format and test verification

---
 .../clients/include/testing_logging.hpp       |  3 ++-
 .../rocblas/library/src/include/handle.hpp    | 23 ++++++++-----------
 2 files changed, 12 insertions(+), 14 deletions(-)

diff --git a/projects/rocblas/clients/include/testing_logging.hpp b/projects/rocblas/clients/include/testing_logging.hpp
index 6a97afc7ed29..051a9561a9e2 100644
--- a/projects/rocblas/clients/include/testing_logging.hpp
+++ b/projects/rocblas/clients/include/testing_logging.hpp
@@ -654,6 +654,8 @@ void testing_logging(const Arguments& arg)
     // Smoke check for ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select). We don't
     // golden-file the line because the selected kernel name varies by GPU arch and backend
     // build — just assert at least one bench-style record with the expected schema was emitted.
+    // Only source=/parent_api= are always present; kernel= is emitted for the Tensile backend
+    // only, so it is intentionally not required here (hipBLASLt-routed GEMMs omit it).
     if constexpr(!rocblas_is_complex<T>)
     {
         const fs::path kernel_select_fspath
@@ -679,7 +681,6 @@ void testing_logging(const Arguments& arg)
         {
             if(line.find("-f gemm_strided_batched_ex") != std::string::npos
                && line.find("# source=") != std::string::npos
-               && line.find("kernel=") != std::string::npos
                && line.find("parent_api=") != std::string::npos)
             {
                 found = true;
diff --git a/projects/rocblas/library/src/include/handle.hpp b/projects/rocblas/library/src/include/handle.hpp
index 6af8dd85aaa7..6727c3aae836 100644
--- a/projects/rocblas/library/src/include/handle.hpp
+++ b/projects/rocblas/library/src/include/handle.hpp
@@ -119,7 +119,7 @@ ROCBLAS_INTERNAL_EXPORT_NOINLINE rocblas_status
 
 // cached device properties for handle device
 ROCBLAS_INTERNAL_EXPORT_NOINLINE const hipDeviceProp_t*
-    rocblas_internal_get_device_prop(rocblas_handle handle);
+                                       rocblas_internal_get_device_prop(rocblas_handle handle);
 
 /*******************************************************************************
  * \brief rocblas_handle is a structure holding the rocblas library context.
@@ -231,7 +231,7 @@ struct _rocblas_handle
     _rocblas_handle();
     ~_rocblas_handle();
 
-    _rocblas_handle(const _rocblas_handle&)            = delete;
+    _rocblas_handle(const _rocblas_handle&) = delete;
     _rocblas_handle& operator=(const _rocblas_handle&) = delete;
 
     // Set the HIP default device ID to the handle's device ID, and restore on exit
@@ -523,10 +523,9 @@ struct _rocblas_handle
     // Sets the optimal size(s) of device memory for a kernel call
     // Maximum size is accumulated in device_memory_query_size
     // Returns rocblas_status_size_increased or rocblas_status_size_unchanged
-    template <
-        typename... Ss,
-        std::enable_if_t<sizeof...(Ss) && conjunction<std::is_convertible<Ss, size_t>...>{}, int>
-        = 0>
+    template <typename... Ss,
+              std::enable_if_t<sizeof...(Ss) && conjunction<std::is_convertible<Ss, size_t>...>{},
+                               int> = 0>
     rocblas_status set_optimal_device_memory_size(Ss... sizes)
     {
         if(!device_memory_size_query)
@@ -906,10 +905,9 @@ struct _rocblas_handle
 
 public:
     // Allocate one or more sizes
-    template <
-        typename... Ss,
-        std::enable_if_t<sizeof...(Ss) && conjunction<std::is_convertible<Ss, size_t>...>{}, int>
-        = 0>
+    template <typename... Ss,
+              std::enable_if_t<sizeof...(Ss) && conjunction<std::is_convertible<Ss, size_t>...>{},
+                               int> = 0>
     auto device_malloc(Ss... sizes)
     {
         return _device_malloc(this, size_t(sizes)...);
@@ -917,8 +915,7 @@ struct _rocblas_handle
 
     template <typename... Ss,
               std::enable_if_t<sizeof...(Ss) && conjunction<std::is_convertible<Ss, size_t>...>{},
-                               int>
-              = 0>
+                               int> = 0>
     auto device_malloc_with_GSU(Ss... sizes) //assume last size is gsu size
     {
         //have to assume it can be called to resize in following iteration (resize if > cur?)
@@ -975,7 +972,7 @@ struct rocblas_internal_api_scope
         if(m_handle)
             m_handle->current_api_name = m_prev;
     }
-    rocblas_internal_api_scope(const rocblas_internal_api_scope&)            = delete;
+    rocblas_internal_api_scope(const rocblas_internal_api_scope&) = delete;
     rocblas_internal_api_scope& operator=(const rocblas_internal_api_scope&) = delete;
 
 private:

From c8ad9592e6d5b34e4faf8479eebd53a80482ba7f Mon Sep 17 00:00:00 2001
From: Estevan Vedovelli <Estevan.Vedovelli@amd.com>
Date: Tue, 18 Aug 2026 15:47:39 -0400
Subject: [PATCH 06/10] Address review comments

---
 .../clients/gtest/kernel_select_test.cpp      |  38 ++++--
 .../src/blas3/Tensile/gemm_templates.cpp      | 129 +++++++-----------
 .../rocblas/library/src/include/logging.hpp   |  96 +++++++++++++
 projects/rocblas/library/src/tensile_host.cpp |  80 +++++------
 4 files changed, 209 insertions(+), 134 deletions(-)

diff --git a/projects/rocblas/clients/gtest/kernel_select_test.cpp b/projects/rocblas/clients/gtest/kernel_select_test.cpp
index c09f1899e00e..489d6ecd5ddb 100644
--- a/projects/rocblas/clients/gtest/kernel_select_test.cpp
+++ b/projects/rocblas/clients/gtest/kernel_select_test.cpp
@@ -27,6 +27,7 @@
 
 #include "client_utility.hpp"
 #include "device_vector.hpp"
+#include "host_vector.hpp"
 #include "rocblas_test.hpp"
 #include <cstdio>
 #include <cstdlib>
@@ -36,8 +37,16 @@
 #include <rocblas/rocblas.h>
 #include <sstream>
 #include <string>
-#ifndef WIN32
-#include <unistd.h>
+#include <vector>
+
+#if __has_include(<filesystem>)
+#include <filesystem>
+namespace fs = std::filesystem;
+#elif __has_include(<experimental/filesystem>)
+#include <experimental/filesystem>
+namespace fs = std::experimental::filesystem;
+#else
+#error no filesystem found
 #endif
 
 #ifdef WIN32
@@ -68,14 +77,14 @@ namespace
         return n;
     }
 
-    // Generates a unique trace-file path under /tmp for this test invocation.
+    // Generates a unique, temp-dir-rooted trace-file path for this test invocation
+    // using rocblas_tempname().
     std::string make_trace_path(const char* test_name)
     {
-        std::string p = "/tmp/rocblas_kernel_select_";
-        p += test_name;
-        p += "_";
-        p += std::to_string(getpid());
-        p += ".log";
+        static const std::string tmp_dir = rocblas_tempname();
+        const fs::path           fspath
+            = tmp_dir + std::string("kernel_select_") + test_name + std::string(".log");
+        const std::string p = fspath.generic_string();
         std::remove(p.c_str());
         return p;
     }
@@ -130,8 +139,18 @@ TEST(KernelSelect, TrsmDecompositionParentApi)
         rocblas_handle handle;
         ASSERT_EQ(rocblas_create_handle(&handle), rocblas_status_success);
 
+        host_vector<float> hA(size_t(m) * m);
+        host_vector<float> hB(size_t(m) * n);
+        for(int col = 0; col < m; ++col)
+            for(int row = 0; row < m; ++row)
+                hA[size_t(col) * m + row] = row == col ? float(m) : (row < col ? 0.5f : 0.0f);
+        for(size_t i = 0; i < size_t(m) * n; ++i)
+            hB[i] = 1.0f;
+
         device_vector<float> dA(size_t(m) * m);
         device_vector<float> dB(size_t(m) * n);
+        ASSERT_EQ(dA.transfer_from(hA), hipSuccess);
+        ASSERT_EQ(dB.transfer_from(hB), hipSuccess);
 
         ASSERT_EQ(rocblas_strsm(handle,
                                 rocblas_side_left,
@@ -156,7 +175,8 @@ TEST(KernelSelect, TrsmDecompositionParentApi)
     const size_t lines = count_occurrences(trace, "parent_api=rocblas_strsm");
     EXPECT_GT(lines, size_t(1)) << "expected multiple sub-problems attributed to rocblas_strsm, "
                                 << "got " << lines << " in " << trace_path;
-    EXPECT_NE(trace.find("-f gemm_strided_batched_ex"), std::string::npos);
+    // Non-batched sub-problems (batch_count == 1) replay as gemm_ex.
+    EXPECT_NE(trace.find("-f gemm_ex"), std::string::npos);
     EXPECT_NE(trace.find("# source="), std::string::npos);
 
     std::remove(trace_path.c_str());
diff --git a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
index 4f465be199b1..c15580eae872 100644
--- a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
+++ b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
@@ -168,82 +168,6 @@ rocblas_status rocblas_internal_gemm(rocblas_handle    handle,
                          ldc);
     }
 
-    // ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select): bench-style line for the
-    // source-GEMM fallback path. Reached when Tensile returned not_implemented (or BUILD_WITH_TENSILE
-    // is off), so fallback_from is "tensile" when Tensile was built in, else "none".
-    if(handle->layer_mode & rocblas_layer_mode_log_kernel_select)
-    {
-#ifdef BUILD_WITH_TENSILE
-        const char* fallback = "tensile";
-#else
-        const char* fallback = "none";
-#endif
-        const char* parent_api = handle->current_api_name ? handle->current_api_name : "unknown";
-
-        rocblas_internal_ostream kernel_field;
-        kernel_field << "kernel=rocblas_gemm_source_solution_64";
-        rocblas_internal_ostream source_field;
-        source_field << "# source=source";
-        rocblas_internal_ostream fallback_field;
-        fallback_field << "fallback_from=" << fallback;
-        rocblas_internal_ostream parent_field;
-        parent_field << "parent_api=" << parent_api;
-
-        rocblas_internal_logger logger;
-        logger.log_kernel_select(handle,
-                                 ROCBLAS_API_BENCH " -f gemm_strided_batched_ex",
-                                 "--transposeA",
-                                 rocblas_transpose_letter(trans_a),
-                                 "--transposeB",
-                                 rocblas_transpose_letter(trans_b),
-                                 "-m",
-                                 m,
-                                 "-n",
-                                 n,
-                                 "-k",
-                                 k,
-                                 LOG_BENCH_SCALAR_VALUE(handle, alpha),
-                                 "--a_type",
-                                 rocblas_precision_string<TScal>,
-                                 "--lda",
-                                 lda,
-                                 "--stride_a",
-                                 stride_a,
-                                 "--b_type",
-                                 rocblas_precision_string<TScal>,
-                                 "--ldb",
-                                 ldb,
-                                 "--stride_b",
-                                 stride_b,
-                                 LOG_BENCH_SCALAR_VALUE(handle, beta),
-                                 "--c_type",
-                                 rocblas_precision_string<TScal>,
-                                 "--ldc",
-                                 ldc,
-                                 "--stride_c",
-                                 stride_c,
-                                 "--d_type",
-                                 rocblas_precision_string<TScal>,
-                                 "--ldd",
-                                 ldc,
-                                 "--stride_d",
-                                 stride_c,
-                                 "--batch_count",
-                                 batch_count,
-                                 "--compute_type",
-                                 rocblas_precision_string<TScal>,
-                                 "--algo",
-                                 0,
-                                 "--solution_index",
-                                 0,
-                                 "--flags",
-                                 0,
-                                 source_field.str(),
-                                 kernel_field.str(),
-                                 fallback_field.str(),
-                                 parent_field.str());
-    }
-
     if(k == 0 || (alpha && *alpha == 0))
     {
         return rocblas_gemm_scale_launcher_64(
@@ -287,6 +211,59 @@ rocblas_status rocblas_internal_gemm(rocblas_handle    handle,
             return status;
     }
 
+    // ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select): bench-style line for the
+    // source-GEMM fallback path. Reached when Tensile returned not_implemented (or
+    // BUILD_WITH_TENSILE is off), so fallback_from is "tensile" when Tensile was built in, else "none".
+    if(status == rocblas_status_success
+       && (handle->layer_mode & rocblas_layer_mode_log_kernel_select))
+    {
+#ifdef BUILD_WITH_TENSILE
+        const char* fallback = "tensile";
+#else
+        const char* fallback = "none";
+#endif
+        const char* parent_api = handle->current_api_name ? handle->current_api_name : "unknown";
+
+        rocblas_internal_ostream metadata_field;
+        metadata_field << "# source=source kernel=rocblas_gemm_source_solution_64 fallback_from="
+                       << fallback << " parent_api=" << parent_api;
+        const std::string metadata = metadata_field.str();
+        const std::string alphas   = LOG_BENCH_SCALAR_VALUE(handle, alpha);
+        const std::string betas    = LOG_BENCH_SCALAR_VALUE(handle, beta);
+
+        rocblas_kernel_select_bench_line line{};
+        line.bench_prefix   = ROCBLAS_API_BENCH;
+        line.function       = batch_count > 1 ? "gemm_strided_batched_ex" : "gemm_ex";
+        line.trans_a        = rocblas_transpose_letter(trans_a);
+        line.trans_b        = rocblas_transpose_letter(trans_b);
+        line.m              = m;
+        line.n              = n;
+        line.k              = k;
+        line.alpha          = alphas.c_str();
+        line.a_type         = rocblas_precision_string<TScal>;
+        line.lda            = lda;
+        line.stride_a       = stride_a;
+        line.b_type         = rocblas_precision_string<TScal>;
+        line.ldb            = ldb;
+        line.stride_b       = stride_b;
+        line.beta           = betas.c_str();
+        line.c_type         = rocblas_precision_string<TScal>;
+        line.ldc            = ldc;
+        line.stride_c       = stride_c;
+        line.d_type         = rocblas_precision_string<TScal>;
+        line.ldd            = ldc;
+        line.stride_d       = stride_c;
+        line.batch_count    = batch_count;
+        line.compute_type   = rocblas_precision_string<TScal>;
+        line.algo           = 0;
+        line.solution_index = 0;
+        line.flags          = 0;
+        line.metadata       = metadata.c_str();
+
+        rocblas_internal_logger logger;
+        logger.log_kernel_select_bench(handle, line);
+    }
+
     return status;
 }
 
diff --git a/projects/rocblas/library/src/include/logging.hpp b/projects/rocblas/library/src/include/logging.hpp
index 97653d94e5dd..51429983551f 100644
--- a/projects/rocblas/library/src/include/logging.hpp
+++ b/projects/rocblas/library/src/include/logging.hpp
@@ -206,6 +206,42 @@ double rocblas_internal_value_category(const T* beta, rocblas_datatype compute_t
 
 extern const char* c_rocblas_internal;
 
+// Fields for one ROCBLAS_LAYER=0x10 kernel-select bench line. Populated by the internal-GEMM
+// paths (tensile_host.cpp, gemm_templates.cpp) and rendered by
+// rocblas_internal_logger::log_kernel_select_bench, which owns the replay-order argument list.
+// Strings are non-owning views; the caller must keep the backing storage alive across the call.
+struct rocblas_kernel_select_bench_line
+{
+    const char* bench_prefix; // e.g. ROCBLAS_API_BENCH
+    // "gemm_ex" for a single sub-problem, "gemm_strided_batched_ex" when batched/strided.
+    const char* function;
+    char        trans_a;
+    char        trans_b;
+    int64_t     m;
+    int64_t     n;
+    int64_t     k;
+    const char* alpha; // preformatted "--alpha <v>"
+    const char* a_type;
+    int64_t     lda;
+    rocblas_stride stride_a;
+    const char* b_type;
+    int64_t     ldb;
+    rocblas_stride stride_b;
+    const char* beta; // preformatted "--beta <v>"
+    const char* c_type;
+    int64_t     ldc;
+    rocblas_stride stride_c;
+    const char* d_type;
+    int64_t     ldd;
+    rocblas_stride stride_d;
+    int64_t     batch_count;
+    const char* compute_type;
+    int32_t     algo;
+    int32_t     solution_index;
+    uint32_t    flags;
+    const char* metadata; // trailing "# source=... kernel=... fallback_from=... parent_api=..."
+};
+
 /******************************************************************
  * ROCBLAS LOGGER *
  ******************************************************************/
@@ -262,6 +298,66 @@ class rocblas_internal_logger
         log_arguments(*handle->log_trace_os, " ", std::forward<Ts>(xs)...);
     }
 
+    // Emits a single rocblas-bench-replayable kernel-select line from a filled-in
+    // rocblas_kernel_select_bench_line. This centralizes the long, replay-order-sensitive
+    // argument list so the Tensile (tensile_host.cpp) and source-GEMM fallback
+    // (gemm_templates.cpp) paths cannot drift out of sync.
+    void log_kernel_select_bench(rocblas_handle handle, const rocblas_kernel_select_bench_line& b)
+    {
+        log_arguments(*handle->log_trace_os,
+                      " ",
+                      b.bench_prefix,
+                      "-f",
+                      b.function,
+                      "--transposeA",
+                      b.trans_a,
+                      "--transposeB",
+                      b.trans_b,
+                      "-m",
+                      b.m,
+                      "-n",
+                      b.n,
+                      "-k",
+                      b.k,
+                      b.alpha,
+                      "--a_type",
+                      b.a_type,
+                      "--lda",
+                      b.lda,
+                      "--stride_a",
+                      b.stride_a,
+                      "--b_type",
+                      b.b_type,
+                      "--ldb",
+                      b.ldb,
+                      "--stride_b",
+                      b.stride_b,
+                      b.beta,
+                      "--c_type",
+                      b.c_type,
+                      "--ldc",
+                      b.ldc,
+                      "--stride_c",
+                      b.stride_c,
+                      "--d_type",
+                      b.d_type,
+                      "--ldd",
+                      b.ldd,
+                      "--stride_d",
+                      b.stride_d,
+                      "--batch_count",
+                      b.batch_count,
+                      "--compute_type",
+                      b.compute_type,
+                      "--algo",
+                      b.algo,
+                      "--solution_index",
+                      b.solution_index,
+                      "--flags",
+                      b.flags,
+                      b.metadata);
+    }
+
     // if profile logging is turned on with
     // (handle->layer_mode & rocblas_layer_mode_log_profile) != 0
     // log_profile will call rocblas_internal_argument_profile to profile actual arguments,
diff --git a/projects/rocblas/library/src/tensile_host.cpp b/projects/rocblas/library/src/tensile_host.cpp
index c141cd3e068d..819f99836abe 100644
--- a/projects/rocblas/library/src/tensile_host.cpp
+++ b/projects/rocblas/library/src/tensile_host.cpp
@@ -1507,57 +1507,39 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
         if(!hipblaslt_backend && !selected_tensile_name.empty())
             metadata_field << " kernel=" << selected_tensile_name;
         metadata_field << " fallback_from=" << fallback << " parent_api=" << parent_api;
+        const std::string metadata = metadata_field.str();
+
+        rocblas_kernel_select_bench_line line{};
+        line.bench_prefix   = ROCBLAS_API_BENCH;
+        line.function       = prob.batch_count > 1 ? "gemm_strided_batched_ex" : "gemm_ex";
+        line.trans_a        = rocblas_transpose_letter(prob.trans_a);
+        line.trans_b        = rocblas_transpose_letter(prob.trans_b);
+        line.m              = prob.m;
+        line.n              = prob.n;
+        line.k              = prob.k;
+        line.alpha          = alphas.c_str();
+        line.a_type         = rocblas_datatype_string(rocblas_datatype_from_type<Ti>);
+        line.lda            = prob.col_stride_a;
+        line.stride_a       = prob.batch_stride_a;
+        line.b_type         = rocblas_datatype_string(rocblas_datatype_from_type<Ti>);
+        line.ldb            = prob.col_stride_b;
+        line.stride_b       = prob.batch_stride_b;
+        line.beta           = betas.c_str();
+        line.c_type         = rocblas_datatype_string(rocblas_datatype_from_type<To>);
+        line.ldc            = prob.col_stride_c;
+        line.stride_c       = prob.batch_stride_c;
+        line.d_type         = rocblas_datatype_string(rocblas_datatype_from_type<To>);
+        line.ldd            = prob.col_stride_d;
+        line.stride_d       = prob.batch_stride_d;
+        line.batch_count    = prob.batch_count;
+        line.compute_type   = rocblas_datatype_string(rocblas_datatype_from_type<Tc>);
+        line.algo           = (int32_t)algo;
+        line.solution_index = solution_index;
+        line.flags          = (uint32_t)prob.flags;
+        line.metadata       = metadata.c_str();
 
         rocblas_internal_logger logger;
-        logger.log_kernel_select(prob.handle,
-                                 ROCBLAS_API_BENCH " -f gemm_strided_batched_ex",
-                                 "--transposeA",
-                                 rocblas_transpose_letter(prob.trans_a),
-                                 "--transposeB",
-                                 rocblas_transpose_letter(prob.trans_b),
-                                 "-m",
-                                 prob.m,
-                                 "-n",
-                                 prob.n,
-                                 "-k",
-                                 prob.k,
-                                 alphas,
-                                 "--a_type",
-                                 rocblas_datatype_string(rocblas_datatype_from_type<Ti>),
-                                 "--lda",
-                                 prob.col_stride_a,
-                                 "--stride_a",
-                                 prob.batch_stride_a,
-                                 "--b_type",
-                                 rocblas_datatype_string(rocblas_datatype_from_type<Ti>),
-                                 "--ldb",
-                                 prob.col_stride_b,
-                                 "--stride_b",
-                                 prob.batch_stride_b,
-                                 betas,
-                                 "--c_type",
-                                 rocblas_datatype_string(rocblas_datatype_from_type<To>),
-                                 "--ldc",
-                                 prob.col_stride_c,
-                                 "--stride_c",
-                                 prob.batch_stride_c,
-                                 "--d_type",
-                                 rocblas_datatype_string(rocblas_datatype_from_type<To>),
-                                 "--ldd",
-                                 prob.col_stride_d,
-                                 "--stride_d",
-                                 prob.batch_stride_d,
-                                 "--batch_count",
-                                 prob.batch_count,
-                                 "--compute_type",
-                                 rocblas_datatype_string(rocblas_datatype_from_type<Tc>),
-                                 "--algo",
-                                 (int)algo,
-                                 "--solution_index",
-                                 solution_index,
-                                 "--flags",
-                                 (uint32_t)prob.flags,
-                                 metadata_field.str());
+        logger.log_kernel_select_bench(prob.handle, line);
     }
 
     return status;

From 7aca35940ee11913e7ac8112c5d0f023c476b954 Mon Sep 17 00:00:00 2001
From: Estevan Vedovelli <Estevan.Vedovelli@amd.com>
Date: Wed, 28 Oct 2026 14:27:06 -0400
Subject: [PATCH 07/10] Fix failing test and small issue in the kernel select
 logging

---
 projects/rocblas/CHANGELOG.md                 |   2 +-
 projects/rocblas/clients/gtest/CMakeLists.txt |   2 +-
 projects/rocblas/clients/gtest/aux_gtest.yaml |   1 +
 .../clients/gtest/kernel_select_gtest.yaml    |  23 +
 .../clients/gtest/kernel_select_test.cpp      | 360 -------------
 .../clients/include/testing_kernel_select.hpp | 471 ++++++++++++++++++
 .../clients/include/testing_logging.hpp       |   2 +-
 .../docs/how-to/logging-in-rocblas.rst        |  31 +-
 .../src/blas3/Tensile/gemm_templates.cpp      |  88 ++--
 .../src/blas3/Tensile/gemm_templates.hpp      |  44 ++
 .../rocblas_gemm_grouped_batched_imp.hpp      |   3 +
 .../library/src/blas_ex/rocblas_gemm_ex.hpp   | 123 +++--
 .../rocblas_gemm_grouped_batched_ex_imp.hpp   |   3 +
 .../src/blas_ex/rocblas_herk_ex_imp.hpp       |   2 +
 .../src/blas_ex/rocblas_syrk_ex_imp.hpp       |   2 +
 .../rocblas/library/src/include/handle.hpp    |  21 +-
 .../rocblas/library/src/include/logging.hpp   | 142 ++----
 projects/rocblas/library/src/logging.cpp      |  58 +++
 .../src64/blas3/rocblas_gemm_kernels_64.cpp   |  25 +
 .../blas_ex/rocblas_gemm_ex_kernels_64.cpp    |  33 ++
 projects/rocblas/library/src/tensile_host.cpp | 108 ++--
 21 files changed, 912 insertions(+), 632 deletions(-)
 create mode 100644 projects/rocblas/clients/gtest/kernel_select_gtest.yaml
 delete mode 100644 projects/rocblas/clients/gtest/kernel_select_test.cpp
 create mode 100644 projects/rocblas/clients/include/testing_kernel_select.hpp

diff --git a/projects/rocblas/CHANGELOG.md b/projects/rocblas/CHANGELOG.md
index 408f59f9ee1d..88d453f72644 100644
--- a/projects/rocblas/CHANGELOG.md
+++ b/projects/rocblas/CHANGELOG.md
@@ -7,6 +7,7 @@ rocBLAS documentation is available at
 
 ### Added
 
+* `ROCBLAS_LAYER=0x10` (`rocblas_layer_mode_log_kernel_select`) writes one `rocblas-bench`-replayable line to the trace log for each internal GEMM sub-problem that launches a kernel. Use it to break higher-level Level 3 calls such as `trsm`, `trmm`, and `syrk` into kernels that can be benchmarked individually. The bench function is `gemm_ex`, `gemm_batched_ex`, or `gemm_strided_batched_ex`, matching the sub-problem layout. A trailing comment names the path that ran it (`hipblaslt`, `tensile`, `source`, or `gemv`), the selected kernel or launcher (except for hipBLASLt), the backends it fell back from, and the calling API.
 * gfx1170 and gfx1171 enabled. Level 3 `gemm`, `gemm_ex`, and the functions that internally use GEMM default to the hipBLASLt backend on these architectures. `ROCBLAS_USE_HIPBLASLT` continues to force or disable the hipBLASLt backend.
 
 ### Changed
@@ -39,7 +40,6 @@ rocBLAS documentation is available at
 
 ### Added
 
-* `ROCBLAS_LAYER=0x10` (`rocblas_layer_mode_log_kernel_select`): emits a `rocblas-bench`-replayable line per internal GEMM sub-problem, annotated with the selected kernel name (Tensile backend only) and the source backend (`hipblaslt` or `tensile`). Useful for decomposing higher-level BLAS3 calls (`trsm`, `trmm`, `syrk`, ...) into individually benchmarkable kernels. Lines also carry `fallback_from=` and `parent_api=` for context.
 * Level 3 grouped batched GEMM functions `rocblas_sgemm_grouped_batched`, `rocblas_dgemm_grouped_batched`, and `rocblas_gemm_grouped_batched_ex` for both C and FORTRAN, including ILP64 API (`_64` name suffix).
 
 
diff --git a/projects/rocblas/clients/gtest/CMakeLists.txt b/projects/rocblas/clients/gtest/CMakeLists.txt
index b4b862faba9e..61c0d97cae0e 100644
--- a/projects/rocblas/clients/gtest/CMakeLists.txt
+++ b/projects/rocblas/clients/gtest/CMakeLists.txt
@@ -47,7 +47,7 @@ set(rocblas_no_tensile_test_source
     set_get_atomics_mode_gtest.cpp
     set_get_alpha_beta_stride_gtest.cpp
     logging_mode_gtest.cpp
-    kernel_select_test.cpp
+    kernel_select_gtest.cpp
     ostream_threadsafety_gtest.cpp
     set_get_vector_gtest.cpp
     set_get_matrix_gtest.cpp
diff --git a/projects/rocblas/clients/gtest/aux_gtest.yaml b/projects/rocblas/clients/gtest/aux_gtest.yaml
index 47319f7e4d24..a63fcd5b4f32 100644
--- a/projects/rocblas/clients/gtest/aux_gtest.yaml
+++ b/projects/rocblas/clients/gtest/aux_gtest.yaml
@@ -2,6 +2,7 @@
 include: set_get_matrix_gtest.yaml
 include: set_get_vector_gtest.yaml
 include: logging_mode_gtest.yaml
+include: kernel_select_gtest.yaml
 include: set_get_pointer_mode_gtest.yaml
 include: set_get_atomics_mode_gtest.yaml
 include: set_get_alpha_beta_stride_gtest.yaml
diff --git a/projects/rocblas/clients/gtest/kernel_select_gtest.yaml b/projects/rocblas/clients/gtest/kernel_select_gtest.yaml
new file mode 100644
index 000000000000..cc234fea898c
--- /dev/null
+++ b/projects/rocblas/clients/gtest/kernel_select_gtest.yaml
@@ -0,0 +1,23 @@
+---
+include: rocblas_common.yaml
+include: known_bugs.yaml
+
+Tests:
+- name: kernel_select
+  category: quick
+  function:
+    - kernel_select_trsm: *single_precision
+    - kernel_select_layout: *single_precision
+    - kernel_select_combined: *single_precision
+    - kernel_select_source: *single_precision
+    - kernel_select_gemv: *single_precision
+  api: C
+
+- name: kernel_select_no_hipblaslt
+  category: quick
+  function:
+    - kernel_select_tensile: *single_precision
+    - kernel_select_check_only: *single_precision
+  api: C
+  use_hipblaslt: 0
+...
diff --git a/projects/rocblas/clients/gtest/kernel_select_test.cpp b/projects/rocblas/clients/gtest/kernel_select_test.cpp
deleted file mode 100644
index 489d6ecd5ddb..000000000000
--- a/projects/rocblas/clients/gtest/kernel_select_test.cpp
+++ /dev/null
@@ -1,360 +0,0 @@
-/* ************************************************************************
- * Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell cop-
- * ies of the Software, and to permit persons to whom the Software is furnished
- * to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IM-
- * PLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
- * FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
- * COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER
- * IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNE-
- * CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
- *
- * ************************************************************************ */
-
-// Targeted tests for ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select).
-// Self-contained: drives the C API directly with hipMalloc/hipMemcpy, sets env vars
-// per-test, and parses the trace file to verify schema-level invariants. The selected
-// kernel name depends on GPU arch and backend build, so we never golden-file the line.
-
-#include "client_utility.hpp"
-#include "device_vector.hpp"
-#include "host_vector.hpp"
-#include "rocblas_test.hpp"
-#include <cstdio>
-#include <cstdlib>
-#include <fstream>
-#include <gtest/gtest.h>
-#include <hip/hip_runtime.h>
-#include <rocblas/rocblas.h>
-#include <sstream>
-#include <string>
-#include <vector>
-
-#if __has_include(<filesystem>)
-#include <filesystem>
-namespace fs = std::filesystem;
-#elif __has_include(<experimental/filesystem>)
-#include <experimental/filesystem>
-namespace fs = std::experimental::filesystem;
-#else
-#error no filesystem found
-#endif
-
-#ifdef WIN32
-#define setenv(A, B, C) _putenv_s(A, B)
-#define unsetenv(A) _putenv_s(A, "")
-#endif
-
-namespace
-{
-    // Reads the entire trace file produced by a test into a single string for substring matching.
-    std::string read_trace(const std::string& path)
-    {
-        std::ifstream     ifs(path);
-        std::stringstream ss;
-        ss << ifs.rdbuf();
-        return ss.str();
-    }
-
-    // Counts occurrences of `needle` in `haystack`. Plain substring scan, no regex.
-    size_t count_occurrences(const std::string& haystack, const std::string& needle)
-    {
-        if(needle.empty())
-            return 0;
-        size_t n = 0;
-        for(size_t pos = 0; (pos = haystack.find(needle, pos)) != std::string::npos;
-            pos += needle.size())
-            ++n;
-        return n;
-    }
-
-    // Generates a unique, temp-dir-rooted trace-file path for this test invocation
-    // using rocblas_tempname().
-    std::string make_trace_path(const char* test_name)
-    {
-        static const std::string tmp_dir = rocblas_tempname();
-        const fs::path           fspath
-            = tmp_dir + std::string("kernel_select_") + test_name + std::string(".log");
-        const std::string p = fspath.generic_string();
-        std::remove(p.c_str());
-        return p;
-    }
-
-    // RAII: snapshot the env vars we mutate, restore on scope exit so test ordering doesn't matter.
-    struct env_guard
-    {
-        env_guard()
-            : m_layer(getenv("ROCBLAS_LAYER") ? getenv("ROCBLAS_LAYER") : "")
-            , m_path(getenv("ROCBLAS_LOG_TRACE_PATH") ? getenv("ROCBLAS_LOG_TRACE_PATH") : "")
-            , m_uselt(getenv("ROCBLAS_USE_HIPBLASLT") ? getenv("ROCBLAS_USE_HIPBLASLT") : "")
-            , m_layer_set(getenv("ROCBLAS_LAYER") != nullptr)
-            , m_path_set(getenv("ROCBLAS_LOG_TRACE_PATH") != nullptr)
-            , m_uselt_set(getenv("ROCBLAS_USE_HIPBLASLT") != nullptr)
-        {
-        }
-        ~env_guard()
-        {
-            auto restore = [](const char* name, const std::string& val, bool was_set) {
-                if(was_set)
-                    setenv(name, val.c_str(), 1);
-                else
-                    unsetenv(name);
-            };
-            restore("ROCBLAS_LAYER", m_layer, m_layer_set);
-            restore("ROCBLAS_LOG_TRACE_PATH", m_path, m_path_set);
-            restore("ROCBLAS_USE_HIPBLASLT", m_uselt, m_uselt_set);
-        }
-
-    private:
-        std::string m_layer, m_path, m_uselt;
-        bool        m_layer_set, m_path_set, m_uselt_set;
-    };
-
-} // namespace
-
-// (1) trsm decomposition produces many sub-problem lines, all attributed to rocblas_strsm.
-// Covers: the rocblas_internal_api_scope at trsm's _impl, multi-line decomposition output,
-// and demonstrates the primary use-case the feature was designed for.
-TEST(KernelSelect, TrsmDecompositionParentApi)
-{
-    env_guard         guard;
-    const std::string trace_path = make_trace_path("trsm");
-    constexpr int     m          = 512;
-    constexpr int     n          = 512;
-    constexpr float   alpha      = 1.0f;
-
-    setenv("ROCBLAS_LAYER", "0x10", 1);
-    setenv("ROCBLAS_LOG_TRACE_PATH", trace_path.c_str(), 1);
-
-    {
-        rocblas_handle handle;
-        ASSERT_EQ(rocblas_create_handle(&handle), rocblas_status_success);
-
-        host_vector<float> hA(size_t(m) * m);
-        host_vector<float> hB(size_t(m) * n);
-        for(int col = 0; col < m; ++col)
-            for(int row = 0; row < m; ++row)
-                hA[size_t(col) * m + row] = row == col ? float(m) : (row < col ? 0.5f : 0.0f);
-        for(size_t i = 0; i < size_t(m) * n; ++i)
-            hB[i] = 1.0f;
-
-        device_vector<float> dA(size_t(m) * m);
-        device_vector<float> dB(size_t(m) * n);
-        ASSERT_EQ(dA.transfer_from(hA), hipSuccess);
-        ASSERT_EQ(dB.transfer_from(hB), hipSuccess);
-
-        ASSERT_EQ(rocblas_strsm(handle,
-                                rocblas_side_left,
-                                rocblas_fill_upper,
-                                rocblas_operation_transpose,
-                                rocblas_diagonal_non_unit,
-                                m,
-                                n,
-                                &alpha,
-                                dA,
-                                m,
-                                dB,
-                                m),
-                  rocblas_status_success);
-
-        ASSERT_EQ(rocblas_destroy_handle(handle), rocblas_status_success);
-    }
-
-    const std::string trace = read_trace(trace_path);
-    // trsm at this size decomposes into >>1 sub-problems; require multiple lines all carrying
-    // the trsm parent. (Exact count depends on internal block size, so we only lower-bound.)
-    const size_t lines = count_occurrences(trace, "parent_api=rocblas_strsm");
-    EXPECT_GT(lines, size_t(1)) << "expected multiple sub-problems attributed to rocblas_strsm, "
-                                << "got " << lines << " in " << trace_path;
-    // Non-batched sub-problems (batch_count == 1) replay as gemm_ex.
-    EXPECT_NE(trace.find("-f gemm_ex"), std::string::npos);
-    EXPECT_NE(trace.find("# source="), std::string::npos);
-
-    std::remove(trace_path.c_str());
-}
-
-// (2) Forcing ROCBLAS_USE_HIPBLASLT=0 should route every GEMM through Tensile.
-// Covers: source=tensile branch, capture of solution->name(), and the Tensile kernel-name
-// format (Cijk_* is Tensile's canonical naming convention).
-TEST(KernelSelect, TensileBackendForced)
-{
-#ifndef BUILD_WITH_TENSILE
-    GTEST_SKIP() << "Tensile not built in; cannot test tensile backend path";
-#else
-    env_guard         guard;
-    const std::string trace_path = make_trace_path("tensile");
-    constexpr int     m          = 256;
-    constexpr int     n          = 256;
-    constexpr int     k          = 256;
-    constexpr float   alpha      = 1.0f;
-    constexpr float   beta       = 0.0f;
-
-    setenv("ROCBLAS_LAYER", "0x10", 1);
-    setenv("ROCBLAS_LOG_TRACE_PATH", trace_path.c_str(), 1);
-    setenv("ROCBLAS_USE_HIPBLASLT", "0", 1);
-
-    {
-        rocblas_handle handle;
-        ASSERT_EQ(rocblas_create_handle(&handle), rocblas_status_success);
-
-        device_vector<float> dA(size_t(m) * k);
-        device_vector<float> dB(size_t(k) * n);
-        device_vector<float> dC(size_t(m) * n);
-
-        ASSERT_EQ(rocblas_sgemm(handle,
-                                rocblas_operation_none,
-                                rocblas_operation_none,
-                                m,
-                                n,
-                                k,
-                                &alpha,
-                                dA,
-                                m,
-                                dB,
-                                k,
-                                &beta,
-                                dC,
-                                m),
-                  rocblas_status_success);
-
-        ASSERT_EQ(rocblas_destroy_handle(handle), rocblas_status_success);
-    }
-
-    const std::string trace = read_trace(trace_path);
-    EXPECT_NE(trace.find("source=tensile"), std::string::npos)
-        << "with ROCBLAS_USE_HIPBLASLT=0 expected source=tensile, got:\n"
-        << trace;
-    // Tensile solution names start with "Cijk_" by convention; if absent the capture is broken.
-    EXPECT_NE(trace.find("kernel=Cijk_"), std::string::npos)
-        << "expected Tensile-style kernel name, got:\n"
-        << trace;
-    EXPECT_EQ(trace.find("source=hipblaslt"), std::string::npos)
-        << "ROCBLAS_USE_HIPBLASLT=0 should suppress hipblaslt source, got:\n"
-        << trace;
-
-    std::remove(trace_path.c_str());
-#endif
-}
-
-// (3) Every emitted line must include the fallback_from= field, regardless of backend.
-// Covers: a weak schema check the existing logging test omits — guards against accidental
-// removal of the field or a code path that forgets to format it.
-TEST(KernelSelect, FallbackFromFieldAlwaysPresent)
-{
-    env_guard         guard;
-    const std::string trace_path = make_trace_path("fallback_field");
-    constexpr int     m          = 256;
-    constexpr int     n          = 256;
-    constexpr int     k          = 256;
-    constexpr float   alpha      = 1.0f;
-    constexpr float   beta       = 0.0f;
-
-    setenv("ROCBLAS_LAYER", "0x10", 1);
-    setenv("ROCBLAS_LOG_TRACE_PATH", trace_path.c_str(), 1);
-
-    {
-        rocblas_handle handle;
-        ASSERT_EQ(rocblas_create_handle(&handle), rocblas_status_success);
-
-        device_vector<float> dA(size_t(m) * k);
-        device_vector<float> dB(size_t(k) * n);
-        device_vector<float> dC(size_t(m) * n);
-
-        ASSERT_EQ(rocblas_sgemm(handle,
-                                rocblas_operation_none,
-                                rocblas_operation_none,
-                                m,
-                                n,
-                                k,
-                                &alpha,
-                                dA,
-                                m,
-                                dB,
-                                k,
-                                &beta,
-                                dC,
-                                m),
-                  rocblas_status_success);
-
-        ASSERT_EQ(rocblas_destroy_handle(handle), rocblas_status_success);
-    }
-
-    const std::string trace = read_trace(trace_path);
-    ASSERT_NE(trace.find("# source="), std::string::npos)
-        << "no kernel-select line emitted, trace was:\n"
-        << trace;
-    // Every kernel-select line should have the field — counting source= vs fallback_from= gives
-    // a uniform per-line check without per-line parsing.
-    const size_t source_lines   = count_occurrences(trace, "# source=");
-    const size_t fallback_lines = count_occurrences(trace, "fallback_from=");
-    EXPECT_EQ(source_lines, fallback_lines)
-        << "every kernel-select line must carry fallback_from=, got " << source_lines
-        << " source= but " << fallback_lines << " fallback_from=";
-
-    std::remove(trace_path.c_str());
-}
-
-// (4) Bits 0x8 and 0x10 are independent and can be combined; the same trace stream carries both.
-// Covers: env-parse OR of multiple bits, and the back-compat invariant that 0x8 still works.
-TEST(KernelSelect, CombinedWithInternalLayer)
-{
-    env_guard         guard;
-    const std::string trace_path = make_trace_path("combined");
-    constexpr int     m          = 256;
-    constexpr int     n          = 256;
-    constexpr int     k          = 256;
-    constexpr float   alpha      = 1.0f;
-    constexpr float   beta       = 0.0f;
-
-    setenv("ROCBLAS_LAYER", "0x18", 1); // 0x8 internal + 0x10 kernel-select
-    setenv("ROCBLAS_LOG_TRACE_PATH", trace_path.c_str(), 1);
-
-    {
-        rocblas_handle handle;
-        ASSERT_EQ(rocblas_create_handle(&handle), rocblas_status_success);
-
-        device_vector<float> dA(size_t(m) * k);
-        device_vector<float> dB(size_t(k) * n);
-        device_vector<float> dC(size_t(m) * n);
-
-        ASSERT_EQ(rocblas_sgemm(handle,
-                                rocblas_operation_none,
-                                rocblas_operation_none,
-                                m,
-                                n,
-                                k,
-                                &alpha,
-                                dA,
-                                m,
-                                dB,
-                                k,
-                                &beta,
-                                dC,
-                                m),
-                  rocblas_status_success);
-
-        ASSERT_EQ(rocblas_destroy_handle(handle), rocblas_status_success);
-    }
-
-    const std::string trace = read_trace(trace_path);
-    // 0x8 emits one of these backend tags per GEMM.
-    const bool has_0x8 = trace.find("rocblas_gemm_hipblaslt_backend") != std::string::npos
-                         || trace.find("rocblas_gemm_tensile_backend") != std::string::npos
-                         || trace.find("rocblas_gemm_source_backend") != std::string::npos;
-    EXPECT_TRUE(has_0x8) << "0x8 internal backend line missing from combined trace:\n" << trace;
-    // 0x10 emits the bench-style line.
-    EXPECT_NE(trace.find("# source="), std::string::npos)
-        << "0x10 kernel-select line missing from combined trace:\n"
-        << trace;
-
-    std::remove(trace_path.c_str());
-}
diff --git a/projects/rocblas/clients/include/testing_kernel_select.hpp b/projects/rocblas/clients/include/testing_kernel_select.hpp
new file mode 100644
index 000000000000..b3eede78a62d
--- /dev/null
+++ b/projects/rocblas/clients/include/testing_kernel_select.hpp
@@ -0,0 +1,471 @@
+/* ************************************************************************
+ * Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell cop-
+ * ies of the Software, and to permit persons to whom the Software is furnished
+ * to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IM-
+ * PLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
+ * FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
+ * COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER
+ * IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNE-
+ * CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+ *
+ * ************************************************************************ */
+
+#pragma once
+
+// Tests for ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select). The selected kernel
+// depends on the GPU and the backends built in, so the lines are checked for schema and for the
+// fields each path controls, never against a golden file.
+
+#include "testing_common.hpp"
+
+#include <cstdio>
+#include <cstdlib>
+#include <fstream>
+#include <string>
+#include <vector>
+
+#ifdef WIN32
+#include <stdlib.h>
+#define setenv(A, B, C) _putenv_s(A, B)
+#define unsetenv(A) _putenv_s(A, "")
+#endif
+
+namespace kernel_select
+{
+    inline bool contains(const std::string& s, const char* needle)
+    {
+        return s.find(needle) != std::string::npos;
+    }
+
+    inline bool ends_with(const std::string& s, const std::string& suffix)
+    {
+        return s.size() >= suffix.size()
+               && s.compare(s.size() - suffix.size(), suffix.size(), suffix) == 0;
+    }
+
+    // Sets an environment variable for the lifetime of the object, then restores it.
+    class scoped_env
+    {
+        std::string m_name;
+        std::string m_old;
+        bool        m_was_set;
+
+    public:
+        scoped_env(const char* name, const std::string& value)
+            : m_name(name)
+            , m_was_set(getenv(name) != nullptr)
+        {
+            if(m_was_set)
+                m_old = getenv(name);
+            setenv(name, value.c_str(), 1);
+        }
+        ~scoped_env()
+        {
+            if(m_was_set)
+                setenv(m_name.c_str(), m_old.c_str(), 1);
+            else
+                unsetenv(m_name.c_str());
+        }
+        scoped_env(const scoped_env&)            = delete;
+        scoped_env& operator=(const scoped_env&) = delete;
+    };
+
+    struct trace
+    {
+        std::string              text;
+        std::vector<std::string> lines; // the kernel-select lines only
+    };
+
+    // Runs fn(handle) on a fresh handle with ROCBLAS_LAYER=layer and returns the trace.
+    template <typename F>
+    trace capture(const Arguments& arg, const char* layer, F&& fn)
+    {
+        static const std::string base  = rocblas_tempname();
+        static int               count = 0;
+        const std::string        path = base + "_kernel_select_" + std::to_string(count++) + ".log";
+        trace                    result;
+        {
+            scoped_env layer_env("ROCBLAS_LAYER", layer);
+            scoped_env path_env("ROCBLAS_LOG_TRACE_PATH", path);
+
+            rocblas_local_handle handle{arg};
+            fn(handle);
+        } // the handle destructor flushes the trace stream
+
+        std::ifstream ifs(path);
+        for(std::string line; std::getline(ifs, line);)
+        {
+            result.text += line + "\n";
+            if(contains(line, "# source="))
+                result.lines.push_back(line);
+        }
+        ifs.close();
+        std::remove(path.c_str());
+        return result;
+    }
+
+    // Checks the fields every kernel-select line must carry.
+    inline void expect_schema(const trace& t)
+    {
+        for(const auto& line : t.lines)
+        {
+            const std::string bench = line.substr(0, line.find('#'));
+            EXPECT_EQ(line.rfind("./rocblas-bench", 0), size_t(0)) << line;
+            EXPECT_TRUE(contains(bench, " -f gemm")) << line;
+            EXPECT_TRUE(contains(line, " fallback_from=")) << line;
+            EXPECT_TRUE(contains(line, " parent_api=")) << line;
+            // Only gemm_strided_batched_ex takes strides, as in rocblas-bench.
+            EXPECT_EQ(contains(bench, "--stride_a"), contains(bench, "-f gemm_strided_batched_ex "))
+                << line;
+        }
+    }
+
+    // Fills A, B and C of an m x n x k non-transposed GEMM with initialized data.
+    template <typename T>
+    struct gemm_data
+    {
+        host_matrix<T>   hA, hB, hC;
+        device_matrix<T> dA, dB, dC;
+
+        gemm_data(const Arguments& arg, int64_t m, int64_t n, int64_t k)
+            : hA(m, k, m)
+            , hB(k, n, k)
+            , hC(m, n, m)
+            , dA(m, k, m)
+            , dB(k, n, k)
+            , dC(m, n, m)
+        {
+            rocblas_init_matrix(
+                hA, arg, rocblas_client_never_set_nan, rocblas_client_general_matrix, true);
+            rocblas_init_matrix(
+                hB, arg, rocblas_client_never_set_nan, rocblas_client_general_matrix);
+            rocblas_init_matrix(
+                hC, arg, rocblas_client_never_set_nan, rocblas_client_general_matrix);
+            EXPECT_EQ(dA.transfer_from(hA), hipSuccess);
+            EXPECT_EQ(dB.transfer_from(hB), hipSuccess);
+            EXPECT_EQ(dC.transfer_from(hC), hipSuccess);
+        }
+    };
+
+    constexpr int64_t c_dim   = 64;
+    constexpr float   c_alpha = 1.0f;
+    constexpr float   c_beta  = 0.0f;
+
+    // Runs rocblas_gemm_ex on a c_dim x n x c_dim float problem with C == D.
+    inline rocblas_status run_sgemm_ex(rocblas_handle    handle,
+                                       gemm_data<float>& d,
+                                       int64_t           n,
+                                       rocblas_gemm_algo algo,
+                                       int32_t           solution_index,
+                                       uint32_t          flags)
+    {
+        return rocblas_gemm_ex(handle,
+                               rocblas_operation_none,
+                               rocblas_operation_none,
+                               c_dim,
+                               n,
+                               c_dim,
+                               &c_alpha,
+                               d.dA,
+                               rocblas_datatype_f32_r,
+                               c_dim,
+                               d.dB,
+                               rocblas_datatype_f32_r,
+                               c_dim,
+                               &c_beta,
+                               d.dC,
+                               rocblas_datatype_f32_r,
+                               c_dim,
+                               d.dC,
+                               rocblas_datatype_f32_r,
+                               c_dim,
+                               rocblas_datatype_f32_r,
+                               algo,
+                               solution_index,
+                               flags);
+    }
+} // namespace kernel_select
+
+// trsm decomposes into several GEMM sub-problems, all attributed to rocblas_strsm.
+inline void testing_kernel_select_trsm(const Arguments& arg)
+{
+    using namespace kernel_select;
+    constexpr int64_t m = 512;
+    constexpr int64_t n = 512;
+
+    HOST_MEMCHECK(host_matrix<float>, hA, (m, m, m));
+    HOST_MEMCHECK(host_matrix<float>, hB, (m, n, m));
+    DEVICE_MEMCHECK(device_matrix<float>, dA, (m, m, m));
+    DEVICE_MEMCHECK(device_matrix<float>, dB, (m, n, m));
+
+    // Diagonally dominant upper-triangular A keeps the solve well conditioned.
+    float* A = hA[0];
+    for(int64_t col = 0; col < m; ++col)
+        for(int64_t row = 0; row < m; ++row)
+            A[col * m + row] = row == col ? float(m) : (row < col ? 0.5f : 0.0f);
+    rocblas_init_matrix(hB, arg, rocblas_client_never_set_nan, rocblas_client_general_matrix, true);
+    CHECK_HIP_ERROR(dA.transfer_from(hA));
+    CHECK_HIP_ERROR(dB.transfer_from(hB));
+
+    const trace t = capture(arg, "0x10", [&](rocblas_handle handle) {
+        CHECK_ROCBLAS_ERROR(rocblas_strsm(handle,
+                                          rocblas_side_left,
+                                          rocblas_fill_upper,
+                                          rocblas_operation_transpose,
+                                          rocblas_diagonal_non_unit,
+                                          m,
+                                          n,
+                                          &c_alpha,
+                                          dA,
+                                          m,
+                                          dB,
+                                          m));
+    });
+
+    expect_schema(t);
+    EXPECT_GT(t.lines.size(), size_t(1)) << t.text;
+    for(const auto& line : t.lines)
+        EXPECT_TRUE(ends_with(line, "parent_api=rocblas_strsm")) << line;
+}
+
+// gemm, gemm_batched and gemm_strided_batched replay as gemm_ex, gemm_batched_ex and
+// gemm_strided_batched_ex respectively.
+inline void testing_kernel_select_layout(const Arguments& arg)
+{
+    using namespace kernel_select;
+    constexpr int64_t d           = c_dim;
+    constexpr int64_t batch_count = 3;
+    constexpr int64_t stride      = d * d;
+
+    {
+        gemm_data<float> g(arg, d, d, d);
+        const trace      t = capture(arg, "0x10", [&](rocblas_handle handle) {
+            CHECK_ROCBLAS_ERROR(rocblas_sgemm(handle,
+                                              rocblas_operation_none,
+                                              rocblas_operation_none,
+                                              d,
+                                              d,
+                                              d,
+                                              &c_alpha,
+                                              g.dA,
+                                              d,
+                                              g.dB,
+                                              d,
+                                              &c_beta,
+                                              g.dC,
+                                              d));
+        });
+        expect_schema(t);
+        ASSERT_EQ(t.lines.size(), size_t(1)) << t.text;
+        EXPECT_TRUE(contains(t.lines[0], "-f gemm_ex ")) << t.lines[0];
+        EXPECT_FALSE(contains(t.lines[0], "--batch_count")) << t.lines[0];
+        EXPECT_TRUE(ends_with(t.lines[0], "parent_api=rocblas_sgemm")) << t.lines[0];
+    }
+
+    {
+        HOST_MEMCHECK(host_batch_matrix<float>, hA, (d, d, d, batch_count));
+        HOST_MEMCHECK(host_batch_matrix<float>, hB, (d, d, d, batch_count));
+        HOST_MEMCHECK(host_batch_matrix<float>, hC, (d, d, d, batch_count));
+        DEVICE_MEMCHECK(device_batch_matrix<float>, dA, (d, d, d, batch_count));
+        DEVICE_MEMCHECK(device_batch_matrix<float>, dB, (d, d, d, batch_count));
+        DEVICE_MEMCHECK(device_batch_matrix<float>, dC, (d, d, d, batch_count));
+        rocblas_init_matrix(
+            hA, arg, rocblas_client_never_set_nan, rocblas_client_general_matrix, true);
+        rocblas_init_matrix(hB, arg, rocblas_client_never_set_nan, rocblas_client_general_matrix);
+        rocblas_init_matrix(hC, arg, rocblas_client_never_set_nan, rocblas_client_general_matrix);
+        CHECK_HIP_ERROR(dA.transfer_from(hA));
+        CHECK_HIP_ERROR(dB.transfer_from(hB));
+        CHECK_HIP_ERROR(dC.transfer_from(hC));
+
+        const trace t = capture(arg, "0x10", [&](rocblas_handle handle) {
+            CHECK_ROCBLAS_ERROR(rocblas_sgemm_batched(handle,
+                                                      rocblas_operation_none,
+                                                      rocblas_operation_none,
+                                                      d,
+                                                      d,
+                                                      d,
+                                                      &c_alpha,
+                                                      dA.ptr_on_device(),
+                                                      d,
+                                                      dB.ptr_on_device(),
+                                                      d,
+                                                      &c_beta,
+                                                      dC.ptr_on_device(),
+                                                      d,
+                                                      batch_count));
+        });
+        expect_schema(t);
+        ASSERT_EQ(t.lines.size(), size_t(1)) << t.text;
+        EXPECT_TRUE(contains(t.lines[0], "-f gemm_batched_ex ")) << t.lines[0];
+        EXPECT_TRUE(contains(t.lines[0], "--batch_count 3 ")) << t.lines[0];
+    }
+
+    {
+        HOST_MEMCHECK(host_strided_batch_matrix<float>, hA, (d, d, d, stride, batch_count));
+        HOST_MEMCHECK(host_strided_batch_matrix<float>, hB, (d, d, d, stride, batch_count));
+        HOST_MEMCHECK(host_strided_batch_matrix<float>, hC, (d, d, d, stride, batch_count));
+        DEVICE_MEMCHECK(device_strided_batch_matrix<float>, dA, (d, d, d, stride, batch_count));
+        DEVICE_MEMCHECK(device_strided_batch_matrix<float>, dB, (d, d, d, stride, batch_count));
+        DEVICE_MEMCHECK(device_strided_batch_matrix<float>, dC, (d, d, d, stride, batch_count));
+        rocblas_init_matrix(
+            hA, arg, rocblas_client_never_set_nan, rocblas_client_general_matrix, true);
+        rocblas_init_matrix(hB, arg, rocblas_client_never_set_nan, rocblas_client_general_matrix);
+        rocblas_init_matrix(hC, arg, rocblas_client_never_set_nan, rocblas_client_general_matrix);
+        CHECK_HIP_ERROR(dA.transfer_from(hA));
+        CHECK_HIP_ERROR(dB.transfer_from(hB));
+        CHECK_HIP_ERROR(dC.transfer_from(hC));
+
+        const trace t = capture(arg, "0x10", [&](rocblas_handle handle) {
+            CHECK_ROCBLAS_ERROR(rocblas_sgemm_strided_batched(handle,
+                                                              rocblas_operation_none,
+                                                              rocblas_operation_none,
+                                                              d,
+                                                              d,
+                                                              d,
+                                                              &c_alpha,
+                                                              dA,
+                                                              d,
+                                                              stride,
+                                                              dB,
+                                                              d,
+                                                              stride,
+                                                              &c_beta,
+                                                              dC,
+                                                              d,
+                                                              stride,
+                                                              batch_count));
+        });
+        expect_schema(t);
+        ASSERT_EQ(t.lines.size(), size_t(1)) << t.text;
+        EXPECT_TRUE(contains(t.lines[0], "-f gemm_strided_batched_ex ")) << t.lines[0];
+        EXPECT_TRUE(contains(t.lines[0], "--stride_a 4096 ")) << t.lines[0];
+        EXPECT_TRUE(contains(t.lines[0], "--batch_count 3 ")) << t.lines[0];
+    }
+}
+
+// Bits 0x8 and 0x10 are independent and share the trace stream.
+inline void testing_kernel_select_combined(const Arguments& arg)
+{
+    using namespace kernel_select;
+    gemm_data<float> g(arg, c_dim, c_dim, c_dim);
+
+    const trace t = capture(arg, "0x18", [&](rocblas_handle handle) {
+        CHECK_ROCBLAS_ERROR(
+            run_sgemm_ex(handle, g, c_dim, rocblas_gemm_algo_standard, 0, rocblas_gemm_flags_none));
+    });
+
+    expect_schema(t);
+    EXPECT_EQ(t.lines.size(), size_t(1)) << t.text;
+    EXPECT_TRUE(contains(t.text, "rocblas_gemm_hipblaslt_backend")
+                || contains(t.text, "rocblas_gemm_tensile_backend")
+                || contains(t.text, "rocblas_gemm_source_backend"))
+        << t.text;
+}
+
+// An explicit source solution index runs the rocBLAS source kernels without trying a backend.
+inline void testing_kernel_select_source(const Arguments& arg)
+{
+    using namespace kernel_select;
+    constexpr int32_t c_source_solution = -2;
+    gemm_data<float>  g(arg, c_dim, c_dim, c_dim);
+
+    const trace t = capture(arg, "0x10", [&](rocblas_handle handle) {
+        CHECK_ROCBLAS_ERROR(run_sgemm_ex(handle,
+                                         g,
+                                         c_dim,
+                                         rocblas_gemm_algo_solution_index,
+                                         c_source_solution,
+                                         rocblas_gemm_flags_none));
+    });
+
+    expect_schema(t);
+    ASSERT_EQ(t.lines.size(), size_t(1)) << t.text;
+    const std::string& line = t.lines[0];
+    EXPECT_TRUE(contains(line, "-f gemm_ex ")) << line;
+    EXPECT_TRUE(contains(line, "--solution_index -2 ")) << line;
+    EXPECT_TRUE(contains(
+        line, "# source=source kernel=rocblas_gemm_source_solution_64 fallback_from=none "))
+        << line;
+    EXPECT_TRUE(ends_with(line, "parent_api=rocblas_gemm_ex")) << line;
+}
+
+// An explicit GEMV solution index runs an n == 1 GEMM through the GEMV kernels.
+inline void testing_kernel_select_gemv(const Arguments& arg)
+{
+#if !defined(BUILD_WITH_TENSILE) && !defined(BUILD_WITH_HIPBLASLT)
+    GTEST_SKIP() << "gemm_ex routes to GEMV only when a GEMM backend is built in";
+#else
+    using namespace kernel_select;
+    constexpr int32_t c_gemv_solution = -9;
+    gemm_data<float>  g(arg, c_dim, 1, c_dim);
+
+    const trace t = capture(arg, "0x10", [&](rocblas_handle handle) {
+        CHECK_ROCBLAS_ERROR(run_sgemm_ex(handle,
+                                         g,
+                                         1,
+                                         rocblas_gemm_algo_solution_index,
+                                         c_gemv_solution,
+                                         rocblas_gemm_flags_none));
+    });
+
+    expect_schema(t);
+    ASSERT_EQ(t.lines.size(), size_t(1)) << t.text;
+    EXPECT_TRUE(contains(t.lines[0], " -n 1 ")) << t.lines[0];
+    EXPECT_TRUE(contains(t.lines[0],
+                         "# source=gemv kernel=rocblas_internal_gemv_launcher fallback_from=none "))
+        << t.lines[0];
+#endif
+}
+
+// With hipBLASLt off, Tensile runs the GEMM and its kernel name is reported.
+inline void testing_kernel_select_tensile(const Arguments& arg)
+{
+#ifndef BUILD_WITH_TENSILE
+    GTEST_SKIP() << "Tensile is not built in";
+#else
+    using namespace kernel_select;
+    gemm_data<float> g(arg, c_dim, c_dim, c_dim);
+
+    const trace t = capture(arg, "0x10", [&](rocblas_handle handle) {
+        CHECK_ROCBLAS_ERROR(
+            run_sgemm_ex(handle, g, c_dim, rocblas_gemm_algo_standard, 0, rocblas_gemm_flags_none));
+    });
+
+    expect_schema(t);
+    ASSERT_EQ(t.lines.size(), size_t(1)) << t.text;
+    EXPECT_TRUE(contains(t.lines[0], "# source=tensile kernel=Cijk_")) << t.lines[0];
+    EXPECT_TRUE(contains(t.lines[0], " fallback_from=none ")) << t.lines[0];
+#endif
+}
+
+// check_solution_index validates a solution without launching it, so nothing is logged.
+inline void testing_kernel_select_check_only(const Arguments& arg)
+{
+#ifndef BUILD_WITH_TENSILE
+    GTEST_SKIP() << "Tensile is not built in";
+#else
+    using namespace kernel_select;
+    gemm_data<float> g(arg, c_dim, c_dim, c_dim);
+
+    const trace t = capture(arg, "0x10", [&](rocblas_handle handle) {
+        CHECK_ROCBLAS_ERROR(run_sgemm_ex(handle,
+                                         g,
+                                         c_dim,
+                                         rocblas_gemm_algo_standard,
+                                         0,
+                                         rocblas_gemm_flags_check_solution_index));
+    });
+
+    EXPECT_TRUE(t.lines.empty()) << t.text;
+#endif
+}
diff --git a/projects/rocblas/clients/include/testing_logging.hpp b/projects/rocblas/clients/include/testing_logging.hpp
index 3e1df89f74c0..7639bf4573ff 100644
--- a/projects/rocblas/clients/include/testing_logging.hpp
+++ b/projects/rocblas/clients/include/testing_logging.hpp
@@ -683,7 +683,7 @@ void testing_logging(const Arguments& arg)
         bool found = false;
         for(std::string line; std::getline(ks_ifs, line);)
         {
-            if(line.find("-f gemm_strided_batched_ex") != std::string::npos
+            if(line.find("-f gemm_ex ") != std::string::npos
                && line.find("# source=") != std::string::npos
                && line.find("parent_api=") != std::string::npos)
             {
diff --git a/projects/rocblas/docs/how-to/logging-in-rocblas.rst b/projects/rocblas/docs/how-to/logging-in-rocblas.rst
index 2d9d8c9a81a9..7119397a3c9b 100644
--- a/projects/rocblas/docs/how-to/logging-in-rocblas.rst
+++ b/projects/rocblas/docs/how-to/logging-in-rocblas.rst
@@ -42,21 +42,28 @@ Internal API logging outputs information like the GEMM backend used for a partic
 Not all internal APIs are logged. The log output goes to the same stream as trace logging.
 
 Kernel-selection logging (``ROCBLAS_LAYER=0x10``) emits one ``rocblas-bench``-replayable line per
-internal GEMM sub-problem dispatched, including those produced by decomposed BLAS3 ops such as
-``trsm``, ``trmm``, ``syrk``, and friends. Each line ends in a ``#`` comment that names the source
-backend (``hipblaslt`` or ``tensile``), the selected kernel name (Tensile backend only), whether
-the call fell back from another backend, and the user-facing API that triggered it. The line
-itself is copy-pasteable into ``rocblas-bench`` to reproduce that single sub-problem in isolation.
-Output goes to the same stream as trace logging. Example:
+internal GEMM sub-problem that launches a kernel, including those produced by decomposed BLAS3 ops
+such as ``trsm``, ``trmm``, ``syrk``, and friends. The bench function is ``gemm_ex``,
+``gemm_batched_ex``, or ``gemm_strided_batched_ex``, matching how the sub-problem is laid out in
+memory. Each line ends in a ``#`` comment with these fields:
+
+* ``source``: the path that ran the sub-problem: ``hipblaslt``, ``tensile``, ``source``
+  (rocBLAS source kernels), or ``gemv`` (a GEMM with ``m == 1`` or ``n == 1`` run as a GEMV).
+* ``kernel``: the selected kernel or launcher. It is omitted for the hipBLASLt backend.
+* ``fallback_from``: the backends that rejected the sub-problem before it ran, for example
+  ``hipblaslt,tensile``, or ``none``.
+* ``parent_api``: the user-facing API that triggered the sub-problem.
+
+The part of the line before the ``#`` can be pasted into ``rocblas-bench`` to reproduce that single
+sub-problem in isolation. Output goes to the same stream as trace logging. Example:
 
 .. code-block:: text
 
-   ./rocblas-bench -f gemm_strided_batched_ex --transposeA N --transposeB N -m 4096 -n 4096 -k 4096 \
-     --alpha 1 --a_type f32_r --lda 4096 --stride_a 0 --b_type f32_r --ldb 4096 --stride_b 0 \
-     --beta 0 --c_type f32_r --ldc 4096 --stride_c 0 --d_type f32_r --ldd 4096 --stride_d 0 \
-     --batch_count 1 --compute_type f32_r --algo 0 --solution_index 0 --flags 0 \
-     # source=tensile kernel=Cijk_Ailk_Bljk_SB_MT128x128x8 \
-     fallback_from=hipblaslt parent_api=rocblas_strsm
+   ./rocblas-bench -f gemm_ex --transposeA N --transposeB N -m 4096 -n 4096 -k 4096 --alpha 1 \
+     --a_type f32_r --lda 4096 --b_type f32_r --ldb 4096 --beta 0 --c_type f32_r --ldc 4096 \
+     --d_type f32_r --ldd 4096 --compute_type f32_r --algo 0 --solution_index 0 --flags 0 \
+     # source=tensile kernel=Cijk_Ailk_Bljk_SB_MT128x128x8 fallback_from=hipblaslt \
+     parent_api=rocblas_strsm
 
 Bits can be combined: ``ROCBLAS_LAYER=0x19`` enables trace (``0x1``), internal (``0x8``), and
 kernel-selection (``0x10``) simultaneously.
diff --git a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
index 40db8a3527bb..eb51227bbeb7 100644
--- a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
+++ b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
@@ -201,6 +201,36 @@ rocblas_status rocblas_internal_gemm(rocblas_handle    handle,
 
 #endif // BUILD_WITH_TENSILE
 
+    const bool kernel_select_logging = handle->layer_mode & rocblas_layer_mode_log_kernel_select;
+#ifdef BUILD_WITH_TENSILE
+    const char* fallback_from
+        = kernel_select_logging ? rocblas_internal_take_kernel_select_fallback(handle) : "none";
+#else
+    const char* fallback_from = "none";
+#endif
+    auto log_kernel_select = [&](const char* kernel) {
+        if(kernel_select_logging)
+            rocblas_internal_log_kernel_select(
+                handle,
+                rocblas_gemm_kernel_select_problem<BATCHED>(trans_a,
+                                                            trans_b,
+                                                            m,
+                                                            n,
+                                                            k,
+                                                            alpha,
+                                                            lda,
+                                                            stride_a,
+                                                            ldb,
+                                                            stride_b,
+                                                            beta,
+                                                            ldc,
+                                                            stride_c,
+                                                            batch_count),
+                "source",
+                kernel,
+                fallback_from);
+    };
+
     bool backend_logging = handle->layer_mode & rocblas_layer_mode_log_internal;
     if(backend_logging)
     {
@@ -225,8 +255,11 @@ rocblas_status rocblas_internal_gemm(rocblas_handle    handle,
 
     if(k == 0 || (alpha && *alpha == 0))
     {
-        return rocblas_gemm_scale_launcher_64(
+        status = rocblas_gemm_scale_launcher_64(
             handle, m, n, *beta, C, offset_c, ldc, stride_c, batch_count);
+        if(status == rocblas_status_success)
+            log_kernel_select("rocblas_gemm_scale_launcher_64");
+        return status;
     }
 
     status = rocblas_status_success;
@@ -266,58 +299,7 @@ rocblas_status rocblas_internal_gemm(rocblas_handle    handle,
             return status;
     }
 
-    // ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select): bench-style line for the
-    // source-GEMM fallback path. Reached when Tensile returned not_implemented (or
-    // BUILD_WITH_TENSILE is off), so fallback_from is "tensile" when Tensile was built in, else "none".
-    if(status == rocblas_status_success
-       && (handle->layer_mode & rocblas_layer_mode_log_kernel_select))
-    {
-#ifdef BUILD_WITH_TENSILE
-        const char* fallback = "tensile";
-#else
-        const char* fallback = "none";
-#endif
-        const char* parent_api = handle->current_api_name ? handle->current_api_name : "unknown";
-
-        rocblas_internal_ostream metadata_field;
-        metadata_field << "# source=source kernel=rocblas_gemm_source_solution_64 fallback_from="
-                       << fallback << " parent_api=" << parent_api;
-        const std::string metadata = metadata_field.str();
-        const std::string alphas   = LOG_BENCH_SCALAR_VALUE(handle, alpha);
-        const std::string betas    = LOG_BENCH_SCALAR_VALUE(handle, beta);
-
-        rocblas_kernel_select_bench_line line{};
-        line.bench_prefix   = ROCBLAS_API_BENCH;
-        line.function       = batch_count > 1 ? "gemm_strided_batched_ex" : "gemm_ex";
-        line.trans_a        = rocblas_transpose_letter(trans_a);
-        line.trans_b        = rocblas_transpose_letter(trans_b);
-        line.m              = m;
-        line.n              = n;
-        line.k              = k;
-        line.alpha          = alphas.c_str();
-        line.a_type         = rocblas_precision_string<TScal>;
-        line.lda            = lda;
-        line.stride_a       = stride_a;
-        line.b_type         = rocblas_precision_string<TScal>;
-        line.ldb            = ldb;
-        line.stride_b       = stride_b;
-        line.beta           = betas.c_str();
-        line.c_type         = rocblas_precision_string<TScal>;
-        line.ldc            = ldc;
-        line.stride_c       = stride_c;
-        line.d_type         = rocblas_precision_string<TScal>;
-        line.ldd            = ldc;
-        line.stride_d       = stride_c;
-        line.batch_count    = batch_count;
-        line.compute_type   = rocblas_precision_string<TScal>;
-        line.algo           = 0;
-        line.solution_index = 0;
-        line.flags          = 0;
-        line.metadata       = metadata.c_str();
-
-        rocblas_internal_logger logger;
-        logger.log_kernel_select_bench(handle, line);
-    }
+    log_kernel_select("rocblas_gemm_source_solution_64");
 
     return status;
 }
diff --git a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.hpp b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.hpp
index c64d88d88820..77743ec7fe6c 100644
--- a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.hpp
+++ b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.hpp
@@ -24,7 +24,51 @@
 
 #include "blas3/rocblas_gemm.hpp"
 #include "int64_helpers.hpp"
+#include "logging.hpp"
 #include "src64/blas3/rocblas_gemm_64.hpp"
+#include "utility.hpp"
+
+// Describes a rocblas_internal_gemm{,_64} problem, whose output D is C, for 0x10 logging.
+// alpha and beta must be host pointers.
+template <bool BATCHED, typename T>
+rocblas_kernel_select_problem rocblas_gemm_kernel_select_problem(rocblas_operation trans_a,
+                                                                 rocblas_operation trans_b,
+                                                                 int64_t           m,
+                                                                 int64_t           n,
+                                                                 int64_t           k,
+                                                                 const T*          alpha,
+                                                                 int64_t           lda,
+                                                                 rocblas_stride    stride_a,
+                                                                 int64_t           ldb,
+                                                                 rocblas_stride    stride_b,
+                                                                 const T*          beta,
+                                                                 int64_t           ldc,
+                                                                 rocblas_stride    stride_c,
+                                                                 int64_t           batch_count)
+{
+    rocblas_kernel_select_problem p;
+    p.trans_a       = trans_a;
+    p.trans_b       = trans_b;
+    p.m             = m;
+    p.n             = n;
+    p.k             = k;
+    p.alpha         = alpha;
+    p.beta          = beta;
+    p.input_type    = rocblas_datatype_from_type<T>;
+    p.output_type   = rocblas_datatype_from_type<T>;
+    p.compute_type  = rocblas_datatype_from_type<T>;
+    p.lda           = lda;
+    p.ldb           = ldb;
+    p.ldc           = ldc;
+    p.ldd           = ldc;
+    p.stride_a      = stride_a;
+    p.stride_b      = stride_b;
+    p.stride_c      = stride_c;
+    p.stride_d      = stride_c;
+    p.batch_count   = batch_count;
+    p.pointer_array = BATCHED;
+    return p;
+}
 
 template <typename T>
 ROCBLAS_INTERNAL_EXPORT_NOINLINE rocblas_status
diff --git a/projects/rocblas/library/src/blas3/rocblas_gemm_grouped_batched_imp.hpp b/projects/rocblas/library/src/blas3/rocblas_gemm_grouped_batched_imp.hpp
index bb6f2e0b66ef..20f7d07742ca 100644
--- a/projects/rocblas/library/src/blas3/rocblas_gemm_grouped_batched_imp.hpp
+++ b/projects/rocblas/library/src/blas3/rocblas_gemm_grouped_batched_imp.hpp
@@ -98,6 +98,9 @@ namespace
     {
         if(!handle)
             return rocblas_status_invalid_handle;
+
+        rocblas_internal_api_scope api_scope(handle, rocblas_gemm_grouped_batched_name<T>);
+
         RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
 
         std::vector<T> alpha_h, beta_h;
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex.hpp b/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex.hpp
index e298135b6608..04417a8ff109 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex.hpp
@@ -316,6 +316,11 @@ bool rocblas_use_gemv_in_gemm(rocblas_handle    handle,
 
     return false;
 }
+// Element type of T*, const T*, T* const* or const T* const*.
+template <typename P>
+using rocblas_gemm_elem_t
+    = std::remove_cv_t<std::remove_pointer_t<std::remove_cv_t<std::remove_pointer_t<P>>>>;
+
 /*
  * Runs a GEMM whose m or n is 1 as a GEMV. The caller must have established
  * that this is legal (rocblas_can_use_gemv_in_gemm) and desirable
@@ -345,6 +350,7 @@ rocblas_status rocblas_gemv_in_gemm_launcher(rocblas_handle    handle,
                                              rocblas_stride    stride_d,
                                              rocblas_int       batch_count)
 {
+    rocblas_status status;
     if(n == 1)
     {
         // If transB is transpose, then just use ldb as increment. Currently can't handle trans_b as conjugate.
@@ -353,27 +359,28 @@ rocblas_status rocblas_gemv_in_gemm_launcher(rocblas_handle    handle,
         // gemm and gemv handle transA differently
         rocblas_int m_gemv = trans_a == rocblas_operation_none ? m : k;
         rocblas_int n_gemv = trans_a == rocblas_operation_none ? k : m;
-        return rocblas_internal_gemv_launcher(handle,
-                                              trans_a,
-                                              m_gemv,
-                                              n_gemv,
-                                              alpha,
-                                              0,
-                                              a,
-                                              offset_a,
-                                              lda,
-                                              stride_a,
-                                              b,
-                                              offset_b,
-                                              incx_gemv,
-                                              stride_b,
-                                              beta,
-                                              0,
-                                              d,
-                                              offset_d,
-                                              1,
-                                              stride_d,
-                                              batch_count);
+
+        status = rocblas_internal_gemv_launcher(handle,
+                                                trans_a,
+                                                m_gemv,
+                                                n_gemv,
+                                                alpha,
+                                                0,
+                                                a,
+                                                offset_a,
+                                                lda,
+                                                stride_a,
+                                                b,
+                                                offset_b,
+                                                incx_gemv,
+                                                stride_b,
+                                                beta,
+                                                0,
+                                                d,
+                                                offset_d,
+                                                1,
+                                                stride_d,
+                                                batch_count);
     }
     else if(m == 1)
     {
@@ -386,30 +393,60 @@ rocblas_status rocblas_gemv_in_gemm_launcher(rocblas_handle    handle,
         // gemm and gemv handle transA differently
         rocblas_int m_gemv = transA_gemv == rocblas_operation_none ? n : k;
         rocblas_int n_gemv = transA_gemv == rocblas_operation_none ? k : n;
-        return rocblas_internal_gemv_launcher(handle,
-                                              transA_gemv,
-                                              m_gemv,
-                                              n_gemv,
-                                              alpha,
-                                              0,
-                                              b,
-                                              offset_b,
-                                              ldb,
-                                              stride_b,
-                                              a,
-                                              offset_a,
-                                              incx_gemv,
-                                              stride_a,
-                                              beta,
-                                              0,
-                                              d,
-                                              offset_d,
-                                              ldd,
-                                              stride_d,
-                                              batch_count);
+
+        status = rocblas_internal_gemv_launcher(handle,
+                                                transA_gemv,
+                                                m_gemv,
+                                                n_gemv,
+                                                alpha,
+                                                0,
+                                                b,
+                                                offset_b,
+                                                ldb,
+                                                stride_b,
+                                                a,
+                                                offset_a,
+                                                incx_gemv,
+                                                stride_a,
+                                                beta,
+                                                0,
+                                                d,
+                                                offset_d,
+                                                ldd,
+                                                stride_d,
+                                                batch_count);
     }
+    else
+        return rocblas_status_internal_error; // rocblas_use_gemv_in_gemm() requires m == 1 || n == 1
 
-    return rocblas_status_internal_error; // rocblas_use_gemv_in_gemm() requires m == 1 || n == 1
+    if(status == rocblas_status_success
+       && (handle->layer_mode & rocblas_layer_mode_log_kernel_select))
+    {
+        rocblas_kernel_select_problem p;
+        p.trans_a       = trans_a;
+        p.trans_b       = trans_b;
+        p.m             = m;
+        p.n             = n;
+        p.k             = k;
+        p.alpha         = alpha;
+        p.beta          = beta;
+        p.input_type    = rocblas_datatype_from_type<rocblas_gemm_elem_t<TiConstPtr>>;
+        p.output_type   = rocblas_datatype_from_type<rocblas_gemm_elem_t<ToPtr>>;
+        p.compute_type  = rocblas_datatype_from_type<rocblas_gemm_elem_t<TScal>>;
+        p.lda           = lda;
+        p.ldb           = ldb;
+        p.ldc           = ldd;
+        p.ldd           = ldd;
+        p.stride_a      = stride_a;
+        p.stride_b      = stride_b;
+        p.stride_c      = stride_d;
+        p.stride_d      = stride_d;
+        p.batch_count   = batch_count;
+        p.pointer_array = std::is_pointer_v<std::remove_pointer_t<TiConstPtr>>;
+        rocblas_internal_log_kernel_select(
+            handle, p, "gemv", "rocblas_internal_gemv_launcher", "none");
+    }
+    return status;
 }
 
 template <bool BATCHED>
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_gemm_grouped_batched_ex_imp.hpp b/projects/rocblas/library/src/blas_ex/rocblas_gemm_grouped_batched_ex_imp.hpp
index d35feae51c1a..5a789694dc35 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_gemm_grouped_batched_ex_imp.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_gemm_grouped_batched_ex_imp.hpp
@@ -99,6 +99,9 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle,
+                                             ROCBLAS_API_STR(rocblas_gemm_grouped_batched_ex));
+
         const bool HPA = compute_type == rocblas_datatype_f32_r
                          && (a_type == rocblas_datatype_f16_r || a_type == rocblas_datatype_bf16_r);
 
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_herk_ex_imp.hpp b/projects/rocblas/library/src/blas_ex/rocblas_herk_ex_imp.hpp
index 3d8bd3efa307..f0eee52cee83 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_herk_ex_imp.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_herk_ex_imp.hpp
@@ -51,6 +51,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, ROCBLAS_API_STR(rocblas_herk_ex));
+
         constexpr bool HERM = true;
 
         //Check if the handle is in the device memory size query, as there are two algorithms one which requires extra workspace memory and one which doesn't
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_syrk_ex_imp.hpp b/projects/rocblas/library/src/blas_ex/rocblas_syrk_ex_imp.hpp
index 79305cfbd932..c8466f9a2855 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_syrk_ex_imp.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_syrk_ex_imp.hpp
@@ -50,6 +50,8 @@ namespace
         if(!handle)
             return rocblas_status_invalid_handle;
 
+        rocblas_internal_api_scope api_scope(handle, ROCBLAS_API_STR(rocblas_syrk_ex));
+
         //Check if the handle is in the device memory size query, as there are two algorithms one which requires extra workspace memory and one which doesn't
         if(handle->is_device_memory_size_query())
         {
diff --git a/projects/rocblas/library/src/include/handle.hpp b/projects/rocblas/library/src/include/handle.hpp
index 6ba0460f4aa1..28e888b64ead 100644
--- a/projects/rocblas/library/src/include/handle.hpp
+++ b/projects/rocblas/library/src/include/handle.hpp
@@ -428,7 +428,13 @@ struct _rocblas_handle
     // at the top of each *_impl and cleared on scope exit. Used by rocblas_layer_mode_log_kernel_select
     // (0x10) to annotate internal-GEMM trace lines with their parent API. Null when no API is active.
     const char* current_api_name = nullptr;
-    void        init_check_numerics();
+
+    // Backends that returned not_implemented for the most recent GEMM sub-problem, e.g.
+    // "hipblaslt,tensile". Written by runContractionProblem when 0x10 logging is on and consumed
+    // by the source-GEMM fallback through rocblas_internal_take_kernel_select_fallback().
+    const char* kernel_select_fallback_from = nullptr;
+
+    void init_check_numerics();
 
     // data pointer for rocSOLVER
     std::shared_ptr<void> data_ptr;
@@ -947,7 +953,10 @@ struct rocblas_internal_api_scope
         , m_prev(m_handle ? m_handle->current_api_name : nullptr)
     {
         if(m_handle)
-            m_handle->current_api_name = name;
+        {
+            m_handle->current_api_name            = name;
+            m_handle->kernel_select_fallback_from = nullptr;
+        }
     }
     ~rocblas_internal_api_scope()
     {
@@ -962,6 +971,14 @@ struct rocblas_internal_api_scope
     const char*    m_prev;
 };
 
+// Returns and clears the backends that rejected the preceding GEMM sub-problem ("none" if none did).
+inline const char* rocblas_internal_take_kernel_select_fallback(rocblas_handle handle)
+{
+    const char* from                    = handle->kernel_select_fallback_from;
+    handle->kernel_select_fallback_from = nullptr;
+    return from ? from : "none";
+}
+
 // For functions which don't use temporary device memory, and won't be likely
 // to use them in the future, the RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle)
 // macro can be used to return from a rocblas function with a requested size of 0.
diff --git a/projects/rocblas/library/src/include/logging.hpp b/projects/rocblas/library/src/include/logging.hpp
index 51429983551f..6395b1b44ab1 100644
--- a/projects/rocblas/library/src/include/logging.hpp
+++ b/projects/rocblas/library/src/include/logging.hpp
@@ -206,42 +206,46 @@ double rocblas_internal_value_category(const T* beta, rocblas_datatype compute_t
 
 extern const char* c_rocblas_internal;
 
-// Fields for one ROCBLAS_LAYER=0x10 kernel-select bench line. Populated by the internal-GEMM
-// paths (tensile_host.cpp, gemm_templates.cpp) and rendered by
-// rocblas_internal_logger::log_kernel_select_bench, which owns the replay-order argument list.
-// Strings are non-owning views; the caller must keep the backing storage alive across the call.
-struct rocblas_kernel_select_bench_line
+// One internal GEMM sub-problem as reported by ROCBLAS_LAYER=0x10
+// (rocblas_layer_mode_log_kernel_select). A and B share input_type; C and D share output_type.
+// alpha and beta must point to host memory of compute_type.
+struct rocblas_kernel_select_problem
 {
-    const char* bench_prefix; // e.g. ROCBLAS_API_BENCH
-    // "gemm_ex" for a single sub-problem, "gemm_strided_batched_ex" when batched/strided.
-    const char* function;
-    char        trans_a;
-    char        trans_b;
-    int64_t     m;
-    int64_t     n;
-    int64_t     k;
-    const char* alpha; // preformatted "--alpha <v>"
-    const char* a_type;
-    int64_t     lda;
-    rocblas_stride stride_a;
-    const char* b_type;
-    int64_t     ldb;
-    rocblas_stride stride_b;
-    const char* beta; // preformatted "--beta <v>"
-    const char* c_type;
-    int64_t     ldc;
-    rocblas_stride stride_c;
-    const char* d_type;
-    int64_t     ldd;
-    rocblas_stride stride_d;
-    int64_t     batch_count;
-    const char* compute_type;
-    int32_t     algo;
-    int32_t     solution_index;
-    uint32_t    flags;
-    const char* metadata; // trailing "# source=... kernel=... fallback_from=... parent_api=..."
+    rocblas_operation trans_a        = rocblas_operation_none;
+    rocblas_operation trans_b        = rocblas_operation_none;
+    int64_t           m              = 0;
+    int64_t           n              = 0;
+    int64_t           k              = 0;
+    const void*       alpha          = nullptr;
+    const void*       beta           = nullptr;
+    rocblas_datatype  input_type     = rocblas_datatype_invalid;
+    rocblas_datatype  output_type    = rocblas_datatype_invalid;
+    rocblas_datatype  compute_type   = rocblas_datatype_invalid;
+    int64_t           lda            = 0;
+    int64_t           ldb            = 0;
+    int64_t           ldc            = 0;
+    int64_t           ldd            = 0;
+    rocblas_stride    stride_a       = 0;
+    rocblas_stride    stride_b       = 0;
+    rocblas_stride    stride_c       = 0;
+    rocblas_stride    stride_d       = 0;
+    int64_t           batch_count    = 1;
+    bool              pointer_array  = false;
+    rocblas_gemm_algo algo           = rocblas_gemm_algo_standard;
+    int32_t           solution_index = 0;
+    uint32_t          flags          = 0;
 };
 
+// Writes one rocblas-bench-replayable line for `prob` to the trace stream, followed by a
+// `# source=... kernel=... fallback_from=... parent_api=...` comment. The bench function is
+// gemm_ex, gemm_batched_ex or gemm_strided_batched_ex, matching how the sub-problem is laid out.
+// `kernel` may be null when the backend does not expose a kernel identifier.
+void rocblas_internal_log_kernel_select(rocblas_handle                       handle,
+                                        const rocblas_kernel_select_problem& prob,
+                                        const char*                          source,
+                                        const char*                          kernel,
+                                        const char*                          fallback_from);
+
 /******************************************************************
  * ROCBLAS LOGGER *
  ******************************************************************/
@@ -286,78 +290,6 @@ class rocblas_internal_logger
         log_arguments(*handle->log_bench_os, " ", std::forward<Ts>(xs)..., atomics_str);
     }
 
-    // if kernel-select logging is turned on with
-    // (handle->layer_mode & rocblas_layer_mode_log_kernel_select) != 0
-    // log_kernel_select emits one space-separated line per internal GEMM sub-problem to the
-    // trace stream. The line is formatted like a rocblas-bench invocation so it can be copied
-    // and replayed directly; trailing kernel-selection metadata is conventionally placed in a
-    // `# source=... kernel=... fallback_from=... parent_api=...` comment by the caller.
-    template <typename... Ts>
-    void log_kernel_select(rocblas_handle handle, Ts&&... xs)
-    {
-        log_arguments(*handle->log_trace_os, " ", std::forward<Ts>(xs)...);
-    }
-
-    // Emits a single rocblas-bench-replayable kernel-select line from a filled-in
-    // rocblas_kernel_select_bench_line. This centralizes the long, replay-order-sensitive
-    // argument list so the Tensile (tensile_host.cpp) and source-GEMM fallback
-    // (gemm_templates.cpp) paths cannot drift out of sync.
-    void log_kernel_select_bench(rocblas_handle handle, const rocblas_kernel_select_bench_line& b)
-    {
-        log_arguments(*handle->log_trace_os,
-                      " ",
-                      b.bench_prefix,
-                      "-f",
-                      b.function,
-                      "--transposeA",
-                      b.trans_a,
-                      "--transposeB",
-                      b.trans_b,
-                      "-m",
-                      b.m,
-                      "-n",
-                      b.n,
-                      "-k",
-                      b.k,
-                      b.alpha,
-                      "--a_type",
-                      b.a_type,
-                      "--lda",
-                      b.lda,
-                      "--stride_a",
-                      b.stride_a,
-                      "--b_type",
-                      b.b_type,
-                      "--ldb",
-                      b.ldb,
-                      "--stride_b",
-                      b.stride_b,
-                      b.beta,
-                      "--c_type",
-                      b.c_type,
-                      "--ldc",
-                      b.ldc,
-                      "--stride_c",
-                      b.stride_c,
-                      "--d_type",
-                      b.d_type,
-                      "--ldd",
-                      b.ldd,
-                      "--stride_d",
-                      b.stride_d,
-                      "--batch_count",
-                      b.batch_count,
-                      "--compute_type",
-                      b.compute_type,
-                      "--algo",
-                      b.algo,
-                      "--solution_index",
-                      b.solution_index,
-                      "--flags",
-                      b.flags,
-                      b.metadata);
-    }
-
     // if profile logging is turned on with
     // (handle->layer_mode & rocblas_layer_mode_log_profile) != 0
     // log_profile will call rocblas_internal_argument_profile to profile actual arguments,
diff --git a/projects/rocblas/library/src/logging.cpp b/projects/rocblas/library/src/logging.cpp
index 03a1d824f477..232ddd11a362 100644
--- a/projects/rocblas/library/src/logging.cpp
+++ b/projects/rocblas/library/src/logging.cpp
@@ -20,6 +20,7 @@
  *
  * ************************************************************************ */
 
+#include <algorithm>
 #include <limits>
 
 #ifndef DISABLE_ROCTX
@@ -29,6 +30,7 @@
 #endif
 
 #include "logging.hpp"
+#include "utility.hpp"
 
 /*************************************************
  * Bench log scalar values pointed to by pointer *
@@ -300,3 +302,59 @@ template double rocblas_internal_value_category(const rocblas_float_complex* bet
                                                 rocblas_datatype             compute_type);
 template double rocblas_internal_value_category(const rocblas_double_complex* beta,
                                                 rocblas_datatype              compute_type);
+
+void rocblas_internal_log_kernel_select(rocblas_handle                       handle,
+                                        const rocblas_kernel_select_problem& p,
+                                        const char*                          source,
+                                        const char*                          kernel,
+                                        const char*                          fallback_from)
+{
+    std::string alphas, betas;
+    (void)rocblas_internal_log_bench_alpha_beta_ex(p.compute_type, p.alpha, p.beta, alphas, betas);
+
+    const bool strided = !p.pointer_array && p.batch_count > 1;
+    const bool batched = p.pointer_array || strided;
+
+    constexpr int64_t i32_max = std::numeric_limits<int32_t>::max();
+    const bool        api_64
+        = std::max({p.m, p.n, p.k, p.lda, p.ldb, p.ldc, p.ldd, p.batch_count}) > i32_max;
+
+    const char* in_type  = rocblas_datatype_string(p.input_type);
+    const char* out_type = rocblas_datatype_string(p.output_type);
+
+    rocblas_internal_ostream bench;
+    bench << "./rocblas-bench" << (api_64 ? " --api 1" : "") << " -f "
+          << (p.pointer_array ? "gemm_batched_ex"
+              : strided       ? "gemm_strided_batched_ex"
+                              : "gemm_ex")
+          << " --transposeA " << rocblas_transpose_letter(p.trans_a) << " --transposeB "
+          << rocblas_transpose_letter(p.trans_b) << " -m " << p.m << " -n " << p.n << " -k " << p.k
+          << " " << alphas << " --a_type " << in_type << " --lda " << p.lda;
+    if(strided)
+        bench << " --stride_a " << p.stride_a;
+    bench << " --b_type " << in_type << " --ldb " << p.ldb;
+    if(strided)
+        bench << " --stride_b " << p.stride_b;
+    bench << " " << betas << " --c_type " << out_type << " --ldc " << p.ldc;
+    if(strided)
+        bench << " --stride_c " << p.stride_c;
+    bench << " --d_type " << out_type << " --ldd " << p.ldd;
+    if(strided)
+        bench << " --stride_d " << p.stride_d;
+    if(batched)
+        bench << " --batch_count " << p.batch_count;
+    bench << " --compute_type " << rocblas_datatype_string(p.compute_type) << " --algo "
+          << int32_t(p.algo) << " --solution_index " << p.solution_index << " --flags " << p.flags;
+    if(handle->atomics_mode == rocblas_atomics_allowed)
+        bench << " --atomics_allowed";
+
+    rocblas_internal_ostream meta;
+    meta << "# source=" << source;
+    if(kernel && *kernel)
+        meta << " kernel=" << kernel;
+    meta << " fallback_from=" << fallback_from
+         << " parent_api=" << (handle->current_api_name ? handle->current_api_name : "unknown");
+
+    rocblas_internal_logger logger;
+    logger.log_arguments(*handle->log_trace_os, " ", bench.str(), meta.str());
+}
diff --git a/projects/rocblas/library/src/src64/blas3/rocblas_gemm_kernels_64.cpp b/projects/rocblas/library/src/src64/blas3/rocblas_gemm_kernels_64.cpp
index 162050fc41a0..bc3eed7ef8f8 100644
--- a/projects/rocblas/library/src/src64/blas3/rocblas_gemm_kernels_64.cpp
+++ b/projects/rocblas/library/src/src64/blas3/rocblas_gemm_kernels_64.cpp
@@ -169,6 +169,31 @@ rocblas_status rocblas_internal_gemm_64(rocblas_handle    handle,
                 }
             }
         }
+
+        if(status == rocblas_status_success
+           && (handle->layer_mode & rocblas_layer_mode_log_kernel_select))
+        {
+            rocblas_internal_log_kernel_select(
+                handle,
+                rocblas_gemm_kernel_select_problem<BATCHED>(trans_a,
+                                                            trans_b,
+                                                            m_64,
+                                                            n_64,
+                                                            k_64,
+                                                            alpha,
+                                                            lda_64,
+                                                            stride_a,
+                                                            ldb_64,
+                                                            stride_b,
+                                                            beta,
+                                                            ldc_64,
+                                                            stride_c,
+                                                            batch_count_64),
+                "source",
+                k_64 == 0 || (alpha && *alpha == 0) ? "rocblas_gemm_scale_launcher_64"
+                                                    : "rocblas_gemm_source_solution_64",
+                "none");
+        }
     }
 
     return status;
diff --git a/projects/rocblas/library/src/src64/blas_ex/rocblas_gemm_ex_kernels_64.cpp b/projects/rocblas/library/src/src64/blas_ex/rocblas_gemm_ex_kernels_64.cpp
index 80ffad8a5320..0af259cb0fcf 100644
--- a/projects/rocblas/library/src/src64/blas_ex/rocblas_gemm_ex_kernels_64.cpp
+++ b/projects/rocblas/library/src/src64/blas_ex/rocblas_gemm_ex_kernels_64.cpp
@@ -206,6 +206,39 @@ rocblas_status rocblas_internal_gemm_ex_typecasting_64(rocblas_handle     handle
             return status;
     }
 
+    if(handle->layer_mode & rocblas_layer_mode_log_kernel_select)
+    {
+        rocblas_kernel_select_problem p;
+        p.trans_a        = trans_a;
+        p.trans_b        = trans_b;
+        p.m              = m_64;
+        p.n              = n_64;
+        p.k              = k_64;
+        p.alpha          = alpha;
+        p.beta           = beta;
+        p.input_type     = rocblas_datatype_from_type<Ti>;
+        p.output_type    = rocblas_datatype_from_type<To>;
+        p.compute_type   = rocblas_datatype_from_type<TScal>;
+        p.lda            = lda_64;
+        p.ldb            = ldb_64;
+        p.ldc            = ldc_64;
+        p.ldd            = ldd_64;
+        p.stride_a       = stride_a;
+        p.stride_b       = stride_b;
+        p.stride_c       = stride_c;
+        p.stride_d       = stride_d;
+        p.batch_count    = batch_count_64;
+        p.pointer_array  = BATCHED;
+        p.algo           = algo;
+        p.solution_index = solution_index;
+        p.flags          = uint32_t(flags);
+        rocblas_internal_log_kernel_select(handle,
+                                           p,
+                                           "source",
+                                           "rocblas_gemm_source_solution_64",
+                                           rocblas_internal_take_kernel_select_fallback(handle));
+    }
+
     if constexpr(check_numerics_supported)
     {
         if(check_numerics)
diff --git a/projects/rocblas/library/src/tensile_host.cpp b/projects/rocblas/library/src/tensile_host.cpp
index 1769a9740e95..18f5cb4b7613 100644
--- a/projects/rocblas/library/src/tensile_host.cpp
+++ b/projects/rocblas/library/src/tensile_host.cpp
@@ -33,7 +33,6 @@
 #include "hipblaslt_host.hpp"
 #endif
 
-#include "int64_helpers.hpp"
 #include "logging.hpp"
 #include "tensile_host.hpp"
 
@@ -1567,60 +1566,61 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
                          prob.col_stride_d);
     }
 
-    // ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select):
-    // one rocblas-bench-replayable line per internal GEMM sub-problem, with kernel-selection
-    // metadata in a trailing `#` comment so the line itself remains copy-pasteable.
-    if(kernel_select_logging && status == rocblas_status_success)
+    if(kernel_select_logging)
     {
-        std::string alphas, betas;
-        (void)rocblas_internal_log_bench_alpha_beta_ex(
-            rocblas_datatype_from_type<Tc>, prob.alpha, prob.beta, alphas, betas);
-
-        const char* source   = hipblaslt_backend ? "hipblaslt" : "tensile";
-        const char* fallback = (hipblaslt_attempted && !hipblaslt_backend) ? "hipblaslt" : "none";
-        const char* parent_api
-            = prob.handle->current_api_name ? prob.handle->current_api_name : "unknown";
-
-        // Trailing kernel-selection metadata, kept in one `#` comment so the bench line stays
-        // copy-pasteable. For tensile we keep the kernel name only.
-        rocblas_internal_ostream metadata_field;
-        metadata_field << "# source=" << source;
-        if(!hipblaslt_backend && !selected_tensile_name.empty())
-            metadata_field << " kernel=" << selected_tensile_name;
-        metadata_field << " fallback_from=" << fallback << " parent_api=" << parent_api;
-        const std::string metadata = metadata_field.str();
-
-        rocblas_kernel_select_bench_line line{};
-        line.bench_prefix   = ROCBLAS_API_BENCH;
-        line.function       = prob.batch_count > 1 ? "gemm_strided_batched_ex" : "gemm_ex";
-        line.trans_a        = rocblas_transpose_letter(prob.trans_a);
-        line.trans_b        = rocblas_transpose_letter(prob.trans_b);
-        line.m              = prob.m;
-        line.n              = prob.n;
-        line.k              = prob.k;
-        line.alpha          = alphas.c_str();
-        line.a_type         = rocblas_datatype_string(rocblas_datatype_from_type<Ti>);
-        line.lda            = prob.col_stride_a;
-        line.stride_a       = prob.batch_stride_a;
-        line.b_type         = rocblas_datatype_string(rocblas_datatype_from_type<Ti>);
-        line.ldb            = prob.col_stride_b;
-        line.stride_b       = prob.batch_stride_b;
-        line.beta           = betas.c_str();
-        line.c_type         = rocblas_datatype_string(rocblas_datatype_from_type<To>);
-        line.ldc            = prob.col_stride_c;
-        line.stride_c       = prob.batch_stride_c;
-        line.d_type         = rocblas_datatype_string(rocblas_datatype_from_type<To>);
-        line.ldd            = prob.col_stride_d;
-        line.stride_d       = prob.batch_stride_d;
-        line.batch_count    = prob.batch_count;
-        line.compute_type   = rocblas_datatype_string(rocblas_datatype_from_type<Tc>);
-        line.algo           = (int32_t)algo;
-        line.solution_index = solution_index;
-        line.flags          = (uint32_t)prob.flags;
-        line.metadata       = metadata.c_str();
-
-        rocblas_internal_logger logger;
-        logger.log_kernel_select_bench(prob.handle, line);
+        // Queries and check_solution_index return success without launching anything. hipBLASLt
+        // honours check_solution_index only for algo == rocblas_gemm_algo_solution_index.
+        const bool check_only = (prob.flags & rocblas_gemm_flags_check_solution_index)
+                                && (!hipblaslt_backend || solutionBased);
+        const bool launched = !check_only && !prob.handle->get_solution_fitness_query()
+                              && !prob.handle->is_device_memory_size_query();
+
+        if(launched && status == rocblas_status_success)
+        {
+            rocblas_kernel_select_problem p;
+            p.trans_a        = prob.trans_a;
+            p.trans_b        = prob.trans_b;
+            p.m              = prob.m;
+            p.n              = prob.n;
+            p.k              = prob.k;
+            p.alpha          = prob.alpha;
+            p.beta           = prob.beta;
+            p.input_type     = rocblas_datatype_from_type<Ti>;
+            p.output_type    = rocblas_datatype_from_type<To>;
+            p.compute_type   = rocblas_datatype_from_type<Tc>;
+            p.lda            = prob.col_stride_a;
+            p.ldb            = prob.col_stride_b;
+            p.ldc            = prob.col_stride_c;
+            p.ldd            = prob.col_stride_d;
+            p.stride_a       = prob.batch_stride_a;
+            p.stride_b       = prob.batch_stride_b;
+            p.stride_c       = prob.batch_stride_c;
+            p.stride_d       = prob.batch_stride_d;
+            p.batch_count    = prob.batch_count;
+            p.pointer_array  = !prob.strided_batch;
+            p.algo           = algo;
+            p.solution_index = solution_index;
+            p.flags          = uint32_t(prob.flags);
+
+            rocblas_internal_log_kernel_select(
+                prob.handle,
+                p,
+                hipblaslt_backend ? "hipblaslt" : "tensile",
+                hipblaslt_backend ? nullptr : selected_tensile_name.c_str(),
+                hipblaslt_attempted && !hipblaslt_backend ? "hipblaslt" : "none");
+        }
+
+        // Tell the source-GEMM fallback which backends rejected this sub-problem.
+        const char* rejected = nullptr;
+        if(status == rocblas_status_not_implemented)
+        {
+#ifdef BUILD_WITH_TENSILE
+            rejected = hipblaslt_attempted ? "hipblaslt,tensile" : "tensile";
+#else
+            rejected = hipblaslt_attempted ? "hipblaslt" : nullptr;
+#endif
+        }
+        prob.handle->kernel_select_fallback_from = rejected;
     }
 
     return status;

From c0d773e2e8fa5fe53c0296290b2578c32676bebb Mon Sep 17 00:00:00 2001
From: Estevan Vedovelli <Estevan.Vedovelli@amd.com>
Date: Wed, 28 Oct 2026 15:42:51 -0400
Subject: [PATCH 08/10] Move tests into existing logging_mode_gtest

---
 projects/rocblas/clients/gtest/CMakeLists.txt |  1 -
 projects/rocblas/clients/gtest/aux_gtest.yaml |  1 -
 .../clients/gtest/kernel_select_gtest.yaml    | 23 ------------
 .../clients/gtest/logging_mode_gtest.cpp      | 36 +++++++++++++++++--
 .../clients/gtest/logging_mode_gtest.yaml     | 17 +++++++++
 5 files changed, 51 insertions(+), 27 deletions(-)
 delete mode 100644 projects/rocblas/clients/gtest/kernel_select_gtest.yaml

diff --git a/projects/rocblas/clients/gtest/CMakeLists.txt b/projects/rocblas/clients/gtest/CMakeLists.txt
index 61c0d97cae0e..94147c66f2b4 100644
--- a/projects/rocblas/clients/gtest/CMakeLists.txt
+++ b/projects/rocblas/clients/gtest/CMakeLists.txt
@@ -47,7 +47,6 @@ set(rocblas_no_tensile_test_source
     set_get_atomics_mode_gtest.cpp
     set_get_alpha_beta_stride_gtest.cpp
     logging_mode_gtest.cpp
-    kernel_select_gtest.cpp
     ostream_threadsafety_gtest.cpp
     set_get_vector_gtest.cpp
     set_get_matrix_gtest.cpp
diff --git a/projects/rocblas/clients/gtest/aux_gtest.yaml b/projects/rocblas/clients/gtest/aux_gtest.yaml
index a63fcd5b4f32..47319f7e4d24 100644
--- a/projects/rocblas/clients/gtest/aux_gtest.yaml
+++ b/projects/rocblas/clients/gtest/aux_gtest.yaml
@@ -2,7 +2,6 @@
 include: set_get_matrix_gtest.yaml
 include: set_get_vector_gtest.yaml
 include: logging_mode_gtest.yaml
-include: kernel_select_gtest.yaml
 include: set_get_pointer_mode_gtest.yaml
 include: set_get_atomics_mode_gtest.yaml
 include: set_get_alpha_beta_stride_gtest.yaml
diff --git a/projects/rocblas/clients/gtest/kernel_select_gtest.yaml b/projects/rocblas/clients/gtest/kernel_select_gtest.yaml
deleted file mode 100644
index cc234fea898c..000000000000
--- a/projects/rocblas/clients/gtest/kernel_select_gtest.yaml
+++ /dev/null
@@ -1,23 +0,0 @@
----
-include: rocblas_common.yaml
-include: known_bugs.yaml
-
-Tests:
-- name: kernel_select
-  category: quick
-  function:
-    - kernel_select_trsm: *single_precision
-    - kernel_select_layout: *single_precision
-    - kernel_select_combined: *single_precision
-    - kernel_select_source: *single_precision
-    - kernel_select_gemv: *single_precision
-  api: C
-
-- name: kernel_select_no_hipblaslt
-  category: quick
-  function:
-    - kernel_select_tensile: *single_precision
-    - kernel_select_check_only: *single_precision
-  api: C
-  use_hipblaslt: 0
-...
diff --git a/projects/rocblas/clients/gtest/logging_mode_gtest.cpp b/projects/rocblas/clients/gtest/logging_mode_gtest.cpp
index 11a65f40d7ba..3a25dc1669a4 100644
--- a/projects/rocblas/clients/gtest/logging_mode_gtest.cpp
+++ b/projects/rocblas/clients/gtest/logging_mode_gtest.cpp
@@ -1,5 +1,5 @@
 /* ************************************************************************
- * Copyright (C) 2018-2025 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (C) 2018-2026 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -23,6 +23,7 @@
 #include "rocblas_data.hpp"
 #include "rocblas_datatype2string.hpp"
 #include "rocblas_test.hpp"
+#include "testing_kernel_select.hpp"
 #include "testing_logging.hpp"
 #include "type_dispatch.hpp"
 #include <cctype>
@@ -53,11 +54,32 @@ namespace
         {
             if(!strcmp(arg.function, "logging"))
                 testing_logging<T>(arg);
+            else if constexpr(std::is_same_v<T, float>)
+            {
+                if(!strcmp(arg.function, "kernel_select_trsm"))
+                    testing_kernel_select_trsm(arg);
+                else if(!strcmp(arg.function, "kernel_select_layout"))
+                    testing_kernel_select_layout(arg);
+                else if(!strcmp(arg.function, "kernel_select_combined"))
+                    testing_kernel_select_combined(arg);
+                else if(!strcmp(arg.function, "kernel_select_source"))
+                    testing_kernel_select_source(arg);
+                else if(!strcmp(arg.function, "kernel_select_gemv"))
+                    testing_kernel_select_gemv(arg);
+                else if(!strcmp(arg.function, "kernel_select_tensile"))
+                    testing_kernel_select_tensile(arg);
+                else if(!strcmp(arg.function, "kernel_select_check_only"))
+                    testing_kernel_select_check_only(arg);
+                else
+                    FAIL() << "Internal error: Test called with unknown function: " << arg.function;
+            }
             else
                 FAIL() << "Internal error: Test called with unknown function: " << arg.function;
         }
     };
 
+    constexpr char c_kernel_select_prefix[] = "kernel_select_";
+
     struct logging : RocBLAS_Test<logging, logging_testing>
     {
         // Filter for which types apply to this suite
@@ -69,13 +91,23 @@ namespace
         // Filter for which functions apply to this suite
         static bool function_filter(const Arguments& arg)
         {
-            return !strcmp(arg.function, "logging");
+            return !strcmp(arg.function, "logging")
+                   || !strncmp(
+                       arg.function, c_kernel_select_prefix, sizeof(c_kernel_select_prefix) - 1);
         }
 
         // Google Test name suffix based on parameters
         static std::string name_suffix(const Arguments& arg)
         {
             RocBLAS_TestName<logging> name(arg.name);
+
+            if(strcmp(arg.function, "logging"))
+            {
+                name << '_' << arg.function + sizeof(c_kernel_select_prefix) - 1 << '_'
+                     << rocblas_datatype2string(arg.a_type);
+                return std::move(name);
+            }
+
             name << rocblas_datatype2string(arg.a_type);
 
             name << (arg.pointer_mode_host ? "_hostptr" : "_devptr");
diff --git a/projects/rocblas/clients/gtest/logging_mode_gtest.yaml b/projects/rocblas/clients/gtest/logging_mode_gtest.yaml
index 9047b25d6ea9..fefbcefd5d04 100644
--- a/projects/rocblas/clients/gtest/logging_mode_gtest.yaml
+++ b/projects/rocblas/clients/gtest/logging_mode_gtest.yaml
@@ -20,5 +20,22 @@ Tests:
   os_flags: [ LINUX ]
   gpu_arch: [ '120?' ] # expected backend for LAYER=8
 
+- name: logging_kernel_select
+  category: quick
+  function:
+    - kernel_select_trsm: *single_precision
+    - kernel_select_layout: *single_precision
+    - kernel_select_combined: *single_precision
+    - kernel_select_source: *single_precision
+    - kernel_select_gemv: *single_precision
+  api: C
+
+- name: logging_kernel_select_no_hipblaslt
+  category: quick
+  function:
+    - kernel_select_tensile: *single_precision
+    - kernel_select_check_only: *single_precision
+  api: C
+  use_hipblaslt: 0
 
 ...

From 370e79cc5705c782505978a89a06d6fe9ee3be75 Mon Sep 17 00:00:00 2001
From: Estevan Vedovelli <Estevan.Vedovelli@amd.com>
Date: Wed, 28 Oct 2026 15:48:41 -0400
Subject: [PATCH 09/10] fix clang-format error

---
 projects/rocblas/clients/include/testing_kernel_select.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/projects/rocblas/clients/include/testing_kernel_select.hpp b/projects/rocblas/clients/include/testing_kernel_select.hpp
index b3eede78a62d..702574f0e9f0 100644
--- a/projects/rocblas/clients/include/testing_kernel_select.hpp
+++ b/projects/rocblas/clients/include/testing_kernel_select.hpp
@@ -76,7 +76,7 @@ namespace kernel_select
             else
                 unsetenv(m_name.c_str());
         }
-        scoped_env(const scoped_env&)            = delete;
+        scoped_env(const scoped_env&) = delete;
         scoped_env& operator=(const scoped_env&) = delete;
     };
 

From 37fdf3bb4ca430998172c110dee9633951589647 Mon Sep 17 00:00:00 2001
From: Estevan Vedovelli <Estevan.Vedovelli@amd.com>
Date: Thu, 29 Oct 2026 11:03:15 -0400
Subject: [PATCH 10/10] Include review suggestions.

---
 .../src/blas3/Tensile/gemm_templates.cpp      | 37 ++++++-----
 .../src/blas3/Tensile/gemm_templates.hpp      | 30 ++++-----
 .../library/src/blas_ex/rocblas_gemm_ex.hpp   |  2 +-
 .../rocblas/library/src/include/logging.hpp   |  4 +-
 projects/rocblas/library/src/logging.cpp      |  2 +-
 .../src64/blas3/rocblas_gemm_kernels_64.cpp   | 40 ++++++------
 .../blas_ex/rocblas_gemm_ex_kernels_64.cpp    |  2 +-
 projects/rocblas/library/src/tensile_host.cpp | 63 +++++++++++--------
 8 files changed, 95 insertions(+), 85 deletions(-)

diff --git a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
index eb51227bbeb7..10848e0b29c9 100644
--- a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
+++ b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp
@@ -210,25 +210,24 @@ rocblas_status rocblas_internal_gemm(rocblas_handle    handle,
 #endif
     auto log_kernel_select = [&](const char* kernel) {
         if(kernel_select_logging)
-            rocblas_internal_log_kernel_select(
-                handle,
-                rocblas_gemm_kernel_select_problem<BATCHED>(trans_a,
-                                                            trans_b,
-                                                            m,
-                                                            n,
-                                                            k,
-                                                            alpha,
-                                                            lda,
-                                                            stride_a,
-                                                            ldb,
-                                                            stride_b,
-                                                            beta,
-                                                            ldc,
-                                                            stride_c,
-                                                            batch_count),
-                "source",
-                kernel,
-                fallback_from);
+            rocblas_internal_log_kernel_select(handle,
+                                               rocblas_make_gemm_problem<BATCHED>(trans_a,
+                                                                                  trans_b,
+                                                                                  m,
+                                                                                  n,
+                                                                                  k,
+                                                                                  alpha,
+                                                                                  lda,
+                                                                                  stride_a,
+                                                                                  ldb,
+                                                                                  stride_b,
+                                                                                  beta,
+                                                                                  ldc,
+                                                                                  stride_c,
+                                                                                  batch_count),
+                                               "source",
+                                               kernel,
+                                               fallback_from);
     };
 
     bool backend_logging = handle->layer_mode & rocblas_layer_mode_log_internal;
diff --git a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.hpp b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.hpp
index 77743ec7fe6c..48d82fa25d94 100644
--- a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.hpp
+++ b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.hpp
@@ -31,22 +31,22 @@
 // Describes a rocblas_internal_gemm{,_64} problem, whose output D is C, for 0x10 logging.
 // alpha and beta must be host pointers.
 template <bool BATCHED, typename T>
-rocblas_kernel_select_problem rocblas_gemm_kernel_select_problem(rocblas_operation trans_a,
-                                                                 rocblas_operation trans_b,
-                                                                 int64_t           m,
-                                                                 int64_t           n,
-                                                                 int64_t           k,
-                                                                 const T*          alpha,
-                                                                 int64_t           lda,
-                                                                 rocblas_stride    stride_a,
-                                                                 int64_t           ldb,
-                                                                 rocblas_stride    stride_b,
-                                                                 const T*          beta,
-                                                                 int64_t           ldc,
-                                                                 rocblas_stride    stride_c,
-                                                                 int64_t           batch_count)
+rocblas_internal_gemm_problem rocblas_make_gemm_problem(rocblas_operation trans_a,
+                                                        rocblas_operation trans_b,
+                                                        int64_t           m,
+                                                        int64_t           n,
+                                                        int64_t           k,
+                                                        const T*          alpha,
+                                                        int64_t           lda,
+                                                        rocblas_stride    stride_a,
+                                                        int64_t           ldb,
+                                                        rocblas_stride    stride_b,
+                                                        const T*          beta,
+                                                        int64_t           ldc,
+                                                        rocblas_stride    stride_c,
+                                                        int64_t           batch_count)
 {
-    rocblas_kernel_select_problem p;
+    rocblas_internal_gemm_problem p;
     p.trans_a       = trans_a;
     p.trans_b       = trans_b;
     p.m             = m;
diff --git a/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex.hpp b/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex.hpp
index 04417a8ff109..fab62966bd4e 100644
--- a/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex.hpp
+++ b/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex.hpp
@@ -422,7 +422,7 @@ rocblas_status rocblas_gemv_in_gemm_launcher(rocblas_handle    handle,
     if(status == rocblas_status_success
        && (handle->layer_mode & rocblas_layer_mode_log_kernel_select))
     {
-        rocblas_kernel_select_problem p;
+        rocblas_internal_gemm_problem p;
         p.trans_a       = trans_a;
         p.trans_b       = trans_b;
         p.m             = m;
diff --git a/projects/rocblas/library/src/include/logging.hpp b/projects/rocblas/library/src/include/logging.hpp
index 6395b1b44ab1..0f16b1e81466 100644
--- a/projects/rocblas/library/src/include/logging.hpp
+++ b/projects/rocblas/library/src/include/logging.hpp
@@ -209,7 +209,7 @@ extern const char* c_rocblas_internal;
 // One internal GEMM sub-problem as reported by ROCBLAS_LAYER=0x10
 // (rocblas_layer_mode_log_kernel_select). A and B share input_type; C and D share output_type.
 // alpha and beta must point to host memory of compute_type.
-struct rocblas_kernel_select_problem
+struct rocblas_internal_gemm_problem
 {
     rocblas_operation trans_a        = rocblas_operation_none;
     rocblas_operation trans_b        = rocblas_operation_none;
@@ -241,7 +241,7 @@ struct rocblas_kernel_select_problem
 // gemm_ex, gemm_batched_ex or gemm_strided_batched_ex, matching how the sub-problem is laid out.
 // `kernel` may be null when the backend does not expose a kernel identifier.
 void rocblas_internal_log_kernel_select(rocblas_handle                       handle,
-                                        const rocblas_kernel_select_problem& prob,
+                                        const rocblas_internal_gemm_problem& prob,
                                         const char*                          source,
                                         const char*                          kernel,
                                         const char*                          fallback_from);
diff --git a/projects/rocblas/library/src/logging.cpp b/projects/rocblas/library/src/logging.cpp
index 232ddd11a362..231f09aed0b3 100644
--- a/projects/rocblas/library/src/logging.cpp
+++ b/projects/rocblas/library/src/logging.cpp
@@ -304,7 +304,7 @@ template double rocblas_internal_value_category(const rocblas_double_complex* be
                                                 rocblas_datatype              compute_type);
 
 void rocblas_internal_log_kernel_select(rocblas_handle                       handle,
-                                        const rocblas_kernel_select_problem& p,
+                                        const rocblas_internal_gemm_problem& p,
                                         const char*                          source,
                                         const char*                          kernel,
                                         const char*                          fallback_from)
diff --git a/projects/rocblas/library/src/src64/blas3/rocblas_gemm_kernels_64.cpp b/projects/rocblas/library/src/src64/blas3/rocblas_gemm_kernels_64.cpp
index bc3eed7ef8f8..cbd285a88fb9 100644
--- a/projects/rocblas/library/src/src64/blas3/rocblas_gemm_kernels_64.cpp
+++ b/projects/rocblas/library/src/src64/blas3/rocblas_gemm_kernels_64.cpp
@@ -173,26 +173,26 @@ rocblas_status rocblas_internal_gemm_64(rocblas_handle    handle,
         if(status == rocblas_status_success
            && (handle->layer_mode & rocblas_layer_mode_log_kernel_select))
         {
-            rocblas_internal_log_kernel_select(
-                handle,
-                rocblas_gemm_kernel_select_problem<BATCHED>(trans_a,
-                                                            trans_b,
-                                                            m_64,
-                                                            n_64,
-                                                            k_64,
-                                                            alpha,
-                                                            lda_64,
-                                                            stride_a,
-                                                            ldb_64,
-                                                            stride_b,
-                                                            beta,
-                                                            ldc_64,
-                                                            stride_c,
-                                                            batch_count_64),
-                "source",
-                k_64 == 0 || (alpha && *alpha == 0) ? "rocblas_gemm_scale_launcher_64"
-                                                    : "rocblas_gemm_source_solution_64",
-                "none");
+            rocblas_internal_log_kernel_select(handle,
+                                               rocblas_make_gemm_problem<BATCHED>(trans_a,
+                                                                                  trans_b,
+                                                                                  m_64,
+                                                                                  n_64,
+                                                                                  k_64,
+                                                                                  alpha,
+                                                                                  lda_64,
+                                                                                  stride_a,
+                                                                                  ldb_64,
+                                                                                  stride_b,
+                                                                                  beta,
+                                                                                  ldc_64,
+                                                                                  stride_c,
+                                                                                  batch_count_64),
+                                               "source",
+                                               k_64 == 0 || (alpha && *alpha == 0)
+                                                   ? "rocblas_gemm_scale_launcher_64"
+                                                   : "rocblas_gemm_source_solution_64",
+                                               "none");
         }
     }
 
diff --git a/projects/rocblas/library/src/src64/blas_ex/rocblas_gemm_ex_kernels_64.cpp b/projects/rocblas/library/src/src64/blas_ex/rocblas_gemm_ex_kernels_64.cpp
index 25a6cc4afddb..445c13ead5e3 100644
--- a/projects/rocblas/library/src/src64/blas_ex/rocblas_gemm_ex_kernels_64.cpp
+++ b/projects/rocblas/library/src/src64/blas_ex/rocblas_gemm_ex_kernels_64.cpp
@@ -219,7 +219,7 @@ rocblas_status rocblas_internal_gemm_ex_typecasting_64(rocblas_handle     handle
 
     if(handle->layer_mode & rocblas_layer_mode_log_kernel_select)
     {
-        rocblas_kernel_select_problem p;
+        rocblas_internal_gemm_problem p;
         p.trans_a        = trans_a;
         p.trans_b        = trans_b;
         p.m              = m_64;
diff --git a/projects/rocblas/library/src/tensile_host.cpp b/projects/rocblas/library/src/tensile_host.cpp
index de603cc6a66c..d5156940afcd 100644
--- a/projects/rocblas/library/src/tensile_host.cpp
+++ b/projects/rocblas/library/src/tensile_host.cpp
@@ -1344,6 +1344,42 @@ bool useHipBLASLt(const RocblasContractionProblem<Ti, To, Tc>& prob)
 #endif
 }
 
+namespace
+{
+    template <typename Ti, typename To, typename Tc>
+    rocblas_internal_gemm_problem
+        make_gemm_problem(const RocblasContractionProblem<Ti, To, Tc>& prob,
+                          rocblas_gemm_algo                            algo,
+                          int32_t                                      solution_index)
+    {
+        rocblas_internal_gemm_problem p;
+        p.trans_a        = prob.trans_a;
+        p.trans_b        = prob.trans_b;
+        p.m              = prob.m;
+        p.n              = prob.n;
+        p.k              = prob.k;
+        p.alpha          = prob.alpha;
+        p.beta           = prob.beta;
+        p.input_type     = rocblas_datatype_from_type<Ti>;
+        p.output_type    = rocblas_datatype_from_type<To>;
+        p.compute_type   = rocblas_datatype_from_type<Tc>;
+        p.lda            = prob.col_stride_a;
+        p.ldb            = prob.col_stride_b;
+        p.ldc            = prob.col_stride_c;
+        p.ldd            = prob.col_stride_d;
+        p.stride_a       = prob.batch_stride_a;
+        p.stride_b       = prob.batch_stride_b;
+        p.stride_c       = prob.batch_stride_c;
+        p.stride_d       = prob.batch_stride_d;
+        p.batch_count    = prob.batch_count;
+        p.pointer_array  = !prob.strided_batch;
+        p.algo           = algo;
+        p.solution_index = solution_index;
+        p.flags          = uint32_t(prob.flags);
+        return p;
+    }
+} // namespace
+
 /******************************************************************************
  * runContractionProblem calls Tensile to run a contraction problem described *
  * by RocblasContractionProblem                                               *
@@ -1578,34 +1614,9 @@ rocblas_status runContractionProblem(const RocblasContractionProblem<Ti, To, Tc>
 
         if(launched && status == rocblas_status_success)
         {
-            rocblas_kernel_select_problem p;
-            p.trans_a        = prob.trans_a;
-            p.trans_b        = prob.trans_b;
-            p.m              = prob.m;
-            p.n              = prob.n;
-            p.k              = prob.k;
-            p.alpha          = prob.alpha;
-            p.beta           = prob.beta;
-            p.input_type     = rocblas_datatype_from_type<Ti>;
-            p.output_type    = rocblas_datatype_from_type<To>;
-            p.compute_type   = rocblas_datatype_from_type<Tc>;
-            p.lda            = prob.col_stride_a;
-            p.ldb            = prob.col_stride_b;
-            p.ldc            = prob.col_stride_c;
-            p.ldd            = prob.col_stride_d;
-            p.stride_a       = prob.batch_stride_a;
-            p.stride_b       = prob.batch_stride_b;
-            p.stride_c       = prob.batch_stride_c;
-            p.stride_d       = prob.batch_stride_d;
-            p.batch_count    = prob.batch_count;
-            p.pointer_array  = !prob.strided_batch;
-            p.algo           = algo;
-            p.solution_index = solution_index;
-            p.flags          = uint32_t(prob.flags);
-
             rocblas_internal_log_kernel_select(
                 prob.handle,
-                p,
+                make_gemm_problem(prob, algo, solution_index),
                 hipblaslt_backend ? "hipblaslt" : "tensile",
                 hipblaslt_backend ? nullptr : selected_tensile_name.c_str(),
                 hipblaslt_attempted && !hipblaslt_backend ? "hipblaslt" : "none");