22949872003

Committed 11 Mar 2026 11:15AM UTC coverage: 63.681% (-0.9%) from 64.6%

Build # 22949872003

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #569 from daisytuner/HIPtarget

ROCmTarget

Coverage Stats

191 of 803 new or added lines in 15 files covered. (23.79%)

3 existing lines in 2 files now uncovered.

24700 of 38787 relevant lines covered (63.68%)

370.4 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

0.0

/opt/src/targets/rocm/rocm_data_offloading_node.cpp

#include "sdfg/targets/rocm/rocm_data_offloading_node.h"

#include <cstddef>
#include <memory>
#include <nlohmann/json_fwd.hpp>

#include "sdfg/analysis/loop_analysis.h"
#include "sdfg/codegen/code_snippet_factory.h"
#include "sdfg/codegen/dispatchers/block_dispatcher.h"
#include "sdfg/codegen/instrumentation/instrumentation_info.h"
#include "sdfg/codegen/language_extension.h"
#include "sdfg/codegen/utils.h"
#include "sdfg/data_flow/data_flow_graph.h"
#include "sdfg/data_flow/data_flow_node.h"
#include "sdfg/data_flow/library_node.h"
#include "sdfg/function.h"
#include "sdfg/graph/graph.h"
#include "sdfg/symbolic/symbolic.h"
#include "sdfg/targets/offloading/data_offloading_node.h"
#include "sdfg/targets/rocm/rocm.h"
#include "symengine/symengine_rcp.h"

