27759308722

Committed 18 Jun 2026 12:26PM UTC coverage: 61.622% (-0.01%) from 61.632%

Build # 27759308722

Build Type

Pull #776

github

Committed by

web-flow

Commit Message

Merge 887424b32 into 6f06476da

Pull Request Pull Request #776: Adds codegen option to emit a global constructor

Coverage Stats

2 of 14 new or added lines in 2 files covered. (14.29%)

2 existing lines in 1 file now uncovered.

36663 of 59497 relevant lines covered (61.62%)

1022.11 hits per line

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

64.36

/sdfg/src/codegen/code_generators/cpp_code_generator.cpp

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

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

namespace sdfg {
namespace codegen {

std::string CPPCodeGenerator::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 \"C\" " + this->language_extension_.declaration("", sdfg_.return_type()) + this->externals_prefix_ +
           sdfg_.name() + "(" + arglist_str + ")";
};

void CPPCodeGenerator::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 __daisy_capture_t* __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 CPPCodeGenerator::dispatch_header_includes(PrettyPrinter& out) {
#if defined(__linux__)
    out << "#include <alloca.h>" << std::endl;
    out << "#include <malloc.h>" << std::endl;
#endif

    out << "#include <cmath>" << std::endl;
    out << "#include <cstdio>" << std::endl;
    out << "#include <cstdlib>" << std::endl;
    out << "#include <cstring>" << std::endl;
    out << "#include <cstdint>" << std::endl;

    out << "#include <daisy_rtl/daisy_rtl.h>" << std::endl;
}

void CPPCodeGenerator::dispatch_header_structures(PrettyPrinter& out) {
    // Forward declarations
    for (auto& structure : sdfg_.structures()) {
        out << "struct " << 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);
        out << "struct ";
        if (definition.is_packed()) {
            out << "__attribute__((packed)) ";
        }
        out << structure << std::endl;
        out << "{\n";

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

        out << "};" << std::endl;
    }
}

