28814257460

Committed 06 Jul 2026 06:31PM UTC coverage: 62.928% (-0.03%) from 62.96%

Build # 28814257460

Build Type

Pull #843

github

Committed by

web-flow

Commit Message

Merge b6b7f08d9 into 95ea95f1f

Pull Request Pull Request #843: adds type_ids for Element classes and a human-readable element_type

Coverage Stats

457 of 599 new or added lines in 121 files covered. (76.29%)

5 existing lines in 3 files now uncovered.

40598 of 64515 relevant lines covered (62.93%)

977.42 hits per line

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

65.22

/opt/src/targets/cuda/cuda_data_offloading_node.cpp

#include "sdfg/targets/cuda/cuda_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/cuda/cuda.h"
#include "sdfg/targets/offloading/data_offloading_node.h"
#include "symengine/symengine_rcp.h"

namespace sdfg {
namespace cuda {

CUDADataOffloadingNode::CUDADataOffloadingNode(
    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_CUDA_Offloading,
          transfer_direction,
          buffer_lifecycle,
          std::move(size)
      ),
      device_id_(std::move(device_id)) {}

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

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

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

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

symbolic::SymbolSet CUDADataOffloadingNode::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 CUDADataOffloadingNode::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 CUDADataOffloadingNode::blocking() const { return true; }

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

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

    return true;
}

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

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

    return true;
}

bool CUDADataOffloadingNode::is_same_target(const offloading::DataOffloadingNode& other) const {
    return dynamic_cast<const CUDADataOffloadingNode*>(&other) != nullptr;
    // TODO check for dev ID
}

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

    // stream << "cudaSetDevice(" << this->language_extension_.expression(offloading_node.device_id()) << ");"
    //        << std::endl;

    out.stream << "cudaError_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 = cudaMalloc(&" << *ptr.local_name << ", "
                   << this->language_extension_.expression(offloading_node.size()) << ");" << std::endl;
        cuda_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 = cudaMemcpy(" << dev_ptr << ", " << inputs.at(offloading_node.host_ptr_input_idx()).expr
                   << ", " << this->language_extension_.expression(offloading_node.size())
                   << ", cudaMemcpyHostToDevice);" << std::endl;
        cuda_error_checking(out.stream, this->language_extension_, "err");
    } else if (offloading_node.is_d2h()) {
        out.stream << "err = cudaMemcpy(" << inputs.at(offloading_node.host_ptr_input_idx()).expr << ", " << dev_ptr
                   << ", " << this->language_extension_.expression(offloading_node.size())
                   << ", cudaMemcpyDeviceToHost);" << std::endl;
        cuda_error_checking(out.stream, this->language_extension_, "err");
    }

    if (offloading_node.is_free()) {
        out.stream << "err = cudaFree(" << dev_ptr << ");" << std::endl;
        cuda_error_checking(out.stream, this->language_extension_, "err");
    }
}

codegen::InstrumentationInfo CUDADataOffloadingNodeDispatcher::instrumentation_info() const {
    auto& cuda_node = static_cast<const CUDADataOffloadingNode&>(node_);
    if (cuda_node.is_d2h()) {
        return codegen::InstrumentationInfo(
            node_.element_id(),
            std::string(cuda_node.element_type()) + ":::" + cuda_node.code().value(),
            TargetType_CUDA,
            analysis::LoopInfo{},
            {{"pcie_bytes", language_extension_.expression(cuda_node.size())}}
        );
    } else if (cuda_node.is_h2d()) {
        return codegen::InstrumentationInfo(
            node_.element_id(),
            std::string(cuda_node.element_type()) + ":::" + cuda_node.code().value(),
            TargetType_CUDA,
            analysis::LoopInfo{},
            {{"pcie_bytes", language_extension_.expression(cuda_node.size())}}
        );
    } else {
        return codegen::LibraryNodeDispatcher::instrumentation_info();
    }
}

nlohmann::json CUDADataOffloadingNodeSerializer::serialize(const sdfg::data_flow::LibraryNode& library_node) {
    auto j = DataOffloadingNodeSerializer::serialize(library_node);

    const auto& node = static_cast<const CUDADataOffloadingNode&>(library_node);

    // Offloading node properties
    serializer::JSONSerializer serializer;

    j["device_id"] = serializer.expression(node.device_id());

    return j;
}

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

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

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

} // namespace cuda
} // namespace sdfg

