diff --git a/projects/rocblas/CHANGELOG.md b/projects/rocblas/CHANGELOG.md index 7e8cfd88fc77..c7bc4e139e51 100644 --- a/projects/rocblas/CHANGELOG.md +++ b/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 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 @@ -53,11 +54,32 @@ namespace { if(!strcmp(arg.function, "logging")) testing_logging(arg); + else if constexpr(std::is_same_v) + { + 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 { // 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 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 ... 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..702574f0e9f0 --- /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 +#include +#include +#include +#include + +#ifdef WIN32 +#include +#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 lines; // the kernel-select lines only + }; + + // Runs fn(handle) on a fresh handle with ROCBLAS_LAYER=layer and returns the trace. + template + 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 + struct gemm_data + { + host_matrix hA, hB, hC; + device_matrix 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& 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, hA, (m, m, m)); + HOST_MEMCHECK(host_matrix, hB, (m, n, m)); + DEVICE_MEMCHECK(device_matrix, dA, (m, m, m)); + DEVICE_MEMCHECK(device_matrix, 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 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, hA, (d, d, d, batch_count)); + HOST_MEMCHECK(host_batch_matrix, hB, (d, d, d, batch_count)); + HOST_MEMCHECK(host_batch_matrix, hC, (d, d, d, batch_count)); + DEVICE_MEMCHECK(device_batch_matrix, dA, (d, d, d, batch_count)); + DEVICE_MEMCHECK(device_batch_matrix, dB, (d, d, d, batch_count)); + DEVICE_MEMCHECK(device_batch_matrix, 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, hA, (d, d, d, stride, batch_count)); + HOST_MEMCHECK(host_strided_batch_matrix, hB, (d, d, d, stride, batch_count)); + HOST_MEMCHECK(host_strided_batch_matrix, hC, (d, d, d, stride, batch_count)); + DEVICE_MEMCHECK(device_strided_batch_matrix, dA, (d, d, d, stride, batch_count)); + DEVICE_MEMCHECK(device_strided_batch_matrix, dB, (d, d, d, stride, batch_count)); + DEVICE_MEMCHECK(device_strided_batch_matrix, 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 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 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 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 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 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 207fac0d0c4b..7639bf4573ff 100644 --- a/projects/rocblas/clients/include/testing_logging.hpp +++ b/projects/rocblas/clients/include/testing_logging.hpp @@ -655,6 +655,45 @@ 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. + // 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) + { + 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( + 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_ex ") != std::string::npos + && line.find("# source=") != 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 8b198af5078b..7119397a3c9b 100644 --- a/projects/rocblas/docs/how-to/logging-in-rocblas.rst +++ b/projects/rocblas/docs/how-to/logging-in-rocblas.rst @@ -41,6 +41,33 @@ 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 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_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. + The default stream for logging output is standard error. :ref:`Four environment variables ` 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..6efcca439ccc 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 (Tensile backend only) and the + * source backend (hipblaslt or tensile). */ + 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 175f618ddfea..10848e0b29c9 100644 --- a/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp +++ b/projects/rocblas/library/src/blas3/Tensile/gemm_templates.cpp @@ -34,7 +34,9 @@ #include "check_numerics_matrix.hpp" #include "handle.hpp" +#include "int64_helpers.hpp" #include "logging.hpp" +#include "utility.hpp" /* * =========================================================================== @@ -199,6 +201,35 @@ 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_make_gemm_problem(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) { @@ -223,8 +254,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; @@ -264,6 +298,8 @@ rocblas_status rocblas_internal_gemm(rocblas_handle handle, return status; } + 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..48d82fa25d94 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 +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_internal_gemm_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; + p.output_type = rocblas_datatype_from_type; + p.compute_type = rocblas_datatype_from_type; + 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 ROCBLAS_INTERNAL_EXPORT_NOINLINE rocblas_status 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); + 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_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); + RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle); std::vector alpha_h, beta_h; 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); + 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); + 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); + 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); + 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); + 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); + 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); + 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); + 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); + //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); + //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); + //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); + 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); + 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); + 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); + 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); + 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); + 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); + 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); + 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); + 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); + //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); + //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); + //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); + 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); + 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); + 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); + 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); + 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); + 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 69240ee207cf..a00f4a1a31ff 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); + 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); + 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); + 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); + 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); + // 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); + // 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.hpp b/projects/rocblas/library/src/blas_ex/rocblas_gemm_ex.hpp index e298135b6608..fab62966bd4e 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 +using rocblas_gemm_elem_t + = std::remove_cv_t>>>; + /* * 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_internal_gemm_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>; + p.output_type = rocblas_datatype_from_type>; + p.compute_type = rocblas_datatype_from_type>; + 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>; + rocblas_internal_log_kernel_select( + handle, p, "gemv", "rocblas_internal_gemv_launcher", "none"); + } + return status; } template 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_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_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); + 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); + 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); + RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle); auto layer_mode = handle->layer_mode; 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/handle.cpp b/projects/rocblas/library/src/handle.cpp index 15098042abd4..99a67cbea8e4 100644 --- a/projects/rocblas/library/src/handle.cpp +++ b/projects/rocblas/library/src/handle.cpp @@ -984,7 +984,9 @@ void _rocblas_handle::init_logging() layer_mode = static_cast(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/include/handle.hpp b/projects/rocblas/library/src/include/handle.hpp index 069fb2d51c16..28e888b64ead 100644 --- a/projects/rocblas/library/src/include/handle.hpp +++ b/projects/rocblas/library/src/include/handle.hpp @@ -423,7 +423,18 @@ struct _rocblas_handle std::unique_ptr log_bench_os; std::unique_ptr log_profile_os; void init_logging(); - void init_check_numerics(); + + // 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; + + // 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 data_ptr; @@ -928,6 +939,46 @@ 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 && (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; + m_handle->kernel_select_fallback_from = nullptr; + } + } + ~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; +}; + +// 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 b5df2254396b..0f16b1e81466 100644 --- a/projects/rocblas/library/src/include/logging.hpp +++ b/projects/rocblas/library/src/include/logging.hpp @@ -206,6 +206,46 @@ double rocblas_internal_value_category(const T* beta, rocblas_datatype compute_t 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_internal_gemm_problem +{ + 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_internal_gemm_problem& prob, + const char* source, + const char* kernel, + const char* fallback_from); + /****************************************************************** * ROCBLAS LOGGER * ******************************************************************/ diff --git a/projects/rocblas/library/src/logging.cpp b/projects/rocblas/library/src/logging.cpp index 03a1d824f477..231f09aed0b3 100644 --- a/projects/rocblas/library/src/logging.cpp +++ b/projects/rocblas/library/src/logging.cpp @@ -20,6 +20,7 @@ * * ************************************************************************ */ +#include #include #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_internal_gemm_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::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..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 @@ -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_make_gemm_problem(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 c819091b8594..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 @@ -217,6 +217,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_internal_gemm_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; + p.output_type = rocblas_datatype_from_type; + p.compute_type = rocblas_datatype_from_type; + 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 583bf405275f..d5156940afcd 100644 --- a/projects/rocblas/library/src/tensile_host.cpp +++ b/projects/rocblas/library/src/tensile_host.cpp @@ -758,7 +758,9 @@ namespace rocblas_layer_mode layer_mode = static_cast(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); @@ -1342,6 +1344,42 @@ bool useHipBLASLt(const RocblasContractionProblem& prob) #endif } +namespace +{ + template + rocblas_internal_gemm_problem + make_gemm_problem(const RocblasContractionProblem& 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; + p.output_type = rocblas_datatype_from_type; + p.compute_type = rocblas_datatype_from_type; + 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 * @@ -1354,6 +1392,13 @@ rocblas_status runContractionProblem(const RocblasContractionProblem 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; + std::string selected_tensile_name; + bool solutionBased = algo == rocblas_gemm_algo_solution_index; bool hipBLASLtOnly = false, TensileOnly = false; bool useDefaultSolution = rocblas_default_solution_index(solution_index); @@ -1369,6 +1414,7 @@ rocblas_status runContractionProblem(const RocblasContractionProblem #ifdef BUILD_WITH_HIPBLASLT if(useHipBLASLt(prob)) { + hipblaslt_attempted = true; try { rocblas_int hipblaslt_idx = map_index_rocblas_to_hipblaslt(solution_index); @@ -1439,6 +1485,9 @@ rocblas_status runContractionProblem(const RocblasContractionProblem if(!solution && fallbackTensileProblem(tensile_prob)) solution = library->findBestSolution(tensile_prob, *hardware, fitness_query); + if(solution && kernel_select_logging) + selected_tensile_name = solution->name(); + if(!solution) { if(!useDefaultSolution) @@ -1554,6 +1603,38 @@ rocblas_status runContractionProblem(const RocblasContractionProblem prob.col_stride_d); } + if(kernel_select_logging) + { + // 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_internal_log_kernel_select( + prob.handle, + 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"); + } + + // 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; }