24845644334

Committed 23 Apr 2026 04:05PM UTC coverage: 64.104% (-0.06%) from 64.167%

Build # 24845644334

Build Type

push

github

Committed by

web-flow

Commit Message

Modular Compile process (#690)

 + DoccTarget to allow modularly adding support for snippets that need to be compiled differently (standalone, cross-compiled kernels or just with special compilers)
 + Builder infrastructure to define C/C++ build options with inheritance, overrides and ability to modify by plugins, extensions and targets as needed
 + The redirected handling of snippets based on the extension can contribute link-options when they have been used (the generic solution to highway support)
 + Model to handle any parallism for compile in its own class without having to touch the definitions and compile options
 + DoccPaths extended to resolve default plugin dirs (as used by et-plugin)
 + PrettyPrinter can write directly to another (file) stream instead of a local buffer
 * extended CodeGenerator to allow generating the header first and direct-to-file (its needed for all snippet & main compile, which could be parallel otherwise)
 + Defaults to parallel builds with as many cores as present
 + DOCC_DEBUG is more defined. A list of key-value pairs separated by ; or : Supported options: "dump" dump sdfgs, "build" add debug symbols, "build_threads=x" override thread count. less then 1 is the default auto-selection.
 + CodegenStats is integrated into the new Compiler concept. To reduce the impact of the asynchronicity, times are always measured, but only recored with the CodegenStats if enabled and later, in sequential places. The items recorded now include the sdfg & snippet names to identify slow parts
 + Switched MacOS tests over to inplace-python with only 1 build

Coverage Stats

224 of 347 new or added lines in 13 files covered. (64.55%)

2 existing lines in 2 files now uncovered.

30800 of 48047 relevant lines covered (64.1%)

573.39 hits per line

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

59.88

/sdfg/src/codegen/code_generators/cuda_code_generator.cpp

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

#include "sdfg/codegen/dispatchers/node_dispatcher_registry.h"
#include "sdfg/codegen/instrumentation/instrumentation_plan.h"
#include "sdfg/helpers/helpers.h"

namespace sdfg {
namespace codegen {

CUDACodeGenerator::CUDACodeGenerator(
    StructuredSDFG& sdfg,
    analysis::AnalysisManager& analysis_manager,
    InstrumentationPlan& instrumentation_plan,
    ArgCapturePlan& arg_capture_plan,
    std::shared_ptr<CodeSnippetFactory> library_snippet_factory
)
    : CodeGenerator(sdfg, analysis_manager, instrumentation_plan, arg_capture_plan, std::move(library_snippet_factory)),
      language_extension_(sdfg) {
    if (sdfg.type() != FunctionType_NV_GLOBAL) {
        throw std::runtime_error("CUDACodeGenerator can only be used for GPU SDFGs");
    }
    if (!arg_capture_plan_.is_empty()) {
        DEBUG_PRINTLN("CUDACodeGenerator does not support capturing args/results!");
    }
};

bool CUDACodeGenerator::generate() {
    this->dispatch_includes();
    this->dispatch_structures();
    this->dispatch_globals();
    this->dispatch_schedule();
    return true;
};

std::string CUDACodeGenerator::function_definition() {
    /********** Arglist **********/
    std::vector<std::string> args;
    for (auto& container : sdfg_.arguments()) {
        args.push_back(language_extension_.declaration(container, sdfg_.type(container)));
    }
    std::stringstream arglist;
    arglist << sdfg::helpers::join(args, ", ");

    return "extern \"C\" __global__ void " + sdfg_.name() + "(" + arglist.str() + ")";
};

bool CUDACodeGenerator::emit_header(PrettyPrinter& out) {
    out << "#pragma once" << std::endl;
    dispatch_header_includes(out);
    dispatch_header_structures(out);
    return true;
}

bool CUDACodeGenerator::emit_main_source(std::ostream& out, const std::filesystem::path& header_path) {
    out << "#include \"" << header_path.filename().string() << "\"" << std::endl;
    dispatch_globals();
    dispatch_schedule();

    out << this->globals_stream_.str() << std::endl;

    append_function_source(out);

    return true;
}

bool CUDACodeGenerator::as_source(const std::filesystem::path& header_path, const std::filesystem::path& source_path) {
    std::ofstream ofs_header(header_path, std::ofstream::out);
    if (!ofs_header.is_open()) {
        return false;
    }

    std::ofstream ofs_source(source_path, std::ofstream::out);
    if (!ofs_source.is_open()) {
        return false;
    }

    ofs_header << "#pragma once" << std::endl;
    ofs_header << this->includes_stream_.str() << std::endl;
    ofs_header << this->classes_stream_.str() << std::endl;
    ofs_header.close();

    ofs_source << "#include \"" << header_path.filename().string() << "\"" << std::endl;

    ofs_source << this->globals_stream_.str() << std::endl;

    append_function_source(ofs_source);

    ofs_source.close();

    return true;
};

void CUDACodeGenerator::append_function_source(std::ostream& ofs_source) {
    ofs_source << this->function_definition() << std::endl;
    ofs_source << "{" << std::endl;
    ofs_source << this->main_stream_.str() << std::endl;
    ofs_source << "}" << std::endl;
}

void CUDACodeGenerator::dispatch_includes() { this->dispatch_header_includes(this->includes_stream_); }

void CUDACodeGenerator::dispatch_header_includes(PrettyPrinter& out) {
    out << "#define "
        << "__DAISY_NVVM__" << std::endl;
    out << "#include <cstdint>" << std::endl;
    out << "#include <daisy_rtl/daisy_rtl.h>" << std::endl;
};

void CUDACodeGenerator::dispatch_structures() { this->dispatch_header_structures(this->classes_stream_); }

void CUDACodeGenerator::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 CUDACodeGenerator::dispatch_globals() {
    for (auto& container : sdfg_.externals()) {
        auto& type = sdfg_.type(container);
        if (type.storage_type().is_nv_global()) {
            auto& base_type = dynamic_cast<const types::Pointer&>(type).pointee_type();
            if (sdfg_.linkage_type(container) == LinkageType_External) {
                this->globals_stream_ << "extern " << language_extension_.declaration(container, base_type) << ";"
                                      << std::endl;
            } else {
                this->globals_stream_ << "static " << language_extension_.declaration(container, base_type);
                if (!type.initializer().empty()) {
                    this->globals_stream_ << " = " << type.initializer();
                }
                this->globals_stream_ << ";" << std::endl;
            }
        }
        if (type.storage_type().is_nv_constant()) {
            assert(type.initializer().empty());
            auto& base_type = dynamic_cast<const types::Pointer&>(type).pointee_type();
            this->globals_stream_ << "__constant__ " << language_extension_.declaration(container, base_type, true)
                                  << ";" << std::endl;
        }
    }
};

void CUDACodeGenerator::dispatch_schedule() {
    // Declare shared memory
    for (auto& container : sdfg_.externals()) {
        auto& type = sdfg_.type(container);
        if (type.storage_type().is_nv_shared()) {
            this->main_stream_ << language_extension_.declaration(container, sdfg_.type(container)) << ";" << std::endl;
        }
    }

    // Declare transient containers
    for (auto& container : sdfg_.containers()) {
        if (!sdfg_.is_transient(container)) {
            continue;
        }

        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;
        }
    }

    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_);
};

} // namespace codegen
} // namespace sdfg