void CPPCodeGenerator::dispatch_globals() {
    // Optional one-shot GPU runtime / context warmup, executed at .so load time
    // via __attribute__((constructor)). Moves the cold CUDA/HIP driver+context
    // init cost out of the first kernel call (and out of any instrumented region
    // wrapping it). Opt-in: callers set GlobalConstructor at construction time.
    if (this->global_constructor_ != GlobalConstructor::None) {
        // No explicit cuda_runtime.h / hip_runtime.h include: the source file
        // is compiled with `-x cuda` / `-x hip`, which provides the runtime
        // API declarations implicitly (matches the rest of the emitted host
        // wrapper, e.g. cudaSetDevice/cudaMalloc below).
        this->globals_stream_ << "static void __attribute__((constructor)) "
                                 "__daisy_gpu_context_warmup(void) {"
                              << std::endl;
        if (this->global_constructor_ == GlobalConstructor::CUDA) {
            this->globals_stream_ << "    (void)cudaSetDevice(0);" << std::endl;
            this->globals_stream_ << "    (void)cudaFree(0);" << std::endl;
        } else if (this->global_constructor_ == GlobalConstructor::ROCm) {
            this->globals_stream_ << "    (void)hipSetDevice(0);" << std::endl;
            this->globals_stream_ << "    (void)hipFree(0);" << std::endl;
        }
        this->globals_stream_ << "}" << std::endl;
    }

    // Declare globals
    for (auto& container : sdfg_.externals()) {
        // Function declarations
        if (dynamic_cast<const types::Function*>(&sdfg_.type(container))) {
            this->globals_stream_ << "extern \"C\" ";
            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 CPPCodeGenerator::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;
            }
        }

        // Allocate if needed
        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/cpp_code_generator.h"
2
3	#include "sdfg/codegen/dispatchers/node_dispatcher_registry.h"
4
5	namespace sdfg {
6	namespace codegen {
7
8	std::string CPPCodeGenerator::function_definition() {	2✔
9	/******** Arglist ********/
10	std::vector<std::string> args;	2✔
11	for (auto& container : this->sdfg_.arguments()) {	3✔
12	args.push_back(language_extension_.declaration(container, this->sdfg_.type(container)));	3✔
13	}	3✔
14
15	std::stringstream arglist;	2✔
16	arglist << sdfg::helpers::join(args, ", ");	2✔
17	std::string arglist_str = arglist.str();	2✔
18	if (arglist_str.empty()) {	2✔
19	arglist_str = "void";	1✔
20	}	1✔
21
22	return "extern \"C\" " + this->language_extension_.declaration("", sdfg_.return_type()) + this->externals_prefix_ +	2✔
23	sdfg_.name() + "(" + arglist_str + ")";	2✔
24	};	2✔
25
26	void CPPCodeGenerator::emit_capture_context_init(std::ostream& ofs_source) const {	×
27	std::string name = sdfg_.name();	×
28	std::string arg_capture_path = sdfg_.metadata().at("arg_capture_path");	×
29
30	ofs_source << "static __daisy_capture_t* __capture_ctx;" << std::endl;	×
31	ofs_source << "static void __attribute__((constructor(1000))) __capture_ctx_init(void) {" << std::endl;	×
32	ofs_source << "\t__capture_ctx = __daisy_capture_init(\"" << name << "\", \"" << arg_capture_path << "\");"	×
33	<< std::endl;	×
34	ofs_source << "}" << std::endl;	×
35	ofs_source << std::endl;	×
36	}	×
37
38	void CPPCodeGenerator::dispatch_header_includes(PrettyPrinter& out) {	7✔
39	#if defined(__linux__)	7✔
40	out << "#include <alloca.h>" << std::endl;	7✔
41	out << "#include <malloc.h>" << std::endl;	7✔
42	#endif	7✔
43
44	out << "#include <cmath>" << std::endl;	7✔
45	out << "#include <cstdio>" << std::endl;	7✔
46	out << "#include <cstdlib>" << std::endl;	7✔
47	out << "#include <cstring>" << std::endl;	7✔
48	out << "#include <cstdint>" << std::endl;	7✔
49
50	out << "#include <daisy_rtl/daisy_rtl.h>" << std::endl;	7✔
51	}	7✔
52
53	void CPPCodeGenerator::dispatch_header_structures(PrettyPrinter& out) {	7✔
54	// Forward declarations
55	for (auto& structure : sdfg_.structures()) {	7✔
56	out << "struct " << structure << ";" << std::endl;	4✔
57	}	4✔
58
59	// Generate topology-sorted structure definitions
60	typedef boost::adjacency_list<boost::vecS, boost::vecS, boost::bidirectionalS> structures_graph;	7✔
61	typedef boost::graph_traits<structures_graph>::vertex_descriptor Vertex;	7✔
62	std::vector<std::string> names;	7✔
63	for (auto& structure : sdfg_.structures()) {	7✔
64	names.push_back(structure);	4✔
65	}	4✔
66	structures_graph graph(names.size());	7✔
67
68	for (auto& structure : names) {	7✔
69	auto& definition = sdfg_.structure(structure);	4✔
70	for (size_t i = 0; i < definition.num_members(); i++) {	8✔
71	auto member_type = &definition.member_type(symbolic::integer(i));	4✔
72	while (dynamic_cast<const types::Array*>(member_type)) {	4✔
73	auto array_type = static_cast<const types::Array*>(member_type);	×
74	member_type = &array_type->element_type();	×
75	}	×
76
77	if (auto member_structure = dynamic_cast<const sdfg::types::Structure*>(member_type)) {	4✔
78	boost::add_edge(	1✔
79	std::find(names.begin(), names.end(), member_structure->name()) - names.begin(),	1✔
80	std::find(names.begin(), names.end(), structure) - names.begin(),	1✔
81	graph	1✔
82	);	1✔
83	}	1✔
84	}	4✔
85	}	4✔
86
87	std::list<Vertex> order;	7✔
88	std::unordered_map<Vertex, boost::default_color_type> vertex_colors;	7✔
89	boost::topological_sort(	7✔
90	graph, std::back_inserter(order), boost::color_map(boost::make_assoc_property_map(vertex_colors))	7✔
91	);	7✔
92	order.reverse();	7✔
93
94	for (auto& structure_index : order) {	7✔
95	std::string structure = names.at(structure_index);	4✔
96	auto& definition = sdfg_.structure(structure);	4✔
97	out << "struct ";	4✔
98	if (definition.is_packed()) {	4✔
99	out << "__attribute__((packed)) ";	1✔
100	}	1✔
101	out << structure << std::endl;	4✔
102	out << "{\n";	4✔
103
104	for (size_t i = 0; i < definition.num_members(); i++) {	8✔
105	auto& member_type = definition.member_type(symbolic::integer(i));	4✔
106	if (dynamic_cast<const sdfg::types::Structure*>(&member_type)) {	4✔
107	out << "struct ";	1✔
108	}	1✔
109	out << language_extension_.declaration("member_" + std::to_string(i), member_type, false, true);	4✔
110	out << ";" << std::endl;	4✔
111	}	4✔
112
113	out << "};" << std::endl;	4✔
114	}	4✔
115	}	7✔
116
117	void CPPCodeGenerator::dispatch_globals() {	7✔
118	// Optional one-shot GPU runtime / context warmup, executed at .so load time
119	// via __attribute__((constructor)). Moves the cold CUDA/HIP driver+context
120	// init cost out of the first kernel call (and out of any instrumented region
121	// wrapping it). Opt-in: callers set GlobalConstructor at construction time.
122	if (this->global_constructor_ != GlobalConstructor::None) {	7✔
123	// No explicit cuda_runtime.h / hip_runtime.h include: the source file
124	// is compiled with `-x cuda` / `-x hip`, which provides the runtime
125	// API declarations implicitly (matches the rest of the emitted host
126	// wrapper, e.g. cudaSetDevice/cudaMalloc below).
NEW 127	this->globals_stream_ << "static void __attribute__((constructor)) "	×
NEW 128	"__daisy_gpu_context_warmup(void) {"	×
NEW 129	<< std::endl;	×
NEW 130	if (this->global_constructor_ == GlobalConstructor::CUDA) {	×
NEW 131	this->globals_stream_ << " (void)cudaSetDevice(0);" << std::endl;	×
NEW 132	this->globals_stream_ << " (void)cudaFree(0);" << std::endl;	×
NEW 133	} else if (this->global_constructor_ == GlobalConstructor::ROCm) {	×
NEW 134	this->globals_stream_ << " (void)hipSetDevice(0);" << std::endl;	×
NEW 135	this->globals_stream_ << " (void)hipFree(0);" << std::endl;	×
NEW 136	}	×
NEW 137	this->globals_stream_ << "}" << std::endl;	×
NEW 138	}	×
139
140	// Declare globals
141	for (auto& container : sdfg_.externals()) {	7✔
142	// Function declarations
143	if (dynamic_cast<const types::Function*>(&sdfg_.type(container))) {	1✔
144	this->globals_stream_ << "extern \"C\" ";	×
145	this->globals_stream_	×
146	<< language_extension_.declaration(this->externals_prefix_ + container, sdfg_.type(container)) << ";"	×
147	<< std::endl;	×
148	continue;	×
149	}	×
150
151	// Other types must be pointers
152	auto& type = dynamic_cast<const types::Pointer&>(sdfg_.type(container));	1✔
153	assert(type.has_pointee_type() && "Externals must have a pointee type");	1✔
154	auto& base_type = type.pointee_type();	1✔
155
156	if (sdfg_.linkage_type(container) == LinkageType_External) {	1✔
157	this->globals_stream_ << "extern "	1✔
158	<< language_extension_.declaration(this->externals_prefix_ + container, base_type)	1✔
159	<< ";" << std::endl;	1✔
160	} else {	1✔
161	this->globals_stream_ << "static "	×
162	<< language_extension_.declaration(this->externals_prefix_ + container, base_type);	×
163	if (!type.initializer().empty()) {	×
164	this->globals_stream_ << " = " << type.initializer();	×
165	}	×
166	this->globals_stream_ << ";" << std::endl;	×
167	}	×
168	}	1✔
169	};	7✔
170
171	void CPPCodeGenerator::dispatch_schedule() {	7✔
172	// Allocate variables
173	for (auto& container : sdfg_.containers()) {	10✔
174	if (sdfg_.is_external(container)) {	10✔
175	continue;	1✔
176	}	1✔
177	auto& type = sdfg_.type(container);	9✔
178
179	// Declare transient
180	if (sdfg_.is_transient(container)) {	9✔
181	std::string val = this->language_extension_.declaration(container, sdfg_.type(container), false, true);	5✔
182	if (!val.empty()) {	5✔
183	this->main_stream_ << val;	5✔
184	this->main_stream_ << ";" << std::endl;	5✔
185	}	5✔
186	}	5✔
187
188	// Allocate if needed
189	if (type.storage_type().allocation() == types::StorageType::AllocationType::Managed) {	9✔
190	assert(	×
191	!type.storage_type().allocation_size().is_null() &&	×
192	"Managed allocations must have a valid allocation size"	×
193	);	×
194	if (type.storage_type().is_cpu_stack()) {	×
195	this->main_stream_ << container << " = ";	×
196	this->main_stream_	×
197	<< "alloca(" << this->language_extension_.expression(type.storage_type().allocation_size()) << ")";	×
198	this->main_stream_ << ";" << std::endl;	×
199	} else if (type.storage_type().is_cpu_heap()) {	×
200	this->main_stream_ << container << " = ";	×
201	this->main_stream_ << this->externals_prefix_ << "malloc("	×
202	<< this->language_extension_.expression(type.storage_type().allocation_size())	×
203	<< ")";	×
204	this->main_stream_ << ";" << std::endl;	×
205	} else if (type.storage_type().is_nv_generic()) {	×
206	this->main_stream_ << "cudaSetDevice(0);" << std::endl;	×
207	this->main_stream_ << "cudaMalloc(&" << container << ", "	×
208	<< this->language_extension_.expression(type.storage_type().allocation_size())	×
209	<< ");" << std::endl;	×
210	}	×
211	}	×
212	}	9✔
213
214	auto dispatcher = create_dispatcher(	7✔
215	language_extension_, sdfg_, analysis_manager_, sdfg_.root(), instrumentation_plan_, arg_capture_plan_	7✔
216	);	7✔
217	dispatcher->dispatch(this->main_stream_, this->globals_stream_, *this->library_snippet_factory_);	7✔
218
219	if (sdfg_.root().size() == 0 \|\|	7✔
220	!dynamic_cast<const structured_control_flow::Return*>(&(sdfg_.root().at(sdfg_.root().size() - 1)).first)) {	7✔
221	// Free heap allocations
222	for (auto& container : sdfg_.containers()) {	9✔
223	if (sdfg_.is_external(container)) {	9✔
224	continue;	1✔
225	}	1✔
226	auto& type = sdfg_.type(container);	8✔
227
228	// Free if needed
229	if (type.storage_type().deallocation() == types::StorageType::AllocationType::Managed) {	8✔
230	if (type.storage_type().is_cpu_heap()) {	×
231	this->main_stream_ << this->externals_prefix_ << "free(" << container << ");" << std::endl;	×
232	} else if (type.storage_type().is_nv_generic()) {	×
233	this->main_stream_ << "cudaSetDevice(0);" << std::endl;	×
234	this->main_stream_ << "cudaFree(" << container << ");" << std::endl;	×
235	}	×
236	}	×
237	}	8✔
238	}	6✔
239	};	7✔
240
241	} // namespace codegen
242	} // namespace sdfg

daisytuner / docc / 27759308722

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