15656007340

Committed 14 Jun 2025 08:51PM UTC coverage: 13.234% (-49.9%) from 63.144%

Build # 15656007340

Build Type

Pull #76

github

Committed by

web-flow

Commit Message

Merge 9586c8161 into 413c53212

Pull Request Pull Request #76: New Loop Dependency Analysis

Run Details

361 of 465 new or added lines in 7 files covered. (77.63%)

6215 existing lines in 110 files now uncovered.

1612 of 12181 relevant lines covered (13.23%)

13.64 hits per line

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

0.0

/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.h"
#include "sdfg/codegen/instrumentation/outermost_loops_instrumentation.h"

namespace sdfg {
namespace codegen {

CUDACodeGenerator::CUDACodeGenerator(StructuredSDFG& sdfg)
    : CodeGenerator(sdfg, InstrumentationStrategy::NONE) {
    if (sdfg.type() != FunctionType_NV_GLOBAL) {
        throw std::runtime_error("CUDACodeGenerator can only be used for GPU SDFGs");
    }
};

CUDACodeGenerator::CUDACodeGenerator(StructuredSDFG& sdfg,
                                     InstrumentationStrategy instrumentation_strategy)
    : CodeGenerator(sdfg, instrumentation_strategy) {
    if (sdfg.type() != FunctionType_NV_GLOBAL) {
        throw std::runtime_error("CUDACodeGenerator can only be used for GPU SDFGs");
    }
};

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,
                                  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 (instrumentation_strategy_ != InstrumentationStrategy::NONE)
        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() {
    // 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) {
            this->globals_stream_ << "extern " << language_extension_.declaration(container, type)
                                  << ";" << std::endl;
        }
        if (type.storage_type() == types::StorageType_NV_Constant) {
            assert(type.initializer().empty());
            this->globals_stream_ << "__constant__ "
                                  << language_extension_.declaration(container, 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;
        }
    }

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

        std::string external_name =
            container.substr(0, container.length() - external_suffix.length());
        this->main_stream_ << language_extension_.declaration(container, sdfg_.type(container));
        this->main_stream_ << " = "
                           << "&" << external_name;
        this->main_stream_ << ";" << 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;
        }
    }

    // Add instrumentation
    auto instrumentation = create_instrumentation(instrumentation_strategy_, sdfg_);

    auto dispatcher = create_dispatcher(language_extension_, sdfg_, sdfg_.root(), *instrumentation);
    dispatcher->dispatch(this->main_stream_, this->globals_stream_, this->library_stream_);
};

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

daisytuner / sdfglib / 15656007340

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