1	#include "sdfg/codegen/code_generators/cuda_code_generator.h"
2
3	#include "sdfg/codegen/dispatchers/node_dispatcher_registry.h"
4	#include "sdfg/codegen/instrumentation/instrumentation_plan.h"
5	#include "sdfg/helpers/helpers.h"
6
7	namespace sdfg {
8	namespace codegen {
9
10	CUDACodeGenerator::CUDACodeGenerator(
11	StructuredSDFG& sdfg,
12	analysis::AnalysisManager& analysis_manager,
13	InstrumentationPlan& instrumentation_plan,
14	ArgCapturePlan& arg_capture_plan,
15	std::shared_ptr<CodeSnippetFactory> library_snippet_factory
16	)
17	: CodeGenerator(sdfg, analysis_manager, instrumentation_plan, arg_capture_plan, std::move(library_snippet_factory)),	4✔
18	language_extension_(sdfg) {	4✔
19	if (sdfg.type() != FunctionType_NV_GLOBAL) {	4✔
20	throw std::runtime_error("CUDACodeGenerator can only be used for GPU SDFGs");	×
21	}	×
22	if (!arg_capture_plan_.is_empty()) {	4✔
23	DEBUG_PRINTLN("CUDACodeGenerator does not support capturing args/results!");	×
24	}	×
25	};	4✔
26
27	bool CUDACodeGenerator::generate() {	3✔
28	this->dispatch_includes();	3✔
29	this->dispatch_structures();	3✔
30	this->dispatch_globals();	3✔
31	this->dispatch_schedule();	3✔
32	return true;	3✔
33	};	3✔
34
35	std::string CUDACodeGenerator::function_definition() {	1✔
36	/******** Arglist ********/
37	std::vector<std::string> args;	1✔
38	for (auto& container : sdfg_.arguments()) {	1✔
39	args.push_back(language_extension_.declaration(container, sdfg_.type(container)));	×
40	}	×
41	std::stringstream arglist;	1✔
42	arglist << sdfg::helpers::join(args, ", ");	1✔
43
44	return "extern \"C\" __global__ void " + sdfg_.name() + "(" + arglist.str() + ")";	1✔
45	};	1✔
46
NEW 47	bool CUDACodeGenerator::emit_header(PrettyPrinter& out) {	×
NEW 48	out << "#pragma once" << std::endl;	×
NEW 49	dispatch_header_includes(out);	×
NEW 50	dispatch_header_structures(out);	×
NEW 51	return true;	×
NEW 52	}	×
53
NEW 54	bool CUDACodeGenerator::emit_main_source(std::ostream& out, const std::filesystem::path& header_path) {	×
NEW 55	out << "#include \"" << header_path.filename().string() << "\"" << std::endl;	×
NEW 56	dispatch_globals();	×
NEW 57	dispatch_schedule();	×
58
NEW 59	out << this->globals_stream_.str() << std::endl;	×
60
NEW 61	append_function_source(out);	×
62
NEW 63	return true;	×
NEW 64	}	×
65
66	bool CUDACodeGenerator::as_source(const std::filesystem::path& header_path, const std::filesystem::path& source_path) {	×
67	std::ofstream ofs_header(header_path, std::ofstream::out);	×
68	if (!ofs_header.is_open()) {	×
69	return false;	×
70	}	×
71
72	std::ofstream ofs_source(source_path, std::ofstream::out);	×
73	if (!ofs_source.is_open()) {	×
74	return false;	×
75	}	×
76
77	ofs_header << "#pragma once" << std::endl;	×
78	ofs_header << this->includes_stream_.str() << std::endl;	×
79	ofs_header << this->classes_stream_.str() << std::endl;	×
80	ofs_header.close();	×
81
82	ofs_source << "#include \"" << header_path.filename().string() << "\"" << std::endl;	×
83
84	ofs_source << this->globals_stream_.str() << std::endl;	×
85
86	append_function_source(ofs_source);	×
87
88	ofs_source.close();	×
89
90	return true;	×
91	};	×
92
NEW 93	void CUDACodeGenerator::append_function_source(std::ostream& ofs_source) {	×
94	ofs_source << this->function_definition() << std::endl;	×
95	ofs_source << "{" << std::endl;	×
96	ofs_source << this->main_stream_.str() << std::endl;	×
97	ofs_source << "}" << std::endl;	×
98	}	×
99
100	void CUDACodeGenerator::dispatch_includes() { this->dispatch_header_includes(this->includes_stream_); }	3✔
101
102	void CUDACodeGenerator::dispatch_header_includes(PrettyPrinter& out) {	3✔
103	out << "#define "	3✔
104	<< "__DAISY_NVVM__" << std::endl;	3✔
105	out << "#include <cstdint>" << std::endl;	3✔
106	out << "#include <daisy_rtl/daisy_rtl.h>" << std::endl;	3✔
107	};	3✔
108
109	void CUDACodeGenerator::dispatch_structures() { this->dispatch_header_structures(this->classes_stream_); }	3✔
110
111	void CUDACodeGenerator::dispatch_header_structures(PrettyPrinter& out) {	3✔
112	// Forward declarations
113	for (auto& structure : sdfg_.structures()) {	3✔
114	out << "struct " << structure << ";" << std::endl;	3✔
115	}	3✔
116
117	// Generate topology-sorted structure definitions
118	typedef boost::adjacency_list<boost::vecS, boost::vecS, boost::bidirectionalS> structures_graph;	3✔
119	typedef boost::graph_traits<structures_graph>::vertex_descriptor Vertex;	3✔
120	std::vector<std::string> names;	3✔
121	for (auto& structure : sdfg_.structures()) {	3✔
122	names.push_back(structure);	3✔
123	}	3✔
124	structures_graph graph(names.size());	3✔
125
126	for (auto& structure : names) {	3✔
127	auto& definition = sdfg_.structure(structure);	3✔
128	for (size_t i = 0; i < definition.num_members(); i++) {	6✔
129	auto member_type = &definition.member_type(symbolic::integer(i));	3✔
130	while (dynamic_cast<const types::Array*>(member_type)) {	3✔
131	auto array_type = static_cast<const types::Array*>(member_type);	×
132	member_type = &array_type->element_type();	×
133	}	×
134
135	if (auto member_structure = dynamic_cast<const sdfg::types::Structure*>(member_type)) {	3✔
136	boost::add_edge(	1✔
137	std::find(names.begin(), names.end(), member_structure->name()) - names.begin(),	1✔
138	std::find(names.begin(), names.end(), structure) - names.begin(),	1✔
139	graph	1✔
140	);	1✔
141	}	1✔
142	}	3✔
143	}	3✔
144
145	std::list<Vertex> order;	3✔
146	std::unordered_map<Vertex, boost::default_color_type> vertex_colors;	3✔
147	boost::topological_sort(	3✔
148	graph, std::back_inserter(order), boost::color_map(boost::make_assoc_property_map(vertex_colors))	3✔
149	);	3✔
150	order.reverse();	3✔
151
152	for (auto& structure_index : order) {	3✔
153	std::string structure = names.at(structure_index);	3✔
154	auto& definition = sdfg_.structure(structure);	3✔
155	out << "struct ";	3✔
156	if (definition.is_packed()) {	3✔
NEW 157	out << "__attribute__((packed)) ";	×
158	}	×
159	out << structure << std::endl;	3✔
160	out << "{\n";	3✔
161
162	for (size_t i = 0; i < definition.num_members(); i++) {	6✔
163	auto& member_type = definition.member_type(symbolic::integer(i));	3✔
164	if (dynamic_cast<const sdfg::types::Structure*>(&member_type)) {	3✔
165	out << "struct ";	1✔
166	}	1✔
167	out << language_extension_.declaration("member_" + std::to_string(i), member_type, false, true);	3✔
168	out << ";" << std::endl;	3✔
169	}	3✔
170
171	out << "};" << std::endl;	3✔
172	}	3✔
173	};	3✔
174
175	void CUDACodeGenerator::dispatch_globals() {	3✔
176	for (auto& container : sdfg_.externals()) {	3✔
177	auto& type = sdfg_.type(container);	1✔
178	if (type.storage_type().is_nv_global()) {	1✔
179	auto& base_type = dynamic_cast<const types::Pointer&>(type).pointee_type();	1✔
180	if (sdfg_.linkage_type(container) == LinkageType_External) {	1✔
181	this->globals_stream_ << "extern " << language_extension_.declaration(container, base_type) << ";"	1✔
182	<< std::endl;	1✔
183	} else {	1✔
184	this->globals_stream_ << "static " << language_extension_.declaration(container, base_type);	×
185	if (!type.initializer().empty()) {	×
186	this->globals_stream_ << " = " << type.initializer();	×
187	}	×
188	this->globals_stream_ << ";" << std::endl;	×
189	}	×
190	}	1✔
191	if (type.storage_type().is_nv_constant()) {	1✔
192	assert(type.initializer().empty());	×
193	auto& base_type = dynamic_cast<const types::Pointer&>(type).pointee_type();	×
194	this->globals_stream_ << "__constant__ " << language_extension_.declaration(container, base_type, true)	×
195	<< ";" << std::endl;	×
196	}	×
197	}	1✔
198	};	3✔
199
200	void CUDACodeGenerator::dispatch_schedule() {	3✔
201	// Declare shared memory
202	for (auto& container : sdfg_.externals()) {	3✔
203	auto& type = sdfg_.type(container);	1✔
204	if (type.storage_type().is_nv_shared()) {	1✔
205	this->main_stream_ << language_extension_.declaration(container, sdfg_.type(container)) << ";" << std::endl;	×
206	}	×
207	}	1✔
208
209	// Declare transient containers
210	for (auto& container : sdfg_.containers()) {	3✔
211	if (!sdfg_.is_transient(container)) {	1✔
212	continue;	1✔
213	}	1✔
214
215	std::string val = this->language_extension_.declaration(container, sdfg_.type(container), false, true);	×
216	if (!val.empty()) {	×
217	this->main_stream_ << val;	×
218	this->main_stream_ << ";" << std::endl;	×
219	}	×
220	}	×
221
222	auto dispatcher = create_dispatcher(	3✔
223	language_extension_, sdfg_, analysis_manager_, sdfg_.root(), instrumentation_plan_, arg_capture_plan_	3✔
224	);	3✔
225	dispatcher->dispatch(this->main_stream_, this->globals_stream_, *this->library_snippet_factory_);	3✔
226	};	3✔
227
228	} // namespace codegen
229	} // namespace sdfg

daisytuner / docc / 24845644334

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