26812260585

Committed 02 Jun 2026 08:49AM UTC coverage: 60.823% (-0.05%) from 60.869%

Build # 26812260585

Build Type

Pull #725

github

Committed by

web-flow

Commit Message

Merge eceff48b9 into cd25c9278

Pull Request Pull Request #725: Tensor node backport

Coverage Stats

599 of 1255 new or added lines in 51 files covered. (47.73%)

541 existing lines in 46 files now uncovered.

35094 of 57699 relevant lines covered (60.82%)

11078.43 hits per line

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

33.33

/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,
    offloading::DataTransferDirection transfer_direction,
    offloading::BufferLifecycle buffer_lifecycle,
    symbolic::Expression size,
    symbolic::Expression device_id
)
    : offloading::DataOffloadingNode(
          element_id,
          debug_info,
          vertex,
          parent,
          LibraryNodeType_ROCM_Offloading,
          transfer_direction,
          buffer_lifecycle,
          size
      ),
      device_id_(device_id) {}

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->transfer_direction(),
        this->buffer_lifecycle(),
        this->size(),
        this->device_id()
    );
}

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;
}

bool ROCMDataOffloadingNode::is_same_target(const DataOffloadingNode& other) const {
    return dynamic_cast<const ROCMDataOffloadingNode*>(&other) != nullptr;
    // TODO check device id
}

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_with_edges(
    codegen::CodegenOutput& out,
    std::vector<codegen::DispatchInput>& inputs,
    std::vector<codegen::DispatchOutput>& outputs
) {
    auto& offloading_node = static_cast<const ROCMDataOffloadingNode&>(this->node_);

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

    std::string dev_ptr;

    if (offloading_node.is_alloc()) {
        auto& ptr = outputs.at(0);
        pre_allocate_output(out, ptr, offloading_node.output(0));
        out.stream << "err = hipMalloc(&" << *ptr.local_name << ", "
                   << this->language_extension_.expression(offloading_node.size()) << ");" << std::endl;
        rocm_error_checking(out.stream, this->language_extension_, "err");
        dev_ptr = *ptr.local_name;
    } else {
        dev_ptr = inputs.at(offloading_node.dev_ptr_input_idx()).expr;
    }

    if (offloading_node.is_h2d()) {
        out.stream << "err = hipMemcpy(" << dev_ptr << ", " << inputs.at(offloading_node.host_ptr_input_idx()).expr
                   << ", " << this->language_extension_.expression(offloading_node.size())
                   << ", hipMemcpyHostToDevice);" << std::endl;
        rocm_error_checking(out.stream, this->language_extension_, "err");
    } else if (offloading_node.is_d2h()) {
        out.stream << "err = hipMemcpy(" << inputs.at(offloading_node.host_ptr_input_idx()).expr << ", " << dev_ptr
                   << ", " << this->language_extension_.expression(offloading_node.size())
                   << ", hipMemcpyHostToDevice);" << std::endl;
        rocm_error_checking(out.stream, this->language_extension_, "err");
    }

    if (offloading_node.is_free()) {
        out.stream << "err = hipFree(" << dev_ptr << ");" << std::endl;
        rocm_error_checking(out.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);
    auto j = offloading::DataOffloadingNodeSerializer::serialize(library_node);

    // Offloading node properties
    sdfg::serializer::JSONSerializer serializer;
    j["device_id"] = serializer.expression(node.device_id());

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

    SymEngine::Expression device_id(j.at("device_id"));

    return offloading::DataOffloadingNodeSerializer::deserialize_generic_offload<
        ROCMDataOffloadingNode>(j, builder, parent, device_id);
}

} // 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	offloading::DataTransferDirection transfer_direction,
32	offloading::BufferLifecycle buffer_lifecycle,
33	symbolic::Expression size,
34	symbolic::Expression device_id
35	)
36	: offloading::DataOffloadingNode(	18✔
37	element_id,	18✔
38	debug_info,	18✔
39	vertex,	18✔
40	parent,	18✔
41	LibraryNodeType_ROCM_Offloading,	18✔
42	transfer_direction,	18✔
43	buffer_lifecycle,	18✔
44	size	18✔
45	),	18✔
46	device_id_(device_id) {}	18✔
47
48	void ROCMDataOffloadingNode::validate(const Function& function) const {	2✔
49	// Prevent copy-in and free
50	if (this->is_h2d() && this->is_free()) {	2✔
51	throw InvalidSDFGException("ROCMDataOffloadingNode: Combination copy-in and free is not allowed");	×
52	}	×
53
54	// Prevent copy-out and alloc
55	if (this->is_d2h() && this->is_alloc()) {	2✔
56	throw InvalidSDFGException("ROCMDataOffloadingNode: Combination copy-out and alloc is not allowed");	×
57	}	×
58	}	2✔
59
60	const symbolic::Expression ROCMDataOffloadingNode::device_id() const { return this->device_id_; }	×
61
62	std::unique_ptr<data_flow::DataFlowNode> ROCMDataOffloadingNode::
63	clone(size_t element_id, const graph::Vertex vertex, data_flow::DataFlowGraph& parent) const {	×
64	return std::make_unique<ROCMDataOffloadingNode>(	×
65	element_id,	×
66	this->debug_info(),	×
67	vertex,	×
68	parent,	×
69	this->transfer_direction(),	×
70	this->buffer_lifecycle(),	×
71	this->size(),	×
72	this->device_id()	×
73	);	×
74	}	×
75
76	symbolic::SymbolSet ROCMDataOffloadingNode::symbols() const {	×
77	if (this->device_id().is_null()) {	×
78	return offloading::DataOffloadingNode::symbols();	×
79	}	×
80	auto symbols = offloading::DataOffloadingNode::symbols();	×
81	auto device_id_atoms = symbolic::atoms(this->device_id());	×
82	symbols.insert(device_id_atoms.begin(), device_id_atoms.end());	×
83	return symbols;	×
84	}	×
85
86	void ROCMDataOffloadingNode::replace(const symbolic::Expression old_expression, const symbolic::Expression new_expression) {	×
87	offloading::DataOffloadingNode::replace(old_expression, new_expression);	×
88	this->device_id_ = symbolic::subs(this->device_id_, old_expression, new_expression);	×
89	}	×
90
91	bool ROCMDataOffloadingNode::blocking() const { return true; }	×
92
93	bool ROCMDataOffloadingNode::redundant_with(const offloading::DataOffloadingNode& other) const {	×
94	if (!offloading::DataOffloadingNode::redundant_with(other)) {	×
95	return false;	×
96	}	×
97
98	auto& other_node = static_cast<const ROCMDataOffloadingNode&>(other);	×
99	if (!symbolic::null_safe_eq(this->device_id(), other_node.device_id())) {	×
100	return false;	×
101	}	×
102
103	return true;	×
104	}	×
105
106	bool ROCMDataOffloadingNode::equal_with(const offloading::DataOffloadingNode& other) const {	×
107	if (!offloading::DataOffloadingNode::equal_with(other)) {	×
108	return false;	×
109	}	×
110
111	auto& other_node = static_cast<const ROCMDataOffloadingNode&>(other);	×
112	if (!symbolic::null_safe_eq(this->device_id(), other_node.device_id())) {	×
113	return false;	×
114	}	×
115
116	return true;	×
117	}	×
118
119	bool ROCMDataOffloadingNode::is_same_target(const DataOffloadingNode& other) const {	×
120	return dynamic_cast<const ROCMDataOffloadingNode*>(&other) != nullptr;	×
121	// TODO check device id
122	}	×
123
124	ROCMDataOffloadingNodeDispatcher::ROCMDataOffloadingNodeDispatcher(
125	codegen::LanguageExtension& language_extension,
126	const Function& function,
127	const data_flow::DataFlowGraph& data_flow_graph,
128	const data_flow::LibraryNode& node
129	)
130	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	4✔
131
132	void ROCMDataOffloadingNodeDispatcher::dispatch_code_with_edges(
133	codegen::CodegenOutput& out,
134	std::vector<codegen::DispatchInput>& inputs,
135	std::vector<codegen::DispatchOutput>& outputs
136	) {	4✔
137	auto& offloading_node = static_cast<const ROCMDataOffloadingNode&>(this->node_);	4✔
138
139	out.stream << "hipError_t err;" << std::endl;	4✔
140
141	std::string dev_ptr;	4✔
142
143	if (offloading_node.is_alloc()) {	4✔
144	auto& ptr = outputs.at(0);	1✔
145	pre_allocate_output(out, ptr, offloading_node.output(0));	1✔
146	out.stream << "err = hipMalloc(&" << *ptr.local_name << ", "	1✔
147	<< this->language_extension_.expression(offloading_node.size()) << ");" << std::endl;	1✔
148	rocm_error_checking(out.stream, this->language_extension_, "err");	1✔
149	dev_ptr = *ptr.local_name;	1✔
150	} else {	3✔
151	dev_ptr = inputs.at(offloading_node.dev_ptr_input_idx()).expr;	3✔
152	}	3✔
153
154	if (offloading_node.is_h2d()) {	4✔
155	out.stream << "err = hipMemcpy(" << dev_ptr << ", " << inputs.at(offloading_node.host_ptr_input_idx()).expr	1✔
156	<< ", " << this->language_extension_.expression(offloading_node.size())	1✔
157	<< ", hipMemcpyHostToDevice);" << std::endl;	1✔
158	rocm_error_checking(out.stream, this->language_extension_, "err");	1✔
159	} else if (offloading_node.is_d2h()) {	3✔
160	out.stream << "err = hipMemcpy(" << inputs.at(offloading_node.host_ptr_input_idx()).expr << ", " << dev_ptr	1✔
161	<< ", " << this->language_extension_.expression(offloading_node.size())	1✔
162	<< ", hipMemcpyHostToDevice);" << std::endl;	1✔
163	rocm_error_checking(out.stream, this->language_extension_, "err");	1✔
164	}	1✔
165
166	if (offloading_node.is_free()) {	4✔
167	out.stream << "err = hipFree(" << dev_ptr << ");" << std::endl;	1✔
168	rocm_error_checking(out.stream, this->language_extension_, "err");	1✔
169	}	1✔
170	}	4✔
171
172	codegen::InstrumentationInfo ROCMDataOffloadingNodeDispatcher::instrumentation_info() const {	×
173	auto& rocm_node = static_cast<const ROCMDataOffloadingNode&>(node_);	×
174	if (rocm_node.is_d2h()) {	×
175	return codegen::InstrumentationInfo(	×
176	node_.element_id(),	×
177	codegen::ElementType_D2HTransfer,	×
178	TargetType_ROCM,	×
179	analysis::LoopInfo{},	×
180	{{"pcie_bytes", language_extension_.expression(rocm_node.size())}}	×
181	);	×
182	} else if (rocm_node.is_h2d()) {	×
183	return codegen::InstrumentationInfo(	×
184	node_.element_id(),	×
185	codegen::ElementType_H2DTransfer,	×
186	TargetType_ROCM,	×
UNCOV 187	analysis::LoopInfo{},	×
188	{{"pcie_bytes", language_extension_.expression(rocm_node.size())}}	×
189	);	×
190	} else {	×
UNCOV 191	return codegen::LibraryNodeDispatcher::instrumentation_info();	×
UNCOV 192	}	×
193	}	×
194
UNCOV 195	nlohmann::json ROCMDataOffloadingNodeSerializer::serialize(const sdfg::data_flow::LibraryNode& library_node) {	×
196	const auto& node = static_cast<const ROCMDataOffloadingNode&>(library_node);	×
197	auto j = offloading::DataOffloadingNodeSerializer::serialize(library_node);	×
198
199	// Offloading node properties
UNCOV 200	sdfg::serializer::JSONSerializer serializer;	×
201	j["device_id"] = serializer.expression(node.device_id());	×
202
203	return j;	×
204	}	×
205
206	data_flow::LibraryNode& ROCMDataOffloadingNodeSerializer::deserialize(
207	const nlohmann::json& j, sdfg::builder::StructuredSDFGBuilder& builder, structured_control_flow::Block& parent
UNCOV 208	) {	×
209	auto code = j["code"].get<std::string>();	×
210	if (code != LibraryNodeType_ROCM_Offloading.value()) {	×
211	throw std::runtime_error("Invalid library node code");	×
UNCOV 212	}	×
213
UNCOV 214	SymEngine::Expression device_id(j.at("device_id"));	×
215
UNCOV 216	return offloading::DataOffloadingNodeSerializer::deserialize_generic_offload<	×
UNCOV 217	ROCMDataOffloadingNode>(j, builder, parent, device_id);	×
UNCOV 218	}	×
219
220	} // namespace rocm
221	} // namespace sdfg

daisytuner / docc / 26812260585

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