ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-13 02:27:33 +00:00
Code Issues Packages Projects Releases Wiki Activity

clang-format and update copyright year

Browse Source
This commit is contained in:
Anders Smedegaard Pedersen
2025-06-04 14:33:16 +02:00
parent e7de038a28
commit 17033fd3bc
6 changed files with 316 additions and 234 deletions
Download Patch File Download Diff File Expand all files Collapse all files

42
cmake/argparse.cmake
View File

@@ -1,32 +1,26 @@
include(FetchContent)
set(ARGPARSE_DIR "" CACHE STRING "Location of local argparse repo to build against")
set(ARGPARSE_DIR "" CACHE STRING "Location of local argparse repo to build against")
if(ARGPARSE_DIR)
set(FETCHCONTENT_SOURCE_DIR_ARGPARSE ${ARGPARSE_DIR} CACHE STRING "argparse source directory override")
endif()
if(ARGPARSE_DIR) set(FETCHCONTENT_SOURCE_DIR_ARGPARSE ${ARGPARSE_DIR} CACHE STRING
"argparse source directory override") endif()
FetchContent_Declare(
argparse
GIT_REPOSITORY https://github.com/p-ranav/argparse.git
GIT_TAG v2.9 # Using a stable version
)
FetchContent_Declare(argparse GIT_REPOSITORY https
: // github.com/p-ranav/argparse.git
GIT_TAG v2 .9 #Using a stable version)
FetchContent_MakeAvailable(argparse)
FetchContent_MakeAvailable(argparse)
# Create a proper modern CMake INTERFACE target
if(NOT TARGET argparse::argparse)
add_library(argparse::argparse INTERFACE)
target_include_directories(argparse::argparse
INTERFACE
$<BUILD_INTERFACE:${argparse_SOURCE_DIR}/include>
$<INSTALL_INTERFACE:include>
)
#Create a proper modern CMake INTERFACE target
if(NOT TARGET argparse::argparse) add_library(argparse::argparse INTERFACE)
target_include_directories(argparse::argparse INTERFACE
$<BUILD_INTERFACE : ${argparse_SOURCE_DIR} /
include> $<INSTALL_INTERFACE : include>)
# Disable the specific warning for code that includes argparse
target_compile_options(argparse::argparse INTERFACE
$<$<CXX_COMPILER_ID:Clang,AppleClang,GNU>:-Wno-missing-noreturn>
)
endif()
#Disable the specific warning for code that includes argparse
target_compile_options(
argparse::argparse INTERFACE
$<$<CXX_COMPILER_ID : Clang, AppleClang, GNU> : -Wno - missing -
noreturn>) endif()
message(STATUS "Argparse source dir: ${argparse_SOURCE_DIR}")
message(STATUS "Argparse source dir: ${argparse_SOURCE_DIR}")

456
profiler/include/profiler/io_profiler.hpp
View File

@@ -1,5 +1,5 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights
#pragma once
@@ -21,20 +21,23 @@ namespace profiler {
namespace io {
// Enum for output types
enum class OutputType {
CONSOLE, // Standard console output
JSONL // JSON Lines file output
enum class OutputType
{
CONSOLE, // Standard console output
JSONL // JSON Lines file output
};
// Layout enum for type safety
enum class LayoutType {
enum class LayoutType
{
RowMajor,
ColumnMajor,
Unknown
};
// ResultEntry represents a single performance result
struct ResultEntry {
struct ResultEntry
{
std::string op_name;
float time_ms;
float tflops;
@@ -42,58 +45,66 @@ struct ResultEntry {
bool is_valid;
std::unordered_map<std::string, std::string> metadata;
ResultEntry(const std::string& name, float time, float tflops_val, float bandwidth, bool valid = true)
: op_name(name), time_ms(time), tflops(tflops_val), gb_per_sec(bandwidth), is_valid(valid) {}
ResultEntry(
const std::string& name, float time, float tflops_val, float bandwidth, bool valid = true)
: op_name(name), time_ms(time), tflops(tflops_val), gb_per_sec(bandwidth), is_valid(valid)
{
}
ResultEntry() : op_name(""), time_ms(0.0f), tflops(0.0f), gb_per_sec(0.0f), is_valid(false) {}
};
// OutputManager - Manages output formatting and destination
class OutputManager {
private:
class OutputManager
{
private:
static OutputManager* instance;
OutputType current_type;
std::string jsonl_filename;
std::ofstream jsonl_file;
ResultEntry best_result;
OutputManager()
: current_type(OutputType::CONSOLE),
jsonl_filename(""),
best_result() {}
OutputManager() : current_type(OutputType::CONSOLE), jsonl_filename(""), best_result() {}
std::string EscapeJsonString(const std::string& str) {
std::string EscapeJsonString(const std::string& str)
{
std::string escaped;
escaped.reserve(str.size() * 2);
for (char c : str) {
switch (c) {
case '"': escaped += "\\\""; break;
case '\\': escaped += "\\\\"; break;
case '\b': escaped += "\\b"; break;
case '\f': escaped += "\\f"; break;
case '\n': escaped += "\\n"; break;
case '\r': escaped += "\\r"; break;
case '\t': escaped += "\\t"; break;
default:
if (c >= 0 && c < 0x20) {
char hex[7];
std::snprintf(hex, sizeof(hex), "\\u%04x", static_cast<unsigned char>(c));
escaped += hex;
} else {
escaped += c;
}
break;
for(char c : str)
{
switch(c)
{
case '"': escaped += "\\\""; break;
case '\\': escaped += "\\\\"; break;
case '\b': escaped += "\\b"; break;
case '\f': escaped += "\\f"; break;
case '\n': escaped += "\\n"; break;
case '\r': escaped += "\\r"; break;
case '\t': escaped += "\\t"; break;
default:
if(c >= 0 && c < 0x20)
{
char hex[7];
std::snprintf(hex, sizeof(hex), "\\u%04x", static_cast<unsigned char>(c));
escaped += hex;
}
else
{
escaped += c;
}
break;
}
}
return escaped;
}
std::string GetCurrentTimestamp() {
auto now = std::chrono::system_clock::now();
std::string GetCurrentTimestamp()
{
auto now = std::chrono::system_clock::now();
auto time_t = std::chrono::system_clock::to_time_t(now);
auto ms = std::chrono::duration_cast<std::chrono::milliseconds>(
now.time_since_epoch()) % 1000;
auto ms =
std::chrono::duration_cast<std::chrono::milliseconds>(now.time_since_epoch()) % 1000;
std::ostringstream oss;
oss << std::put_time(std::gmtime(&time_t), "%Y-%m-%dT%H:%M:%S");
@@ -101,46 +112,59 @@ private:
return oss.str();
}
bool IsNumeric(const std::string& str) {
if (str.empty()) return false;
bool IsNumeric(const std::string& str)
{
if(str.empty())
return false;
bool has_dot = false;
size_t start = 0;
if (str[0] == '-') {
if (str.length() == 1) return false;
if(str[0] == '-')
{
if(str.length() == 1)
return false;
start = 1;
}
for (size_t i = start; i < str.length(); ++i) {
for(size_t i = start; i < str.length(); ++i)
{
char c = str[i];
if (std::isdigit(c)) {
if(std::isdigit(c))
{
continue;
} else if (c == '.' && !has_dot) {
}
else if(c == '.' && !has_dot)
{
has_dot = true;
} else {
}
else
{
return false;
}
}
return true;
}
void OutputToConsole(const ResultEntry& result, bool is_best) {
if (is_best) {
std::cout << "Best Perf: " << result.time_ms << " ms, "
<< result.tflops << " TFlops, "
<< result.gb_per_sec << " GB/s, "
<< result.op_name << std::endl;
} else {
std::cout << "Perf: " << std::setw(10) << result.time_ms << " ms, "
<< result.tflops << " TFlops, "
<< result.gb_per_sec << " GB/s, "
<< result.op_name << std::endl;
void OutputToConsole(const ResultEntry& result, bool is_best)
{
if(is_best)
{
std::cout << "Best Perf: " << result.time_ms << " ms, " << result.tflops << " TFlops, "
<< result.gb_per_sec << " GB/s, " << result.op_name << std::endl;
}
else
{
std::cout << "Perf: " << std::setw(10) << result.time_ms << " ms, " << result.tflops
<< " TFlops, " << result.gb_per_sec << " GB/s, " << result.op_name
<< std::endl;
}
}
void OutputToJsonl(const ResultEntry& result, bool is_best) {
if (!jsonl_file.is_open()) {
void OutputToJsonl(const ResultEntry& result, bool is_best)
{
if(!jsonl_file.is_open())
{
OutputToConsole(result, is_best);
return;
}
@@ -153,26 +177,36 @@ private:
jsonl_file << "\"is_best\":" << (is_best ? "true" : "false") << ",";
jsonl_file << "\"timestamp\":\"" << GetCurrentTimestamp() << "\"";
if (!result.metadata.empty()) {
for (const auto& [key, value] : result.metadata) {
if(!result.metadata.empty())
{
for(const auto& [key, value] : result.metadata)
{
jsonl_file << ",\"" << EscapeJsonString(key) << "\":";
if (IsNumeric(value)) {
if(IsNumeric(value))
{
jsonl_file << value;
} else {
}
else
{
jsonl_file << "\"" << EscapeJsonString(value) << "\"";
}
}
jsonl_file << ",\"operation_params\":{";
bool first = true;
for (const auto& [key, value] : result.metadata) {
if (!first) jsonl_file << ",";
for(const auto& [key, value] : result.metadata)
{
if(!first)
jsonl_file << ",";
jsonl_file << "\"" << EscapeJsonString(key) << "\":";
if (IsNumeric(value)) {
if(IsNumeric(value))
{
jsonl_file << value;
} else {
}
else
{
jsonl_file << "\"" << EscapeJsonString(value) << "\"";
}
first = false;
@@ -184,26 +218,32 @@ private:
jsonl_file.flush();
}
public:
static OutputManager& GetInstance() {
if (instance == nullptr) {
public:
static OutputManager& GetInstance()
{
if(instance == nullptr)
{
instance = new OutputManager();
}
return *instance;
}
void SetOutputType(OutputType type, const std::string& filename = "") {
void SetOutputType(OutputType type, const std::string& filename = "")
{
current_type = type;
if (jsonl_file.is_open()) {
if(jsonl_file.is_open())
{
jsonl_file.close();
}
if (type == OutputType::JSONL && !filename.empty()) {
if(type == OutputType::JSONL && !filename.empty())
{
jsonl_filename = filename;
jsonl_file.open(filename);
if (!jsonl_file.is_open()) {
if(!jsonl_file.is_open())
{
std::cerr << "Error: Could not open JSONL file: " << filename << std::endl;
std::cerr << "Falling back to console output." << std::endl;
current_type = OutputType::CONSOLE;
@@ -211,44 +251,42 @@ public:
}
}
OutputType GetOutputType() const {
return current_type;
}
OutputType GetOutputType() const { return current_type; }
void OutputResult(const ResultEntry& result) {
if (!best_result.is_valid || result.tflops > best_result.tflops) {
void OutputResult(const ResultEntry& result)
{
if(!best_result.is_valid || result.tflops > best_result.tflops)
{
best_result = result;
}
switch (current_type) {
case OutputType::JSONL:
OutputToJsonl(result, false);
break;
case OutputType::CONSOLE:
default:
OutputToConsole(result, false);
break;
switch(current_type)
{
case OutputType::JSONL: OutputToJsonl(result, false); break;
case OutputType::CONSOLE:
default: OutputToConsole(result, false); break;
}
}
void OutputBestResult() {
if (!best_result.is_valid) {
void OutputBestResult()
{
if(!best_result.is_valid)
{
return;
}
switch (current_type) {
case OutputType::JSONL:
OutputToJsonl(best_result, true);
break;
case OutputType::CONSOLE:
default:
OutputToConsole(best_result, true);
break;
switch(current_type)
{
case OutputType::JSONL: OutputToJsonl(best_result, true); break;
case OutputType::CONSOLE:
default: OutputToConsole(best_result, true); break;
}
}
~OutputManager() {
if (jsonl_file.is_open()) {
~OutputManager()
{
if(jsonl_file.is_open())
{
jsonl_file.close();
}
}
@@ -258,50 +296,69 @@ public:
namespace utils {
// Convert DataTypeEnum to string
inline std::string DataTypeToString(DataTypeEnum data_type) {
switch (data_type) {
case DataTypeEnum::Half: return "f16";
case DataTypeEnum::Float: return "f32";
case DataTypeEnum::Int32: return "i32";
case DataTypeEnum::Int8: return "i8";
case DataTypeEnum::Int8x4: return "i8x4";
case DataTypeEnum::BFloat16: return "bf16";
case DataTypeEnum::Double: return "f64";
case DataTypeEnum::Float8: return "fp8";
case DataTypeEnum::Unknown:
default: return "unknown";
inline std::string DataTypeToString(DataTypeEnum data_type)
{
switch(data_type)
{
case DataTypeEnum::Half: return "f16";
case DataTypeEnum::Float: return "f32";
case DataTypeEnum::Int32: return "i32";
case DataTypeEnum::Int8: return "i8";
case DataTypeEnum::Int8x4: return "i8x4";
case DataTypeEnum::BFloat16: return "bf16";
case DataTypeEnum::Double: return "f64";
case DataTypeEnum::Float8: return "fp8";
case DataTypeEnum::Unknown:
default: return "unknown";
}
}
// Convert LayoutType to string
inline std::string LayoutToString(LayoutType layout) {
switch (layout) {
case LayoutType::RowMajor: return "RowMajor";
case LayoutType::ColumnMajor: return "ColumnMajor";
case LayoutType::Unknown:
default: return "unknown";
inline std::string LayoutToString(LayoutType layout)
{
switch(layout)
{
case LayoutType::RowMajor: return "RowMajor";
case LayoutType::ColumnMajor: return "ColumnMajor";
case LayoutType::Unknown:
default: return "unknown";
}
}
// Template function to automatically detect data type from C++ type
template<typename DataType>
DataTypeEnum GetDataTypeEnum() {
if constexpr(std::is_same_v<DataType, float>) {
template <typename DataType>
DataTypeEnum GetDataTypeEnum()
{
if constexpr(std::is_same_v<DataType, float>)
{
return DataTypeEnum::Float;
} else if constexpr(std::is_same_v<DataType, double>) {
}
else if constexpr(std::is_same_v<DataType, double>)
{
return DataTypeEnum::Double;
} else if constexpr(std::is_same_v<DataType, int8_t>) {
}
else if constexpr(std::is_same_v<DataType, int8_t>)
{
return DataTypeEnum::Int8;
} else if constexpr(std::is_same_v<DataType, int32_t>) {
}
else if constexpr(std::is_same_v<DataType, int32_t>)
{
return DataTypeEnum::Int32;
} else {
}
else
{
// Use typeid to try to identify CK-specific types
std::string type_name = typeid(DataType).name();
if (type_name.find("half") != std::string::npos) {
if(type_name.find("half") != std::string::npos)
{
return DataTypeEnum::Half;
} else if (type_name.find("bhalf") != std::string::npos) {
}
else if(type_name.find("bhalf") != std::string::npos)
{
return DataTypeEnum::BFloat16;
} else if (sizeof(DataType) == 1 && !std::is_same_v<DataType, int8_t>) {
}
else if(sizeof(DataType) == 1 && !std::is_same_v<DataType, int8_t>)
{
return DataTypeEnum::Float8;
}
return DataTypeEnum::Unknown;
@@ -309,12 +366,16 @@ DataTypeEnum GetDataTypeEnum() {
}
// Template function to automatically detect layout type
template<typename Layout>
LayoutType GetLayoutType() {
template <typename Layout>
LayoutType GetLayoutType()
{
std::string type_name = typeid(Layout).name();
if (type_name.find("RowMajor") != std::string::npos) {
if(type_name.find("RowMajor") != std::string::npos)
{
return LayoutType::RowMajor;
} else if (type_name.find("ColumnMajor") != std::string::npos) {
}
else if(type_name.find("ColumnMajor") != std::string::npos)
{
return LayoutType::ColumnMajor;
}
return LayoutType::Unknown;
@@ -323,9 +384,14 @@ LayoutType GetLayoutType() {
} // namespace utils
// Template functions for automatic metadata creation - ALL LOGIC STAYS HERE
template<typename ALayout, typename BLayout, typename CLayout,
typename ADataType, typename BDataType, typename CDataType>
inline std::unordered_map<std::string, std::string> CreateGemmMetadata(int M, int N, int K) {
template <typename ALayout,
typename BLayout,
typename CLayout,
typename ADataType,
typename BDataType,
typename CDataType>
inline std::unordered_map<std::string, std::string> CreateGemmMetadata(int M, int N, int K)
{
std::unordered_map<std::string, std::string> metadata;
// Automatically detect types using template introspection
@@ -337,62 +403,63 @@ inline std::unordered_map<std::string, std::string> CreateGemmMetadata(int M, in
auto b_layout = utils::GetLayoutType<BLayout>();
auto c_layout = utils::GetLayoutType<CLayout>();
metadata["operation_type"] = "gemm";
metadata["datatype"] = utils::DataTypeToString(c_datatype);
metadata["input_datatype"] = utils::DataTypeToString(a_datatype);
metadata["operation_type"] = "gemm";
metadata["datatype"] = utils::DataTypeToString(c_datatype);
metadata["input_datatype"] = utils::DataTypeToString(a_datatype);
metadata["weight_datatype"] = utils::DataTypeToString(b_datatype);
metadata["output_datatype"] = utils::DataTypeToString(c_datatype);
metadata["layout_a"] = utils::LayoutToString(a_layout);
metadata["layout_b"] = utils::LayoutToString(b_layout);
metadata["layout_c"] = utils::LayoutToString(c_layout);
metadata["M"] = std::to_string(M);
metadata["N"] = std::to_string(N);
metadata["K"] = std::to_string(K);
metadata["layout_a"] = utils::LayoutToString(a_layout);
metadata["layout_b"] = utils::LayoutToString(b_layout);
metadata["layout_c"] = utils::LayoutToString(c_layout);
metadata["M"] = std::to_string(M);
metadata["N"] = std::to_string(N);
metadata["K"] = std::to_string(K);
return metadata;
}
// Template function for convolution metadata
template<typename InputDataType, typename WeightDataType, typename OutputDataType>
template <typename InputDataType, typename WeightDataType, typename OutputDataType>
inline std::unordered_map<std::string, std::string> CreateConvMetadata(
int N, int C, int H, int W, int K, int Y, int X,
const std::string& conv_type = "conv2d") {
int N, int C, int H, int W, int K, int Y, int X, const std::string& conv_type = "conv2d")
{
std::unordered_map<std::string, std::string> metadata;
auto input_dtype = utils::GetDataTypeEnum<InputDataType>();
auto input_dtype = utils::GetDataTypeEnum<InputDataType>();
auto weight_dtype = utils::GetDataTypeEnum<WeightDataType>();
auto output_dtype = utils::GetDataTypeEnum<OutputDataType>();
metadata["operation_type"] = conv_type;
metadata["input_datatype"] = utils::DataTypeToString(input_dtype);
metadata["operation_type"] = conv_type;
metadata["input_datatype"] = utils::DataTypeToString(input_dtype);
metadata["weight_datatype"] = utils::DataTypeToString(weight_dtype);
metadata["output_datatype"] = utils::DataTypeToString(output_dtype);
metadata["N"] = std::to_string(N);
metadata["C"] = std::to_string(C);
metadata["H"] = std::to_string(H);
metadata["W"] = std::to_string(W);
metadata["K"] = std::to_string(K);
metadata["Y"] = std::to_string(Y);
metadata["X"] = std::to_string(X);
metadata["N"] = std::to_string(N);
metadata["C"] = std::to_string(C);
metadata["H"] = std::to_string(H);
metadata["W"] = std::to_string(W);
metadata["K"] = std::to_string(K);
metadata["Y"] = std::to_string(Y);
metadata["X"] = std::to_string(X);
return metadata;
}
// Template function for normalization metadata
template<typename DataType>
inline std::unordered_map<std::string, std::string> CreateNormMetadata(
const std::vector<int>& dimensions,
const std::string& norm_type) {
template <typename DataType>
inline std::unordered_map<std::string, std::string>
CreateNormMetadata(const std::vector<int>& dimensions, const std::string& norm_type)
{
std::unordered_map<std::string, std::string> metadata;
auto dtype = utils::GetDataTypeEnum<DataType>();
metadata["operation_type"] = norm_type;
metadata["datatype"] = utils::DataTypeToString(dtype);
metadata["datatype"] = utils::DataTypeToString(dtype);
for (size_t i = 0; i < dimensions.size(); ++i) {
for(size_t i = 0; i < dimensions.size(); ++i)
{
metadata["dim" + std::to_string(i)] = std::to_string(dimensions[i]);
}
@@ -400,14 +467,16 @@ inline std::unordered_map<std::string, std::string> CreateNormMetadata(
}
// Generic operation metadata
inline std::unordered_map<std::string, std::string> CreateOperationMetadata(
const std::string& operation_type,
const std::unordered_map<std::string, std::string>& additional_fields = {}) {
inline std::unordered_map<std::string, std::string>
CreateOperationMetadata(const std::string& operation_type,
const std::unordered_map<std::string, std::string>& additional_fields = {})
{
std::unordered_map<std::string, std::string> metadata;
metadata["operation_type"] = operation_type;
for (const auto& [key, value] : additional_fields) {
for(const auto& [key, value] : additional_fields)
{
metadata[key] = value;
}
@@ -415,60 +484,73 @@ inline std::unordered_map<std::string, std::string> CreateOperationMetadata(
}
// Main reporting functions - kernel implementations only call these
inline void SetOutputType(OutputType type, const std::string& filename = "") {
inline void SetOutputType(OutputType type, const std::string& filename = "")
{
OutputManager::GetInstance().SetOutputType(type, filename);
}
inline OutputType GetOutputType() {
return OutputManager::GetInstance().GetOutputType();
}
inline OutputType GetOutputType() { return OutputManager::GetInstance().GetOutputType(); }
inline void ReportResult(const std::string& op_name, float time_ms, float tflops, float gb_per_sec,
const std::unordered_map<std::string, std::string>& metadata) {
inline void ReportResult(const std::string& op_name,
float time_ms,
float tflops,
float gb_per_sec,
const std::unordered_map<std::string, std::string>& metadata)
{
ResultEntry result(op_name, time_ms, tflops, gb_per_sec);
result.metadata = metadata;
OutputManager::GetInstance().OutputResult(result);
}
inline void ReportResult(const std::string& op_name, float time_ms, float tflops, float gb_per_sec) {
inline void ReportResult(const std::string& op_name, float time_ms, float tflops, float gb_per_sec)
{
std::unordered_map<std::string, std::string> empty_metadata;
ReportResult(op_name, time_ms, tflops, gb_per_sec, empty_metadata);
}
inline void ReportBestResult() {
OutputManager::GetInstance().OutputBestResult();
}
inline void ReportBestResult() { OutputManager::GetInstance().OutputBestResult(); }
inline void AddOutputArguments(argparse::ArgumentParser& parser) {
inline void AddOutputArguments(argparse::ArgumentParser& parser)
{
parser.add_argument("-o", "--output")
.help("Output format (console, jsonl=<filename>)")
.default_value(std::string("console"));
}
inline void ParseOutputArguments(const argparse::ArgumentParser& parser) {
try {
inline void ParseOutputArguments(const argparse::ArgumentParser& parser)
{
try
{
std::string output_format = parser.get<std::string>("--output");
if(output_format == "console") {
if(output_format == "console")
{
SetOutputType(OutputType::CONSOLE);
}
else if(output_format.compare(0, 6, "jsonl=") == 0) {
else if(output_format.compare(0, 6, "jsonl=") == 0)
{
std::string filename = output_format.substr(6);
if(!filename.empty()) {
if(!filename.empty())
{
SetOutputType(OutputType::JSONL, filename);
}
else {
std::cerr << "Error: JSONL output requires a filename (use --output=jsonl=filename.jsonl)" << std::endl;
else
{
std::cerr
<< "Error: JSONL output requires a filename (use --output=jsonl=filename.jsonl)"
<< std::endl;
SetOutputType(OutputType::CONSOLE);
}
}
else {
std::cerr << "Warning: Unknown output format '" << output_format
<< "', using console" << std::endl;
else
{
std::cerr << "Warning: Unknown output format '" << output_format << "', using console"
<< std::endl;
SetOutputType(OutputType::CONSOLE);
}
}
catch(const std::exception& err) {
catch(const std::exception& err)
{
std::cerr << "Error parsing output arguments: " << err.what() << std::endl;
SetOutputType(OutputType::CONSOLE);
}

2
profiler/include/profiler/profile_batched_gemm_impl.hpp
View File

@@ -1,5 +1,5 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once

14
profiler/include/profiler/profile_gemm_impl.hpp
View File

@@ -1,5 +1,5 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
@@ -24,7 +24,6 @@
#include "ck/library/utility/fill.hpp"
#include "profiler/io_profiler.hpp"
namespace ck {
namespace profiler {
@@ -181,10 +180,11 @@ int profile_gemm_impl(int do_verification,
float gb_per_sec = num_btype / 1.E6 / avg_time;
auto metadata = ck::profiler::io::CreateGemmMetadata<ALayout, BLayout, CLayout, ADataType, BDataType, CDataType>(M, N, K);
auto metadata = ck::profiler::io::
CreateGemmMetadata<ALayout, BLayout, CLayout, ADataType, BDataType, CDataType>(
M, N, K);
ck::profiler::io::ReportResult(op_name, avg_time, tflops, gb_per_sec, metadata);
if(tflops > best_tflops)
{
best_instance_id = instance_id;
@@ -254,11 +254,11 @@ int profile_gemm_impl(int do_verification,
float gb_per_sec = num_btype / 1.E6 / avg_time;
// Report the more accurate best result with metadata
auto metadata = ck::profiler::io::CreateGemmMetadata<ALayout, BLayout, CLayout, ADataType, BDataType, CDataType>(M, N, K);
auto metadata = ck::profiler::io::
CreateGemmMetadata<ALayout, BLayout, CLayout, ADataType, BDataType, CDataType>(
M, N, K);
ck::profiler::io::ReportResult(op_name, avg_time, tflops, gb_per_sec, metadata);
ck::profiler::io::ReportBestResult();
}
}

4
profiler/src/io_profiler.cpp
View File

@@ -1,12 +1,12 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include "profiler/io_profiler.hpp"
namespace ck {
namespace profiler {
namespace io {
OutputManager* OutputManager::instance = nullptr;
OutputManager* OutputManager::instance = nullptr;
} // namespace io
} // namespace profiler
} // namespace ck

32
profiler/src/profiler.cpp
View File

@@ -1,5 +1,5 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <cstring>
@@ -56,21 +56,25 @@ std::tuple<std::string, int, char**> parse_global_arguments(int argc, char* argv
args_to_keep.push_back(0); // Always keep program name
// If we have an operation name, keep it
if(!operation_name.empty() && argc > 1) {
if(!operation_name.empty() && argc > 1)
{
args_to_keep.push_back(1); // Keep the operation name
}
// Check remaining arguments
for(int i = 1; i < argc; i++) {
for(int i = 1; i < argc; i++)
{
std::string arg = argv[i];
// Skip the operation name (already handled)
if(i == 1 && !operation_name.empty()) {
if(i == 1 && !operation_name.empty())
{
continue;
}
// Skip output arguments and their values
if(arg == "-o" || arg == "--output") {
if(arg == "-o" || arg == "--output")
{
i++; // Skip the value too
continue;
}
@@ -80,21 +84,21 @@ std::tuple<std::string, int, char**> parse_global_arguments(int argc, char* argv
}
// Create new argv array
int new_argc = args_to_keep.size();
int new_argc = args_to_keep.size();
char** new_argv = new char*[new_argc];
for(int i = 0; i < new_argc; i++) {
for(int i = 0; i < new_argc; i++)
{
// Make a copy of each string we want to keep
int original_index = args_to_keep[i];
size_t len = std::strlen(argv[original_index]) + 1; // +1 for null terminator
new_argv[i] = new char[len];
size_t len = std::strlen(argv[original_index]) + 1; // +1 for null terminator
new_argv[i] = new char[len];
std::strcpy(new_argv[i], argv[original_index]);
}
return {operation_name, new_argc, new_argv};
}
int main(int argc, char* argv[])
{
// Check if no arguments were provided
@@ -112,8 +116,10 @@ int main(int argc, char* argv[])
// Helper function to clean up memory
auto cleanup = [new_argv, new_argc]() {
if(new_argv) {
for(int i = 0; i < new_argc; i++) {
if(new_argv)
{
for(int i = 0; i < new_argc; i++)
{
delete[] new_argv[i];
}
delete[] new_argv;
@@ -130,7 +136,7 @@ int main(int argc, char* argv[])
// Look up the operation and run it if found
if(const auto operation = ProfilerOperationRegistry::GetInstance().Get(operation_name);
operation.has_value())
operation.has_value())
{
int result = (*operation)(new_argc, new_argv);
cleanup();