namespace sdfg {
namespace rocm {

ROCMDataOffloadingNode::ROCMDataOffloadingNode(
    size_t element_id,
    const DebugInfo& debug_info,
    const graph::Vertex vertex,
    data_flow::DataFlowGraph& parent,
    symbolic::Expression size,
    symbolic::Expression device_id,
    offloading::DataTransferDirection transfer_direction,
    offloading::BufferLifecycle buffer_lifecycle
)
    : offloading::DataOffloadingNode(
          element_id,
          debug_info,
          vertex,
          parent,
          LibraryNodeType_ROCM_Offloading,
          {},
          {},
          transfer_direction,
          buffer_lifecycle,
          size
      ),
      device_id_(device_id) {
    if (!is_NONE(transfer_direction)) {
        this->inputs_.push_back("_src");
        this->outputs_.push_back("_dst");
    } else if (is_ALLOC(buffer_lifecycle)) {
        this->outputs_.push_back("_ret");
    } else if (is_FREE(buffer_lifecycle)) {
        this->inputs_.push_back("_ptr");
        this->outputs_.push_back("_ptr");
    }
}

void ROCMDataOffloadingNode::validate(const Function& function) const {
    // Prevent copy-in and free
    if (this->is_h2d() && this->is_free()) {
        throw InvalidSDFGException("ROCMDataOffloadingNode: Combination copy-in and free is not allowed");
    }

    // Prevent copy-out and alloc
    if (this->is_d2h() && this->is_alloc()) {
        throw InvalidSDFGException("ROCMDataOffloadingNode: Combination copy-out and alloc is not allowed");
    }
}

const symbolic::Expression ROCMDataOffloadingNode::device_id() const { return this->device_id_; }

std::unique_ptr<data_flow::DataFlowNode> ROCMDataOffloadingNode::
    clone(size_t element_id, const graph::Vertex vertex, data_flow::DataFlowGraph& parent) const {
    return std::make_unique<ROCMDataOffloadingNode>(
        element_id,
        this->debug_info(),
        vertex,
        parent,
        this->size(),
        this->device_id(),
        this->transfer_direction(),
        this->buffer_lifecycle()
    );
}

symbolic::SymbolSet ROCMDataOffloadingNode::symbols() const {
    if (this->device_id().is_null()) {
        return offloading::DataOffloadingNode::symbols();
    }
    auto symbols = offloading::DataOffloadingNode::symbols();
    auto device_id_atoms = symbolic::atoms(this->device_id());
    symbols.insert(device_id_atoms.begin(), device_id_atoms.end());
    return symbols;
}

void ROCMDataOffloadingNode::replace(const symbolic::Expression old_expression, const symbolic::Expression new_expression) {
    offloading::DataOffloadingNode::replace(old_expression, new_expression);
    this->device_id_ = symbolic::subs(this->device_id_, old_expression, new_expression);
}

bool ROCMDataOffloadingNode::blocking() const { return true; }

bool ROCMDataOffloadingNode::redundant_with(const offloading::DataOffloadingNode& other) const {
    if (!offloading::DataOffloadingNode::redundant_with(other)) {
        return false;
    }

    auto& other_node = static_cast<const ROCMDataOffloadingNode&>(other);
    if (!symbolic::null_safe_eq(this->device_id(), other_node.device_id())) {
        return false;
    }

    return true;
}

bool ROCMDataOffloadingNode::equal_with(const offloading::DataOffloadingNode& other) const {
    if (!offloading::DataOffloadingNode::equal_with(other)) {
        return false;
    }

    auto& other_node = static_cast<const ROCMDataOffloadingNode&>(other);
    if (!symbolic::null_safe_eq(this->device_id(), other_node.device_id())) {
        return false;
    }

    return true;
}

ROCMDataOffloadingNodeDispatcher::ROCMDataOffloadingNodeDispatcher(
    codegen::LanguageExtension& language_extension,
    const Function& function,
    const data_flow::DataFlowGraph& data_flow_graph,
    const data_flow::LibraryNode& node
)
    : codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}

void ROCMDataOffloadingNodeDispatcher::dispatch_code(
    codegen::PrettyPrinter& stream,
    codegen::PrettyPrinter& globals_stream,
    codegen::CodeSnippetFactory& library_snippet_factory
) {
    auto& offloading_node = static_cast<const ROCMDataOffloadingNode&>(this->node_);

    stream << "hipError_t err;" << std::endl;

    if (offloading_node.is_alloc()) {
        stream << "err = hipMalloc(&" << offloading_node.output(0) << ", "
               << this->language_extension_.expression(offloading_node.size()) << ");" << std::endl;
        rocm_error_checking(stream, this->language_extension_, "err");
    }

    if (offloading_node.is_h2d()) {
        stream << "err = hipMemcpy(" << offloading_node.output(0) << ", " << offloading_node.input(0) << ", "
               << this->language_extension_.expression(offloading_node.size()) << ", hipMemcpyHostToDevice);"
               << std::endl;
        rocm_error_checking(stream, this->language_extension_, "err");
    } else if (offloading_node.is_d2h()) {
        stream << "err = hipMemcpy(" << offloading_node.output(0) << ", " << offloading_node.input(0) << ", "
               << this->language_extension_.expression(offloading_node.size()) << ", hipMemcpyDeviceToHost);"
               << std::endl;
        rocm_error_checking(stream, this->language_extension_, "err");
    }

    if (offloading_node.is_free()) {
        stream << "err = hipFree(" << offloading_node.input(0) << ");" << std::endl;
        rocm_error_checking(stream, this->language_extension_, "err");
    }
}

codegen::InstrumentationInfo ROCMDataOffloadingNodeDispatcher::instrumentation_info() const {
    auto& rocm_node = static_cast<const ROCMDataOffloadingNode&>(node_);
    if (rocm_node.is_d2h()) {
        return codegen::InstrumentationInfo(
            node_.element_id(),
            codegen::ElementType_D2HTransfer,
            TargetType_ROCM,
            analysis::LoopInfo{},
            {{"pcie_bytes", language_extension_.expression(rocm_node.size())}}
        );
    } else if (rocm_node.is_h2d()) {
        return codegen::InstrumentationInfo(
            node_.element_id(),
            codegen::ElementType_H2DTransfer,
            TargetType_ROCM,
            analysis::LoopInfo{},
            {{"pcie_bytes", language_extension_.expression(rocm_node.size())}}
        );
    } else {
        return codegen::LibraryNodeDispatcher::instrumentation_info();
    }
}

nlohmann::json ROCMDataOffloadingNodeSerializer::serialize(const sdfg::data_flow::LibraryNode& library_node) {
    const auto& node = static_cast<const ROCMDataOffloadingNode&>(library_node);
    nlohmann::json j;

    // Library node
    j["type"] = "library_node";
    j["element_id"] = library_node.element_id();

    // Debug info
    auto& debug_info = library_node.debug_info();
    j["has"] = debug_info.has();
    j["filename"] = debug_info.filename();
    j["start_line"] = debug_info.start_line();
    j["start_column"] = debug_info.start_column();
    j["end_line"] = debug_info.end_line();
    j["end_column"] = debug_info.end_column();

    // Library node properties
    j["code"] = std::string(library_node.code().value());

    // Offloading node properties
    sdfg::serializer::JSONSerializer serializer;
    if (node.size().is_null()) {
        j["size"] = nlohmann::json::value_t::null;
    } else {
        j["size"] = serializer.expression(node.size());
    }
    j["device_id"] = serializer.expression(node.device_id());
    j["transfer_direction"] = static_cast<int8_t>(node.transfer_direction());
    j["buffer_lifecycle"] = static_cast<int8_t>(node.buffer_lifecycle());

    return j;
}

data_flow::LibraryNode& ROCMDataOffloadingNodeSerializer::deserialize(
    const nlohmann::json& j, sdfg::builder::StructuredSDFGBuilder& builder, structured_control_flow::Block& parent
) {
    auto code = j["code"].get<std::string>();
    if (code != LibraryNodeType_ROCM_Offloading.value()) {
        throw std::runtime_error("Invalid library node code");
    }

    sdfg::serializer::JSONSerializer serializer;
    DebugInfo debug_info = serializer.json_to_debug_info(j["debug_info"]);

    symbolic::Expression size;
    if (!j.contains("size") || j.at("size").is_null()) {
        size = SymEngine::null;
    } else {
        size = symbolic::parse(j.at("size"));
    }
    SymEngine::Expression device_id(j.at("device_id"));
    auto transfer_direction = static_cast<offloading::DataTransferDirection>(j["transfer_direction"].get<int8_t>());
    auto buffer_lifecycle = static_cast<offloading::BufferLifecycle>(j["buffer_lifecycle"].get<int8_t>());

    return builder.add_library_node<
        ROCMDataOffloadingNode>(parent, debug_info, size, device_id, transfer_direction, buffer_lifecycle);
}

} // namespace rocm
} // namespace sdfg

