16779684622

Committed 06 Aug 2025 02:21PM UTC coverage: 64.3% (-1.0%) from 65.266%

Build # 16779684622

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #172 from daisytuner/opaque-pointers

Opaque pointers, typed memlets, untyped tasklet connectors

Run Details

330 of 462 new or added lines in 38 files covered. (71.43%)

382 existing lines in 30 files now uncovered.

8865 of 13787 relevant lines covered (64.3%)

116.73 hits per line

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

66.15

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

namespace sdfg {
namespace codegen {

CUDACodeGenerator::CUDACodeGenerator(
    StructuredSDFG& sdfg,
    InstrumentationPlan& instrumentation_plan,
    bool capture_args_results,
    const std::pair<std::filesystem::path, std::filesystem::path>* output_and_header_paths
)
    : CodeGenerator(sdfg, instrumentation_plan, capture_args_results, output_and_header_paths),
      language_extension_(sdfg.externals()) {
    if (sdfg.type() != FunctionType_NV_GLOBAL) {
        throw std::runtime_error("CUDACodeGenerator can only be used for GPU SDFGs");
    }
    if (capture_args_results) {
        std::cerr << "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::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::ofstream& 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->includes_stream_ << "#define "
                           << "__DAISY_NVVM__" << std::endl;
    this->includes_stream_ << "#include <daisy_rtl/daisy_rtl.h>" << std::endl;
};

void CUDACodeGenerator::dispatch_structures() {
    // Forward declarations
    for (auto& structure : sdfg_.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 : 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_ << "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, false, true);
            this->classes_stream_ << ";" << std::endl;
        }

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

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

daisytuner / sdfglib / 16779684622

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