15003172388

Committed 13 May 2025 05:40PM UTC coverage: 57.685% (-4.0%) from 61.65%

Build # 15003172388

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #10 from daisytuner/variadic-functions

Support for Packed Structs and Function Types

Run Details

33 of 110 new or added lines in 10 files covered. (30.0%)

457 existing lines in 6 files now uncovered.

7161 of 12414 relevant lines covered (57.68%)

526.43 hits per line

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

63.09

/src/codegen/code_generators/cuda_code_generator.cpp

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

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

namespace sdfg {
namespace codegen {

CUDACodeGenerator::CUDACodeGenerator(ConditionalSchedule& schedule, bool instrumented)
    : CodeGenerator(schedule, instrumented){

      };

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

std::string CUDACodeGenerator::function_definition() {
    // Define SDFG as a function
    auto& function = this->schedule_.schedule(0).sdfg();

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

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

bool CUDACodeGenerator::as_source(const std::filesystem::path& header_path,
                                  const std::filesystem::path& source_path,
                                  const std::filesystem::path& library_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;
    }

    std::ofstream ofs_library(library_path, std::ofstream::out);
    if (!ofs_library.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;
    ofs_source << this->function_definition() << std::endl;
    ofs_source << "{" << std::endl;
    ofs_source << this->main_stream_.str() << std::endl;
    ofs_source << "}" << std::endl;
    ofs_source.close();

    auto library_content = this->library_stream_.str();
    if (library_content.empty()) {
        ofs_library.close();
        return true;
    }

    ofs_library << "#include \"" << header_path.filename().string() << "\"" << std::endl;
    ofs_library << library_content << std::endl;
    ofs_library.close();

    return true;
};

void CUDACodeGenerator::dispatch_includes() {
    this->includes_stream_ << "#define "
                           << "__DAISY_NVVM__" << std::endl;
    this->includes_stream_ << "#include "
                           << "\"daisyrtl.h\"" << std::endl;
    if (instrumented_) this->includes_stream_ << "#include <daisy_rtl.h>" << std::endl;

    this->includes_stream_ << "#define __daisy_min(a,b) ((a)<(b)?(a):(b))" << std::endl;
    this->includes_stream_ << "#define __daisy_max(a,b) ((a)>(b)?(a):(b))" << std::endl;
    this->includes_stream_ << "#define __daisy_fma(a,b,c) a * b + c" << std::endl;
};

void CUDACodeGenerator::dispatch_structures() {
    auto& function = this->schedule_.schedule(0).sdfg();

    // Forward declarations
    for (auto& structure : function.structures()) {
        this->classes_stream_ << "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 : function.structures()) {
        names.push_back(structure);
    }
    structures_graph graph(names.size());

    for (auto& structure : names) {
        auto& definition = function.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 = function.structure(structure);
        this->classes_stream_ << "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 (dynamic_cast<const sdfg::types::Structure*>(&member_type)) {
                this->classes_stream_ << "struct ";
            }
            this->classes_stream_ << language_extension_.declaration("member_" + std::to_string(i),
                                                                     member_type);
            this->classes_stream_ << ";" << std::endl;
        }

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

void CUDACodeGenerator::dispatch_globals() {
    auto& function = schedule_.schedule(0).sdfg();
    for (auto& container : function.externals()) {
        auto& type = function.type(container);
        if (type.address_space() != 3 && type.address_space() != 4) {
            this->globals_stream_ << "extern " << language_extension_.declaration(container, type)
                                  << ";" << std::endl;
        }
        if (type.address_space() == 4) {
            assert(type.initializer().empty());
            this->globals_stream_ << "__constant__ "
                                  << language_extension_.declaration(container, type, true) << ";"
                                  << std::endl;
        }
    }
};

void CUDACodeGenerator::dispatch_schedule() {
    auto& function = schedule_.schedule(0).sdfg();

    // Declare shared memory
    for (auto& container : function.externals()) {
        auto& type = function.type(container);
        if (type.address_space() == 3) {
            this->main_stream_ << language_extension_.declaration(container,
                                                                  function.type(container))
                               << ";" << std::endl;
        }
    }

    // Map external variables to internal variables
    for (auto& container : function.containers()) {
        if (!function.is_internal(container)) {
            continue;
        }

        std::string external_name =
            container.substr(0, container.length() - external_suffix.length());
        this->main_stream_ << language_extension_.declaration(container, function.type(container));
        this->main_stream_ << " = "
                           << "&" << external_name;
        this->main_stream_ << ";" << std::endl;
    }

    for (size_t i = 0; i < schedule_.size(); i++) {
        if (i == 0) {
            this->main_stream_ << "if (" << language_extension_.expression(schedule_.condition(i))
                               << ") {\n";

            auto& function_i = schedule_.schedule(i).builder().subject();
            auto dispatcher = create_dispatcher(language_extension_, schedule_.schedule(i),
                                                function_i.root(), this->instrumented_);
            dispatcher->dispatch(this->main_stream_, this->globals_stream_, this->library_stream_);

            this->main_stream_ << "}\n";
        } else {
            this->main_stream_ << "else if ("
                               << language_extension_.expression(schedule_.condition(i)) << ") {\n";

            auto& function_i = schedule_.schedule(i).builder().subject();
            auto dispatcher = create_dispatcher(language_extension_, schedule_.schedule(i),
                                                function_i.root(), this->instrumented_);
            dispatcher->dispatch(this->main_stream_, this->globals_stream_, this->library_stream_);

            this->main_stream_ << "}\n";
        }
    }
};

}  // namespace codegen
}  // namespace sdfg

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

daisytuner / sdfglib / 15003172388

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