26463753889

Committed 26 May 2026 05:18PM UTC coverage: 60.864% (-0.02%) from 60.886%

Build # 26463753889

Build Type

Pull #719

github

Committed by

web-flow

Commit Message

Merge 0b90ddd88 into 707dadcf8

Pull Request Pull Request #719: Libnode ptr edges

Coverage Stats

961 of 1749 new or added lines in 52 files covered. (54.95%)

90 existing lines in 29 files now uncovered.

35222 of 57870 relevant lines covered (60.86%)

11043.61 hits per line

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

11.36

/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(
    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);
    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(	14✔
37	element_id,	14✔
38	debug_info,	14✔
39	vertex,	14✔
40	parent,	14✔
41	LibraryNodeType_ROCM_Offloading,	14✔
42	transfer_direction,	14✔
43	buffer_lifecycle,	14✔
44	size	14✔
45	),	14✔
46	device_id_(device_id) {}	14✔
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(),	×
NEW 70	this->buffer_lifecycle(),	×
NEW 71	this->size(),	×
NEW 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) {}	×
131
132	void ROCMDataOffloadingNodeDispatcher::dispatch_code(
133	codegen::PrettyPrinter& stream,
134	codegen::PrettyPrinter& globals_stream,
135	codegen::CodeSnippetFactory& library_snippet_factory
136	) {	×
137	auto& offloading_node = static_cast<const ROCMDataOffloadingNode&>(this->node_);	×
138
139	stream << "hipError_t err;" << std::endl;	×
140
141	if (offloading_node.is_alloc()) {	×
142	stream << "err = hipMalloc(&" << offloading_node.output(0) << ", "	×
143	<< this->language_extension_.expression(offloading_node.size()) << ");" << std::endl;	×
144	rocm_error_checking(stream, this->language_extension_, "err");	×
145	}	×
146
147	if (offloading_node.is_h2d()) {	×
148	stream << "err = hipMemcpy(" << offloading_node.output(0) << ", " << offloading_node.input(0) << ", "	×
149	<< this->language_extension_.expression(offloading_node.size()) << ", hipMemcpyHostToDevice);"	×
150	<< std::endl;	×
151	rocm_error_checking(stream, this->language_extension_, "err");	×
152	} else if (offloading_node.is_d2h()) {	×
153	stream << "err = hipMemcpy(" << offloading_node.output(0) << ", " << offloading_node.input(0) << ", "	×
154	<< this->language_extension_.expression(offloading_node.size()) << ", hipMemcpyDeviceToHost);"	×
155	<< std::endl;	×
156	rocm_error_checking(stream, this->language_extension_, "err");	×
157	}	×
158
159	if (offloading_node.is_free()) {	×
160	stream << "err = hipFree(" << offloading_node.input(0) << ");" << std::endl;	×
161	rocm_error_checking(stream, this->language_extension_, "err");	×
162	}	×
163	}	×
164
165	codegen::InstrumentationInfo ROCMDataOffloadingNodeDispatcher::instrumentation_info() const {	×
166	auto& rocm_node = static_cast<const ROCMDataOffloadingNode&>(node_);	×
167	if (rocm_node.is_d2h()) {	×
168	return codegen::InstrumentationInfo(	×
169	node_.element_id(),	×
170	codegen::ElementType_D2HTransfer,	×
171	TargetType_ROCM,	×
172	analysis::LoopInfo{},	×
173	{{"pcie_bytes", language_extension_.expression(rocm_node.size())}}	×
174	);	×
175	} else if (rocm_node.is_h2d()) {	×
176	return codegen::InstrumentationInfo(	×
177	node_.element_id(),	×
178	codegen::ElementType_H2DTransfer,	×
179	TargetType_ROCM,	×
180	analysis::LoopInfo{},	×
181	{{"pcie_bytes", language_extension_.expression(rocm_node.size())}}	×
182	);	×
183	} else {	×
184	return codegen::LibraryNodeDispatcher::instrumentation_info();	×
185	}	×
186	}	×
187
188	nlohmann::json ROCMDataOffloadingNodeSerializer::serialize(const sdfg::data_flow::LibraryNode& library_node) {	×
189	const auto& node = static_cast<const ROCMDataOffloadingNode&>(library_node);	×
NEW 190	auto j = offloading::DataOffloadingNodeSerializer::serialize(library_node);	×
191
192	// Offloading node properties
193	sdfg::serializer::JSONSerializer serializer;	×
194	j["device_id"] = serializer.expression(node.device_id());	×
195
196	return j;	×
197	}	×
198
199	data_flow::LibraryNode& ROCMDataOffloadingNodeSerializer::deserialize(
200	const nlohmann::json& j, sdfg::builder::StructuredSDFGBuilder& builder, structured_control_flow::Block& parent
201	) {	×
202	auto code = j["code"].get<std::string>();	×
203	if (code != LibraryNodeType_ROCM_Offloading.value()) {	×
204	throw std::runtime_error("Invalid library node code");	×
205	}	×
206
207	SymEngine::Expression device_id(j.at("device_id"));	×
208
NEW 209	return offloading::DataOffloadingNodeSerializer::deserialize_generic_offload<	×
NEW 210	ROCMDataOffloadingNode>(j, builder, parent, device_id);	×
UNCOV 211	}	×
212
213	} // namespace rocm
214	} // namespace sdfg

daisytuner / docc / 26463753889

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