feat(rocblas): Kernel Selection trace logging option

projects/rocblas/CHANGELOG.md

@@ -9,6 +9,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

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");

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
 
 ...