24844557075

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

Build # 24844557075

Build Type

Pull #690

github

Committed by

web-flow

Commit Message

Merge f9cbf3e49 into 3762dada8

Pull Request Pull Request #690: Modular Compile process

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 / 24844557075

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