26556322966

Committed 27 May 2026 03:45PM UTC coverage: 60.869% (-0.02%) from 60.886%

Build # 26556322966

Build Type

push

github

Committed by

web-flow

Commit Message

Libnode ptr edges (#719)

Migrating SDFGs to treat pointers as inputs to libNodes / Calls as scalars.
A pointer will only appear in an output edge if its actually returned from the function (like malloc).

* Stdlib, Blas and Tensor Matmul nodes were migrated to this new format. Other, currently transitory Tensor Nodes are not yet migrated.
* DOCC version was bumped to incorporate previous docc-llvm versions (up to 0.4.0) that had been counted separately.
! Until all passes consider the use / leak of pointers as uncertainty / hiding potential writes, TensorNodes are declared as general side-effect.
* Lots of utility functions to centralize the creation (and edges) of various libNodes that needed to be changed.
* Fixed & unified docc paths across python and llvm front-ends.
* Skip BlockFusion test that fails to its libNodes currently having side effects
~ Prevent a crash in DotViz when using symbolic offsets into structs
* Removing old ConstProp pass, it is not safe for the new pointer representation and should not be all too critical

Coverage Stats

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

87 existing lines in 28 files now uncovered.

35225 of 57870 relevant lines covered (60.87%)

11046.32 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 / 26556322966

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