1	#include "sdfg/targets/rocm/rocm_data_offloading_node.h"
2
3	#include <cstddef>
4	#include <memory>
5	#include <nlohmann/json_fwd.hpp>
6
7	#include "sdfg/analysis/loop_analysis.h"
8	#include "sdfg/codegen/code_snippet_factory.h"
9	#include "sdfg/codegen/dispatchers/block_dispatcher.h"
10	#include "sdfg/codegen/instrumentation/instrumentation_info.h"
11	#include "sdfg/codegen/language_extension.h"
12	#include "sdfg/codegen/utils.h"
13	#include "sdfg/data_flow/data_flow_graph.h"
14	#include "sdfg/data_flow/data_flow_node.h"
15	#include "sdfg/data_flow/library_node.h"
16	#include "sdfg/function.h"
17	#include "sdfg/graph/graph.h"
18	#include "sdfg/symbolic/symbolic.h"
19	#include "sdfg/targets/offloading/data_offloading_node.h"
20	#include "sdfg/targets/rocm/rocm.h"
21	#include "symengine/symengine_rcp.h"
22
23	namespace sdfg {
24	namespace rocm {
25
26	ROCMDataOffloadingNode::ROCMDataOffloadingNode(
27	size_t element_id,
28	const DebugInfo& debug_info,
29	const graph::Vertex vertex,
30	data_flow::DataFlowGraph& parent,
31	symbolic::Expression size,
32	symbolic::Expression device_id,
33	offloading::DataTransferDirection transfer_direction,
34	offloading::BufferLifecycle buffer_lifecycle
35	)
NEW 36	: offloading::DataOffloadingNode(	×
NEW 37	element_id,	×
NEW 38	debug_info,	×
NEW 39	vertex,	×
NEW 40	parent,	×
NEW 41	LibraryNodeType_ROCM_Offloading,	×
NEW 42	{},	×
NEW 43	{},	×
NEW 44	transfer_direction,	×
NEW 45	buffer_lifecycle,	×
NEW 46	size	×
NEW 47	),	×
NEW 48	device_id_(device_id) {	×
NEW 49	if (!is_NONE(transfer_direction)) {	×
NEW 50	this->inputs_.push_back("_src");	×
NEW 51	this->outputs_.push_back("_dst");	×
NEW 52	} else if (is_ALLOC(buffer_lifecycle)) {	×
NEW 53	this->outputs_.push_back("_ret");	×
NEW 54	} else if (is_FREE(buffer_lifecycle)) {	×
NEW 55	this->inputs_.push_back("_ptr");	×
NEW 56	this->outputs_.push_back("_ptr");	×
NEW 57	}	×
NEW 58	}	×
59
NEW 60	void ROCMDataOffloadingNode::validate(const Function& function) const {	×
61	// Prevent copy-in and free
NEW 62	if (this->is_h2d() && this->is_free()) {	×
NEW 63	throw InvalidSDFGException("ROCMDataOffloadingNode: Combination copy-in and free is not allowed");	×
NEW 64	}	×
65
66	// Prevent copy-out and alloc
NEW 67	if (this->is_d2h() && this->is_alloc()) {	×
NEW 68	throw InvalidSDFGException("ROCMDataOffloadingNode: Combination copy-out and alloc is not allowed");	×
NEW 69	}	×
NEW 70	}	×
71
NEW 72	const symbolic::Expression ROCMDataOffloadingNode::device_id() const { return this->device_id_; }	×
73
74	std::unique_ptr<data_flow::DataFlowNode> ROCMDataOffloadingNode::
NEW 75	clone(size_t element_id, const graph::Vertex vertex, data_flow::DataFlowGraph& parent) const {	×
NEW 76	return std::make_unique<ROCMDataOffloadingNode>(	×
NEW 77	element_id,	×
NEW 78	this->debug_info(),	×
NEW 79	vertex,	×
NEW 80	parent,	×
NEW 81	this->size(),	×
NEW 82	this->device_id(),	×
NEW 83	this->transfer_direction(),	×
NEW 84	this->buffer_lifecycle()	×
NEW 85	);	×
NEW 86	}	×
87
NEW 88	symbolic::SymbolSet ROCMDataOffloadingNode::symbols() const {	×
NEW 89	if (this->device_id().is_null()) {	×
NEW 90	return offloading::DataOffloadingNode::symbols();	×
NEW 91	}	×
NEW 92	auto symbols = offloading::DataOffloadingNode::symbols();	×
NEW 93	auto device_id_atoms = symbolic::atoms(this->device_id());	×
NEW 94	symbols.insert(device_id_atoms.begin(), device_id_atoms.end());	×
NEW 95	return symbols;	×
NEW 96	}	×
97
NEW 98	void ROCMDataOffloadingNode::replace(const symbolic::Expression old_expression, const symbolic::Expression new_expression) {	×
NEW 99	offloading::DataOffloadingNode::replace(old_expression, new_expression);	×
NEW 100	this->device_id_ = symbolic::subs(this->device_id_, old_expression, new_expression);	×
NEW 101	}	×
102
NEW 103	bool ROCMDataOffloadingNode::blocking() const { return true; }	×
104
NEW 105	bool ROCMDataOffloadingNode::redundant_with(const offloading::DataOffloadingNode& other) const {	×
NEW 106	if (!offloading::DataOffloadingNode::redundant_with(other)) {	×
NEW 107	return false;	×
NEW 108	}	×
109
NEW 110	auto& other_node = static_cast<const ROCMDataOffloadingNode&>(other);	×
NEW 111	if (!symbolic::null_safe_eq(this->device_id(), other_node.device_id())) {	×
NEW 112	return false;	×
NEW 113	}	×
114
NEW 115	return true;	×
NEW 116	}	×
117
NEW 118	bool ROCMDataOffloadingNode::equal_with(const offloading::DataOffloadingNode& other) const {	×
NEW 119	if (!offloading::DataOffloadingNode::equal_with(other)) {	×
NEW 120	return false;	×
NEW 121	}	×
122
NEW 123	auto& other_node = static_cast<const ROCMDataOffloadingNode&>(other);	×
NEW 124	if (!symbolic::null_safe_eq(this->device_id(), other_node.device_id())) {	×
NEW 125	return false;	×
NEW 126	}	×
127
NEW 128	return true;	×
NEW 129	}	×
130
131	ROCMDataOffloadingNodeDispatcher::ROCMDataOffloadingNodeDispatcher(
132	codegen::LanguageExtension& language_extension,
133	const Function& function,
134	const data_flow::DataFlowGraph& data_flow_graph,
135	const data_flow::LibraryNode& node
136	)
NEW 137	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
138
139	void ROCMDataOffloadingNodeDispatcher::dispatch_code(
140	codegen::PrettyPrinter& stream,
141	codegen::PrettyPrinter& globals_stream,
142	codegen::CodeSnippetFactory& library_snippet_factory
NEW 143	) {	×
NEW 144	auto& offloading_node = static_cast<const ROCMDataOffloadingNode&>(this->node_);	×
145
NEW 146	stream << "hipError_t err;" << std::endl;	×
147
NEW 148	if (offloading_node.is_alloc()) {	×
NEW 149	stream << "err = hipMalloc(&" << offloading_node.output(0) << ", "	×
NEW 150	<< this->language_extension_.expression(offloading_node.size()) << ");" << std::endl;	×
NEW 151	rocm_error_checking(stream, this->language_extension_, "err");	×
NEW 152	}	×
153
NEW 154	if (offloading_node.is_h2d()) {	×
NEW 155	stream << "err = hipMemcpy(" << offloading_node.output(0) << ", " << offloading_node.input(0) << ", "	×
NEW 156	<< this->language_extension_.expression(offloading_node.size()) << ", hipMemcpyHostToDevice);"	×
NEW 157	<< std::endl;	×
NEW 158	rocm_error_checking(stream, this->language_extension_, "err");	×
NEW 159	} else if (offloading_node.is_d2h()) {	×
NEW 160	stream << "err = hipMemcpy(" << offloading_node.output(0) << ", " << offloading_node.input(0) << ", "	×
NEW 161	<< this->language_extension_.expression(offloading_node.size()) << ", hipMemcpyDeviceToHost);"	×
NEW 162	<< std::endl;	×
NEW 163	rocm_error_checking(stream, this->language_extension_, "err");	×
NEW 164	}	×
165
NEW 166	if (offloading_node.is_free()) {	×
NEW 167	stream << "err = hipFree(" << offloading_node.input(0) << ");" << std::endl;	×
NEW 168	rocm_error_checking(stream, this->language_extension_, "err");	×
NEW 169	}	×
NEW 170	}	×
171
NEW 172	codegen::InstrumentationInfo ROCMDataOffloadingNodeDispatcher::instrumentation_info() const {	×
NEW 173	auto& rocm_node = static_cast<const ROCMDataOffloadingNode&>(node_);	×
NEW 174	if (rocm_node.is_d2h()) {	×
NEW 175	return codegen::InstrumentationInfo(	×
NEW 176	node_.element_id(),	×
NEW 177	codegen::ElementType_D2HTransfer,	×
NEW 178	TargetType_ROCM,	×
NEW 179	analysis::LoopInfo{},	×
NEW 180	{{"pcie_bytes", language_extension_.expression(rocm_node.size())}}	×
NEW 181	);	×
NEW 182	} else if (rocm_node.is_h2d()) {	×
NEW 183	return codegen::InstrumentationInfo(	×
NEW 184	node_.element_id(),	×
NEW 185	codegen::ElementType_H2DTransfer,	×
NEW 186	TargetType_ROCM,	×
NEW 187	analysis::LoopInfo{},	×
NEW 188	{{"pcie_bytes", language_extension_.expression(rocm_node.size())}}	×
NEW 189	);	×
NEW 190	} else {	×
NEW 191	return codegen::LibraryNodeDispatcher::instrumentation_info();	×
NEW 192	}	×
NEW 193	}	×
194
NEW 195	nlohmann::json ROCMDataOffloadingNodeSerializer::serialize(const sdfg::data_flow::LibraryNode& library_node) {	×
NEW 196	const auto& node = static_cast<const ROCMDataOffloadingNode&>(library_node);	×
NEW 197	nlohmann::json j;	×
198
199	// Library node
NEW 200	j["type"] = "library_node";	×
NEW 201	j["element_id"] = library_node.element_id();	×
202
203	// Debug info
NEW 204	auto& debug_info = library_node.debug_info();	×
NEW 205	j["has"] = debug_info.has();	×
NEW 206	j["filename"] = debug_info.filename();	×
NEW 207	j["start_line"] = debug_info.start_line();	×
NEW 208	j["start_column"] = debug_info.start_column();	×
NEW 209	j["end_line"] = debug_info.end_line();	×
NEW 210	j["end_column"] = debug_info.end_column();	×
211
212	// Library node properties
NEW 213	j["code"] = std::string(library_node.code().value());	×
214
215	// Offloading node properties
NEW 216	sdfg::serializer::JSONSerializer serializer;	×
NEW 217	if (node.size().is_null()) {	×
NEW 218	j["size"] = nlohmann::json::value_t::null;	×
NEW 219	} else {	×
NEW 220	j["size"] = serializer.expression(node.size());	×
NEW 221	}	×
NEW 222	j["device_id"] = serializer.expression(node.device_id());	×
NEW 223	j["transfer_direction"] = static_cast<int8_t>(node.transfer_direction());	×
NEW 224	j["buffer_lifecycle"] = static_cast<int8_t>(node.buffer_lifecycle());	×
225
NEW 226	return j;	×
NEW 227	}	×
228
229	data_flow::LibraryNode& ROCMDataOffloadingNodeSerializer::deserialize(
230	const nlohmann::json& j, sdfg::builder::StructuredSDFGBuilder& builder, structured_control_flow::Block& parent
NEW 231	) {	×
NEW 232	auto code = j["code"].get<std::string>();	×
NEW 233	if (code != LibraryNodeType_ROCM_Offloading.value()) {	×
NEW 234	throw std::runtime_error("Invalid library node code");	×
NEW 235	}	×
236
NEW 237	sdfg::serializer::JSONSerializer serializer;	×
NEW 238	DebugInfo debug_info = serializer.json_to_debug_info(j["debug_info"]);	×
239
NEW 240	symbolic::Expression size;	×
NEW 241	if (!j.contains("size") \|\| j.at("size").is_null()) {	×
NEW 242	size = SymEngine::null;	×
NEW 243	} else {	×
NEW 244	size = symbolic::parse(j.at("size"));	×
NEW 245	}	×
NEW 246	SymEngine::Expression device_id(j.at("device_id"));	×
NEW 247	auto transfer_direction = static_cast<offloading::DataTransferDirection>(j["transfer_direction"].get<int8_t>());	×
NEW 248	auto buffer_lifecycle = static_cast<offloading::BufferLifecycle>(j["buffer_lifecycle"].get<int8_t>());	×
249
NEW 250	return builder.add_library_node<	×
NEW 251	ROCMDataOffloadingNode>(parent, debug_info, size, device_id, transfer_direction, buffer_lifecycle);	×
NEW 252	}	×
253
254	} // namespace rocm
255	} // namespace sdfg

daisytuner / docc / 22949872003

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous