19162352523

Committed 07 Nov 2025 08:16AM UTC coverage: 61.556% (-0.4%) from 61.911%

Build # 19162352523

Build Type

push

github

Committed by

web-flow

Commit Message

rewrite of arg captures to work on a scope level (#319)

* rewrite of arg captures to work on a scope level

* adding first unit tests for arg capturing

* Fix scoped arg captures

* Fix debug output

* Only write capture fila name to index

* Switch to element-id based capture storage

* Reenable debug prints

* Fix serialization deserialization of arg capture index

* Use fake node ids in rtl test

* Add debug output for capture file creation

* Boost coverage

---------

Co-authored-by: Nora Hagmeyer <nora.hagmeyer@daisytuner.com>

Run Details

171 of 320 new or added lines in 18 files covered. (53.44%)

17 existing lines in 5 files now uncovered.

10265 of 16676 relevant lines covered (61.56%)

100.85 hits per line

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

71.79

/src/codegen/code_generators/c_code_generator.cpp

#include "sdfg/codegen/code_generators/c_code_generator.h"

#include "sdfg/codegen/dispatchers/node_dispatcher_registry.h"

namespace sdfg {
namespace codegen {

std::string CCodeGenerator::function_definition() {
    /********** Arglist **********/
    std::vector<std::string> args;
    for (auto& container : this->sdfg_.arguments()) {
        args.push_back(language_extension_.declaration(container, this->sdfg_.type(container)));
    }

    std::stringstream arglist;
    arglist << sdfg::helpers::join(args, ", ");
    std::string arglist_str = arglist.str();
    if (arglist_str.empty()) {
        arglist_str = "void";
    }

    return "extern " + this->language_extension_.declaration("", sdfg_.return_type()) + this->externals_prefix_ +
           sdfg_.name() + "(" + arglist_str + ")";
};

void CCodeGenerator::emit_capture_context_init(std::ostream& ofs_source) const {
    std::string name = sdfg_.name();
    std::string arg_capture_path = sdfg_.metadata().at("arg_capture_path");

    ofs_source << "static void* __capture_ctx;" << std::endl;
    ofs_source << "static void __attribute__((constructor(1000))) __capture_ctx_init(void) {" << std::endl;
    ofs_source << "\t__capture_ctx = __daisy_capture_init(\"" << name << "\", \"" << arg_capture_path << "\");"
               << std::endl;
    ofs_source << "}" << std::endl;
    ofs_source << std::endl;
}

void CCodeGenerator::dispatch_includes() {
    this->includes_stream_ << "#include <math.h>" << std::endl;
    this->includes_stream_ << "#include <alloca.h>" << std::endl;
    this->includes_stream_ << "#include <stdbool.h>" << std::endl;
    this->includes_stream_ << "#include <stdio.h>" << std::endl;
    this->includes_stream_ << "#include <stdlib.h>" << std::endl;
    this->includes_stream_ << "#include <string.h>" << std::endl;
    this->includes_stream_ << "#include <cblas.h>" << std::endl;
    this->includes_stream_ << "#include <daisy_rtl/daisy_rtl.h>" << std::endl;
};

void CCodeGenerator::dispatch_structures() {
    // Forward declarations
    for (auto& structure : sdfg_.structures()) {
        this->classes_stream_ << "typedef struct " << structure << " " << structure << ";" << std::endl;
    }

    // Generate topology-sorted structure definitions
    typedef boost::adjacency_list<boost::vecS, boost::vecS, boost::bidirectionalS> structures_graph;
    typedef boost::graph_traits<structures_graph>::vertex_descriptor Vertex;
    std::vector<std::string> names;
    for (auto& structure : sdfg_.structures()) {
        names.push_back(structure);
    }
    structures_graph graph(names.size());

    for (auto& structure : names) {
        auto& definition = sdfg_.structure(structure);
        for (size_t i = 0; i < definition.num_members(); i++) {
            auto member_type = &definition.member_type(symbolic::integer(i));
            while (dynamic_cast<const types::Array*>(member_type)) {
                auto array_type = static_cast<const types::Array*>(member_type);
                member_type = &array_type->element_type();
            }

            if (auto member_structure = dynamic_cast<const sdfg::types::Structure*>(member_type)) {
                boost::add_edge(
                    std::find(names.begin(), names.end(), member_structure->name()) - names.begin(),
                    std::find(names.begin(), names.end(), structure) - names.begin(),
                    graph
                );
            }
        }
    }

    std::list<Vertex> order;
    std::unordered_map<Vertex, boost::default_color_type> vertex_colors;
    boost::topological_sort(
        graph, std::back_inserter(order), boost::color_map(boost::make_assoc_property_map(vertex_colors))
    );
    order.reverse();

    for (auto& structure_index : order) {
        std::string structure = names.at(structure_index);
        auto& definition = sdfg_.structure(structure);
        this->classes_stream_ << "typedef struct ";
        if (definition.is_packed()) {
            this->classes_stream_ << "__attribute__((packed)) ";
        }
        this->classes_stream_ << structure << std::endl;
        this->classes_stream_ << "{\n";

        for (size_t i = 0; i < definition.num_members(); i++) {
            auto& member_type = definition.member_type(symbolic::integer(i));
            if (auto pointer_type = dynamic_cast<const sdfg::types::Pointer*>(&member_type)) {
                if (pointer_type->has_pointee_type() &&
                    dynamic_cast<const sdfg::types::Structure*>(&pointer_type->pointee_type())) {
                    this->classes_stream_ << "struct ";
                }
            }
            this->classes_stream_
                << language_extension_.declaration("member_" + std::to_string(i), member_type, false, true);
            this->classes_stream_ << ";" << std::endl;
        }

        this->classes_stream_ << "} " << structure << ";" << std::endl;
    }
};

void CCodeGenerator::dispatch_globals() {
    // Declare globals
    for (auto& container : sdfg_.externals()) {
        // Function declarations
        if (dynamic_cast<const types::Function*>(&sdfg_.type(container))) {
            // Declare function
            this->globals_stream_ << "extern ";
            this->globals_stream_
                << language_extension_.declaration(this->externals_prefix_ + container, sdfg_.type(container)) << ";"
                << std::endl;
            continue;
        }

        // Other types must be pointers
        auto& type = dynamic_cast<const types::Pointer&>(sdfg_.type(container));
        assert(type.has_pointee_type() && "Externals must have a pointee type");
        auto& base_type = type.pointee_type();

        if (sdfg_.linkage_type(container) == LinkageType_External) {
            this->globals_stream_ << "extern "
                                  << language_extension_.declaration(this->externals_prefix_ + container, base_type)
                                  << ";" << std::endl;
        } else {
            this->globals_stream_ << "static "
                                  << language_extension_.declaration(this->externals_prefix_ + container, base_type);
            if (!type.initializer().empty()) {
                this->globals_stream_ << " = " << type.initializer();
            }
            this->globals_stream_ << ";" << std::endl;
        }
    }
};

void CCodeGenerator::dispatch_schedule() {
    // Allocate variables
    for (auto& container : sdfg_.containers()) {
        if (sdfg_.is_external(container)) {
            continue;
        }
        auto& type = sdfg_.type(container);

        // Declare transient
        if (sdfg_.is_transient(container)) {
            std::string val = this->language_extension_.declaration(container, sdfg_.type(container), false, true);
            if (!val.empty()) {
                this->main_stream_ << val;
                this->main_stream_ << ";" << std::endl;
            }
        }

        if (type.storage_type().allocation() == types::StorageType::AllocationType::Managed) {
            assert(
                !type.storage_type().allocation_size().is_null() &&
                "Managed allocations must have a valid allocation size"
            );
            if (type.storage_type().is_cpu_stack()) {
                this->main_stream_ << container << " = ";
                this->main_stream_
                    << "alloca(" << this->language_extension_.expression(type.storage_type().allocation_size()) << ")";
                this->main_stream_ << ";" << std::endl;
            } else if (type.storage_type().is_cpu_heap()) {
                this->main_stream_ << container << " = ";
                this->main_stream_ << this->externals_prefix_ << "malloc("
                                   << this->language_extension_.expression(type.storage_type().allocation_size())
                                   << ")";
                this->main_stream_ << ";" << std::endl;
            } else if (type.storage_type().is_nv_generic()) {
                this->main_stream_ << "cudaSetDevice(0);" << std::endl;
                this->main_stream_ << "cudaMalloc(&" << container << ", "
                                   << this->language_extension_.expression(type.storage_type().allocation_size())
                                   << ");" << std::endl;
            }
        }
    }

    auto dispatcher = create_dispatcher(
        language_extension_, sdfg_, analysis_manager_, sdfg_.root(), instrumentation_plan_, arg_capture_plan_
    );
    dispatcher->dispatch(this->main_stream_, this->globals_stream_, this->library_snippet_factory_);

    if (sdfg_.root().size() == 0 ||
        !dynamic_cast<const structured_control_flow::Return*>(&(sdfg_.root().at(sdfg_.root().size() - 1)).first)) {
        // Free heap allocations
        for (auto& container : sdfg_.containers()) {
            if (sdfg_.is_external(container)) {
                continue;
            }
            auto& type = sdfg_.type(container);

            // Free if needed
            if (type.storage_type().deallocation() == types::StorageType::AllocationType::Managed) {
                if (type.storage_type().is_cpu_heap()) {
                    this->main_stream_ << this->externals_prefix_ << "free(" << container << ");" << std::endl;
                } else if (type.storage_type().is_nv_generic()) {
                    this->main_stream_ << "cudaSetDevice(0);" << std::endl;
                    this->main_stream_ << "cudaFree(" << container << ");" << std::endl;
                }
            }
        }
    }
};

} // namespace codegen
} // namespace sdfg

1	#include "sdfg/codegen/code_generators/c_code_generator.h"
2
3	#include "sdfg/codegen/dispatchers/node_dispatcher_registry.h"
4
5	namespace sdfg {
6	namespace codegen {
7
8	std::string CCodeGenerator::function_definition() {	1✔
9	/******** Arglist ********/
10	std::vector<std::string> args;	1✔
11	for (auto& container : this->sdfg_.arguments()) {	1✔
12	args.push_back(language_extension_.declaration(container, this->sdfg_.type(container)));	×
13	}
14
15	std::stringstream arglist;	1✔
16	arglist << sdfg::helpers::join(args, ", ");	1✔
17	std::string arglist_str = arglist.str();	1✔
18	if (arglist_str.empty()) {	1✔
19	arglist_str = "void";	1✔
20	}	1✔
21
22	return "extern " + this->language_extension_.declaration("", sdfg_.return_type()) + this->externals_prefix_ +	3✔
23	sdfg_.name() + "(" + arglist_str + ")";	2✔
24	};	1✔
25
UNCOV 26	void CCodeGenerator::emit_capture_context_init(std::ostream& ofs_source) const {	×
UNCOV 27	std::string name = sdfg_.name();	×
UNCOV 28	std::string arg_capture_path = sdfg_.metadata().at("arg_capture_path");	×
29
UNCOV 30	ofs_source << "static void* __capture_ctx;" << std::endl;	×
UNCOV 31	ofs_source << "static void __attribute__((constructor(1000))) __capture_ctx_init(void) {" << std::endl;	×
UNCOV 32	ofs_source << "\t__capture_ctx = __daisy_capture_init(\"" << name << "\", \"" << arg_capture_path << "\");"	×
UNCOV 33	<< std::endl;	×
UNCOV 34	ofs_source << "}" << std::endl;	×
UNCOV 35	ofs_source << std::endl;	×
UNCOV 36	}	×
37
38	void CCodeGenerator::dispatch_includes() {	10✔
39	this->includes_stream_ << "#include <math.h>" << std::endl;	10✔
40	this->includes_stream_ << "#include <alloca.h>" << std::endl;	10✔
41	this->includes_stream_ << "#include <stdbool.h>" << std::endl;	10✔
42	this->includes_stream_ << "#include <stdio.h>" << std::endl;	10✔
43	this->includes_stream_ << "#include <stdlib.h>" << std::endl;	10✔
44	this->includes_stream_ << "#include <string.h>" << std::endl;	10✔
45	this->includes_stream_ << "#include <cblas.h>" << std::endl;	10✔
46	this->includes_stream_ << "#include <daisy_rtl/daisy_rtl.h>" << std::endl;	10✔
47	};	10✔
48
49	void CCodeGenerator::dispatch_structures() {	10✔
50	// Forward declarations
51	for (auto& structure : sdfg_.structures()) {	13✔
52	this->classes_stream_ << "typedef struct " << structure << " " << structure << ";" << std::endl;	3✔
53	}
54
55	// Generate topology-sorted structure definitions
56	typedef boost::adjacency_list<boost::vecS, boost::vecS, boost::bidirectionalS> structures_graph;
57	typedef boost::graph_traits<structures_graph>::vertex_descriptor Vertex;
58	std::vector<std::string> names;	10✔
59	for (auto& structure : sdfg_.structures()) {	13✔
60	names.push_back(structure);	3✔
61	}
62	structures_graph graph(names.size());	10✔
63
64	for (auto& structure : names) {	13✔
65	auto& definition = sdfg_.structure(structure);	3✔
66	for (size_t i = 0; i < definition.num_members(); i++) {	6✔
67	auto member_type = &definition.member_type(symbolic::integer(i));	3✔
68	while (dynamic_cast<const types::Array*>(member_type)) {	3✔
69	auto array_type = static_cast<const types::Array*>(member_type);	×
70	member_type = &array_type->element_type();	×
71	}
72
73	if (auto member_structure = dynamic_cast<const sdfg::types::Structure*>(member_type)) {	3✔
74	boost::add_edge(	1✔
75	std::find(names.begin(), names.end(), member_structure->name()) - names.begin(),	1✔
76	std::find(names.begin(), names.end(), structure) - names.begin(),	1✔
77	graph
78	);
79	}	1✔
80	}	3✔
81	}
82
83	std::list<Vertex> order;	10✔
84	std::unordered_map<Vertex, boost::default_color_type> vertex_colors;	10✔
85	boost::topological_sort(	10✔
86	graph, std::back_inserter(order), boost::color_map(boost::make_assoc_property_map(vertex_colors))	10✔
87	);
88	order.reverse();	10✔
89
90	for (auto& structure_index : order) {	13✔
91	std::string structure = names.at(structure_index);	3✔
92	auto& definition = sdfg_.structure(structure);	3✔
93	this->classes_stream_ << "typedef struct ";	3✔
94	if (definition.is_packed()) {	3✔
95	this->classes_stream_ << "__attribute__((packed)) ";	×
96	}	×
97	this->classes_stream_ << structure << std::endl;	3✔
98	this->classes_stream_ << "{\n";	3✔
99
100	for (size_t i = 0; i < definition.num_members(); i++) {	6✔
101	auto& member_type = definition.member_type(symbolic::integer(i));	3✔
102	if (auto pointer_type = dynamic_cast<const sdfg::types::Pointer*>(&member_type)) {	3✔
103	if (pointer_type->has_pointee_type() &&	×
104	dynamic_cast<const sdfg::types::Structure*>(&pointer_type->pointee_type())) {	×
105	this->classes_stream_ << "struct ";	×
106	}	×
107	}	×
108	this->classes_stream_	6✔
109	<< language_extension_.declaration("member_" + std::to_string(i), member_type, false, true);	3✔
110	this->classes_stream_ << ";" << std::endl;	3✔
111	}	3✔
112
113	this->classes_stream_ << "} " << structure << ";" << std::endl;	3✔
114	}	3✔
115	};	10✔
116
117	void CCodeGenerator::dispatch_globals() {	10✔
118	// Declare globals
119	for (auto& container : sdfg_.externals()) {	11✔
120	// Function declarations
121	if (dynamic_cast<const types::Function*>(&sdfg_.type(container))) {	1✔
122	// Declare function
123	this->globals_stream_ << "extern ";	×
124	this->globals_stream_	×
125	<< language_extension_.declaration(this->externals_prefix_ + container, sdfg_.type(container)) << ";"	×
126	<< std::endl;	×
127	continue;	×
128	}
129
130	// Other types must be pointers
131	auto& type = dynamic_cast<const types::Pointer&>(sdfg_.type(container));	1✔
132	assert(type.has_pointee_type() && "Externals must have a pointee type");	2✔
133	auto& base_type = type.pointee_type();	1✔
134
135	if (sdfg_.linkage_type(container) == LinkageType_External) {	1✔
136	this->globals_stream_ << "extern "	2✔
137	<< language_extension_.declaration(this->externals_prefix_ + container, base_type)	1✔
138	<< ";" << std::endl;	1✔
139	} else {	1✔
140	this->globals_stream_ << "static "	×
141	<< language_extension_.declaration(this->externals_prefix_ + container, base_type);	×
142	if (!type.initializer().empty()) {	×
143	this->globals_stream_ << " = " << type.initializer();	×
144	}	×
145	this->globals_stream_ << ";" << std::endl;	×
146	}
147	}
148	};	10✔
149
150	void CCodeGenerator::dispatch_schedule() {	10✔
151	// Allocate variables
152	for (auto& container : sdfg_.containers()) {	18✔
153	if (sdfg_.is_external(container)) {	8✔
154	continue;	1✔
155	}
156	auto& type = sdfg_.type(container);	7✔
157
158	// Declare transient
159	if (sdfg_.is_transient(container)) {	7✔
160	std::string val = this->language_extension_.declaration(container, sdfg_.type(container), false, true);	3✔
161	if (!val.empty()) {	3✔
162	this->main_stream_ << val;	3✔
163	this->main_stream_ << ";" << std::endl;	3✔
164	}	3✔
165	}	3✔
166
167	if (type.storage_type().allocation() == types::StorageType::AllocationType::Managed) {	7✔
168	assert(	8✔
169	!type.storage_type().allocation_size().is_null() &&
170	"Managed allocations must have a valid allocation size"
171	);
172	if (type.storage_type().is_cpu_stack()) {	4✔
173	this->main_stream_ << container << " = ";	×
174	this->main_stream_	×
175	<< "alloca(" << this->language_extension_.expression(type.storage_type().allocation_size()) << ")";	×
176	this->main_stream_ << ";" << std::endl;	×
177	} else if (type.storage_type().is_cpu_heap()) {	4✔
178	this->main_stream_ << container << " = ";	4✔
179	this->main_stream_ << this->externals_prefix_ << "malloc("	8✔
180	<< this->language_extension_.expression(type.storage_type().allocation_size())	4✔
181	<< ")";	4✔
182	this->main_stream_ << ";" << std::endl;	4✔
183	} else if (type.storage_type().is_nv_generic()) {	4✔
184	this->main_stream_ << "cudaSetDevice(0);" << std::endl;	×
185	this->main_stream_ << "cudaMalloc(&" << container << ", "	×
186	<< this->language_extension_.expression(type.storage_type().allocation_size())	×
187	<< ");" << std::endl;	×
188	}	×
189	}	4✔
190	}
191
192	auto dispatcher = create_dispatcher(	10✔
193	language_extension_, sdfg_, analysis_manager_, sdfg_.root(), instrumentation_plan_, arg_capture_plan_	10✔
194	);
195	dispatcher->dispatch(this->main_stream_, this->globals_stream_, this->library_snippet_factory_);	10✔
196
197	if (sdfg_.root().size() == 0 \|\|	10✔
198	!dynamic_cast<const structured_control_flow::Return*>(&(sdfg_.root().at(sdfg_.root().size() - 1)).first)) {	×
199	// Free heap allocations
200	for (auto& container : sdfg_.containers()) {	18✔
201	if (sdfg_.is_external(container)) {	8✔
202	continue;	1✔
203	}
204	auto& type = sdfg_.type(container);	7✔
205
206	// Free if needed
207	if (type.storage_type().deallocation() == types::StorageType::AllocationType::Managed) {	7✔
208	if (type.storage_type().is_cpu_heap()) {	3✔
209	this->main_stream_ << this->externals_prefix_ << "free(" << container << ");" << std::endl;	3✔
210	} else if (type.storage_type().is_nv_generic()) {	3✔
211	this->main_stream_ << "cudaSetDevice(0);" << std::endl;	×
212	this->main_stream_ << "cudaFree(" << container << ");" << std::endl;	×
213	}	×
214	}	3✔
215	}
216	}	10✔
217	};	10✔
218
219	} // namespace codegen
220	} // namespace sdfg

daisytuner / sdfglib / 19162352523

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