1	#include "sdfg/targets/cuda/cuda_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/cuda/cuda.h"
20	#include "sdfg/targets/offloading/data_offloading_node.h"
21	#include "symengine/symengine_rcp.h"
22
23	namespace sdfg {
24	namespace cuda {
25
26	CUDADataOffloadingNode::CUDADataOffloadingNode(
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(	211✔
37	element_id,	211✔
38	debug_info,	211✔
39	vertex,	211✔
40	parent,	211✔
41	LibraryNodeType_CUDA_Offloading,	211✔
42	transfer_direction,	211✔
43	buffer_lifecycle,	211✔
44	std::move(size)	211✔
45	),	211✔
46	device_id_(std::move(device_id)) {}	211✔
47
48	void CUDADataOffloadingNode::validate(const Function& function) const {	270✔
49	// Prevent copy-in and free
50	if (this->is_h2d() && this->is_free()) {	270✔
51	throw InvalidSDFGException("CUDADataOffloadingNode: Combination copy-in and free is not allowed");	×
52	}	×
53
54	// Prevent copy-out and alloc
55	if (this->is_d2h() && this->is_alloc()) {	270✔
56	throw InvalidSDFGException("CUDADataOffloadingNode: Combination copy-out and alloc is not allowed");	×
57	}	×
58	}	270✔
59
60	const symbolic::Expression CUDADataOffloadingNode::device_id() const { return this->device_id_; }	398✔
61
62	std::unique_ptr<data_flow::DataFlowNode> CUDADataOffloadingNode::
63	clone(size_t element_id, const graph::Vertex vertex, data_flow::DataFlowGraph& parent) const {	4✔
64	return std::make_unique<CUDADataOffloadingNode>(	4✔
65	element_id,	4✔
66	this->debug_info(),	4✔
67	vertex,	4✔
68	parent,	4✔
69	this->transfer_direction(),	4✔
70	this->buffer_lifecycle(),	4✔
71	this->size(),	4✔
72	this->device_id()	4✔
73	);	4✔
74	}	4✔
75
76	symbolic::SymbolSet CUDADataOffloadingNode::symbols() const {	185✔
77	if (this->device_id().is_null()) {	185✔
78	return offloading::DataOffloadingNode::symbols();	×
79	}	×
80	auto symbols = offloading::DataOffloadingNode::symbols();	185✔
81	auto device_id_atoms = symbolic::atoms(this->device_id());	185✔
82	symbols.insert(device_id_atoms.begin(), device_id_atoms.end());	185✔
83	return symbols;	185✔
84	}	185✔
85
86	void CUDADataOffloadingNode::replace(const symbolic::Expression old_expression, const symbolic::Expression new_expression) {	79✔
87	offloading::DataOffloadingNode::replace(old_expression, new_expression);	79✔
88	this->device_id_ = symbolic::subs(this->device_id_, old_expression, new_expression);	79✔
89	}	79✔
90
91	bool CUDADataOffloadingNode::blocking() const { return true; }	×
92
93	bool CUDADataOffloadingNode::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 CUDADataOffloadingNode&>(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 CUDADataOffloadingNode::equal_with(const offloading::DataOffloadingNode& other) const {	6✔
107	if (!offloading::DataOffloadingNode::equal_with(other)) {	6✔
108	return false;	×
109	}	×
110
111	auto& other_node = static_cast<const CUDADataOffloadingNode&>(other);	6✔
112	if (!symbolic::null_safe_eq(this->device_id(), other_node.device_id())) {	6✔
113	return false;	×
114	}	×
115
116	return true;	6✔
117	}	6✔
118
119	bool CUDADataOffloadingNode::is_same_target(const offloading::DataOffloadingNode& other) const {	×
120	return dynamic_cast<const CUDADataOffloadingNode*>(&other) != nullptr;	×
121	// TODO check for dev ID
122	}	×
123
124	CUDADataOffloadingNodeDispatcher::CUDADataOffloadingNodeDispatcher(
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) {}	6✔
131
132	void CUDADataOffloadingNodeDispatcher::dispatch_code_with_edges(
133	codegen::CodegenOutput& out,
134	std::vector<codegen::DispatchInput>& inputs,
135	std::vector<codegen::DispatchOutput>& outputs
136	) {	6✔
137	auto& offloading_node = static_cast<const CUDADataOffloadingNode&>(this->node_);	6✔
138
139	// stream << "cudaSetDevice(" << this->language_extension_.expression(offloading_node.device_id()) << ");"
140	// << std::endl;
141
142	out.stream << "cudaError_t err;" << std::endl;	6✔
143
144	std::string dev_ptr;	6✔
145	if (offloading_node.is_alloc()) {	6✔
146	auto& ptr = outputs.at(0);	2✔
147	pre_allocate_output(out, ptr, offloading_node.output(0));	2✔
148	out.stream << "err = cudaMalloc(&" << *ptr.local_name << ", "	2✔
149	<< this->language_extension_.expression(offloading_node.size()) << ");" << std::endl;	2✔
150	cuda_error_checking(out.stream, this->language_extension_, "err");	2✔
151	dev_ptr = *ptr.local_name;	2✔
152	} else {	4✔
153	dev_ptr = inputs.at(offloading_node.dev_ptr_input_idx()).expr;	4✔
154	}	4✔
155
156	if (offloading_node.is_h2d()) {	6✔
157	out.stream << "err = cudaMemcpy(" << dev_ptr << ", " << inputs.at(offloading_node.host_ptr_input_idx()).expr	1✔
158	<< ", " << this->language_extension_.expression(offloading_node.size())	1✔
159	<< ", cudaMemcpyHostToDevice);" << std::endl;	1✔
160	cuda_error_checking(out.stream, this->language_extension_, "err");	1✔
161	} else if (offloading_node.is_d2h()) {	5✔
162	out.stream << "err = cudaMemcpy(" << inputs.at(offloading_node.host_ptr_input_idx()).expr << ", " << dev_ptr	2✔
163	<< ", " << this->language_extension_.expression(offloading_node.size())	2✔
164	<< ", cudaMemcpyDeviceToHost);" << std::endl;	2✔
165	cuda_error_checking(out.stream, this->language_extension_, "err");	2✔
166	}	2✔
167
168	if (offloading_node.is_free()) {	6✔
169	out.stream << "err = cudaFree(" << dev_ptr << ");" << std::endl;	2✔
170	cuda_error_checking(out.stream, this->language_extension_, "err");	2✔
171	}	2✔
172	}	6✔
173
174	codegen::InstrumentationInfo CUDADataOffloadingNodeDispatcher::instrumentation_info() const {	×
175	auto& cuda_node = static_cast<const CUDADataOffloadingNode&>(node_);	×
176	if (cuda_node.is_d2h()) {	×
177	return codegen::InstrumentationInfo(	×
178	node_.element_id(),	×
NEW 179	std::string(cuda_node.element_type()) + ":::" + cuda_node.code().value(),	×
180	TargetType_CUDA,	×
181	analysis::LoopInfo{},	×
182	{{"pcie_bytes", language_extension_.expression(cuda_node.size())}}	×
183	);	×
184	} else if (cuda_node.is_h2d()) {	×
185	return codegen::InstrumentationInfo(	×
186	node_.element_id(),	×
NEW 187	std::string(cuda_node.element_type()) + ":::" + cuda_node.code().value(),	×
188	TargetType_CUDA,	×
189	analysis::LoopInfo{},	×
190	{{"pcie_bytes", language_extension_.expression(cuda_node.size())}}	×
191	);	×
192	} else {	×
193	return codegen::LibraryNodeDispatcher::instrumentation_info();	×
194	}	×
195	}	×
196
197	nlohmann::json CUDADataOffloadingNodeSerializer::serialize(const sdfg::data_flow::LibraryNode& library_node) {	4✔
198	auto j = DataOffloadingNodeSerializer::serialize(library_node);	4✔
199
200	const auto& node = static_cast<const CUDADataOffloadingNode&>(library_node);	4✔
201
202	// Offloading node properties
203	serializer::JSONSerializer serializer;	4✔
204
205	j["device_id"] = serializer.expression(node.device_id());	4✔
206
207	return j;	4✔
208	}	4✔
209
210	data_flow::LibraryNode& CUDADataOffloadingNodeSerializer::deserialize(
211	const nlohmann::json& j, sdfg::builder::StructuredSDFGBuilder& builder, structured_control_flow::Block& parent
212	) {	4✔
213	auto code = j["code"].get<std::string>();	4✔
214	if (code != LibraryNodeType_CUDA_Offloading.value()) {	4✔
215	throw std::runtime_error("Invalid library node code");	×
216	}	×
217
218	SymEngine::Expression device_id(j.at("device_id"));	4✔
219
220	return offloading::DataOffloadingNodeSerializer::deserialize_generic_offload<	4✔
221	CUDADataOffloadingNode>(j, builder, parent, device_id);	4✔
222	}	4✔
223
224	} // namespace cuda
225	} // namespace sdfg

daisytuner / docc / 28814257460

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