20764569418

Committed 06 Jan 2026 10:50PM UTC coverage: 62.168% (+21.4%) from 40.764%

Build # 20764569418

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #433 from daisytuner/clang-coverage

updates clang coverage flags

Run Details

14988 of 24109 relevant lines covered (62.17%)

88.57 hits per line

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

30.27

/src/data_flow/library_nodes/math/blas/dot_node.cpp

#include "sdfg/data_flow/library_nodes/math/blas/dot_node.h"
#include <stdexcept>
#include <string>

#include "sdfg/analysis/analysis.h"
#include "sdfg/builder/structured_sdfg_builder.h"

#include "sdfg/analysis/scope_analysis.h"
#include "sdfg/symbolic/symbolic.h"

namespace sdfg {
namespace math {
namespace blas {

DotNode::DotNode(
    size_t element_id,
    const DebugInfo& debug_info,
    const graph::Vertex vertex,
    data_flow::DataFlowGraph& parent,
    const data_flow::ImplementationType& implementation_type,
    const BLAS_Precision& precision,
    symbolic::Expression n,
    symbolic::Expression incx,
    symbolic::Expression incy
)
    : BLASNode(
          element_id,
          debug_info,
          vertex,
          parent,
          LibraryNodeType_DOT,
          {"__out"},
          {"__x", "__y"},
          implementation_type,
          precision
      ),
      n_(n), incx_(incx), incy_(incy) {}

symbolic::Expression DotNode::n() const { return this->n_; };

symbolic::Expression DotNode::incx() const { return this->incx_; };

symbolic::Expression DotNode::incy() const { return this->incy_; };

symbolic::SymbolSet DotNode::symbols() const {
    symbolic::SymbolSet syms;

    for (auto& atom : symbolic::atoms(this->n_)) {
        syms.insert(atom);
    }
    for (auto& atom : symbolic::atoms(this->incx_)) {
        syms.insert(atom);
    }
    for (auto& atom : symbolic::atoms(this->incy_)) {
        syms.insert(atom);
    }

    return syms;
};

void DotNode::replace(const symbolic::Expression old_expression, const symbolic::Expression new_expression) {
    this->n_ = symbolic::subs(this->n_, old_expression, new_expression);
    this->incx_ = symbolic::subs(this->incx_, old_expression, new_expression);
    this->incy_ = symbolic::subs(this->incy_, old_expression, new_expression);
};

void DotNode::validate(const Function& function) const {}

bool DotNode::expand(builder::StructuredSDFGBuilder& builder, analysis::AnalysisManager& analysis_manager) {
    auto& scope_analysis = analysis_manager.get<analysis::ScopeAnalysis>();

    auto& dataflow = this->get_parent();
    auto& block = static_cast<structured_control_flow::Block&>(*dataflow.get_parent());
    auto& parent = static_cast<structured_control_flow::Sequence&>(*scope_analysis.parent_scope(&block));
    int index = parent.index(block);
    auto& transition = parent.at(index).second;

    const data_flow::Memlet* iedge_x = nullptr;
    const data_flow::Memlet* iedge_y = nullptr;
    for (const auto& iedge : dataflow.in_edges(*this)) {
        if (iedge.dst_conn() == "__x") {
            iedge_x = &iedge;
        } else if (iedge.dst_conn() == "__y") {
            iedge_y = &iedge;
        }
    }

    const data_flow::Memlet* oedge_res = nullptr;
    for (const auto& oedge : dataflow.out_edges(*this)) {
        if (oedge.src_conn() == "__out") {
            oedge_res = &oedge;
            break;
        }
    }

    // Check if legal
    auto& input_node_x = static_cast<const data_flow::AccessNode&>(iedge_x->src());
    auto& input_node_y = static_cast<const data_flow::AccessNode&>(iedge_y->src());
    auto& output_node_res = static_cast<const data_flow::AccessNode&>(oedge_res->dst());
    if (dataflow.in_degree(input_node_x) != 0 || dataflow.in_degree(input_node_y) != 0 ||
        dataflow.out_degree(output_node_res) != 0) {
        return false;
    }

    auto& new_sequence = builder.add_sequence_before(parent, block, transition.assignments(), block.debug_info());

    std::string loop_var = builder.find_new_name("_i");
    builder.add_container(loop_var, types::Scalar(types::PrimitiveType::UInt64));

    auto loop_indvar = symbolic::symbol(loop_var);
    auto loop_init = symbolic::integer(0);
    auto loop_condition = symbolic::Lt(loop_indvar, this->n_);
    auto loop_update = symbolic::add(loop_indvar, symbolic::integer(1));

    auto& loop =
        builder.add_for(new_sequence, loop_indvar, loop_condition, loop_init, loop_update, {}, block.debug_info());
    auto& body = loop.root();

    auto& new_block = builder.add_block(body);

    auto& res_in = builder.add_access(new_block, output_node_res.data());
    auto& res_out = builder.add_access(new_block, output_node_res.data());
    auto& x = builder.add_access(new_block, input_node_x.data());
    auto& y = builder.add_access(new_block, input_node_y.data());

    auto& tasklet = builder.add_tasklet(new_block, data_flow::TaskletCode::fp_fma, "__out", {"_in1", "_in2", "_in3"});

    builder.add_computational_memlet(
        new_block,
        x,
        tasklet,
        "_in1",
        {symbolic::mul(loop_indvar, this->incx_)},
        iedge_x->base_type(),
        iedge_x->debug_info()
    );
    builder.add_computational_memlet(
        new_block,
        y,
        tasklet,
        "_in2",
        {symbolic::mul(loop_indvar, this->incy_)},
        iedge_y->base_type(),
        iedge_y->debug_info()
    );
    builder
        .add_computational_memlet(new_block, res_in, tasklet, "_in3", {}, oedge_res->base_type(), oedge_res->debug_info());
    builder.add_computational_memlet(
        new_block, tasklet, "__out", res_out, {}, oedge_res->base_type(), oedge_res->debug_info()
    );

    // Clean up
    builder.remove_memlet(block, *iedge_x);
    builder.remove_memlet(block, *iedge_y);
    builder.remove_memlet(block, *oedge_res);
    builder.remove_node(block, input_node_x);
    builder.remove_node(block, input_node_y);
    builder.remove_node(block, output_node_res);
    builder.remove_node(block, *this);
    builder.remove_child(parent, index + 1);

    return true;
}

symbolic::Expression DotNode::flop() const {
    auto muls = this->n_;
    auto adds = symbolic::sub(this->n_, symbolic::one());
    return symbolic::add(muls, adds);
}

std::unique_ptr<data_flow::DataFlowNode> DotNode::
    clone(size_t element_id, const graph::Vertex vertex, data_flow::DataFlowGraph& parent) const {
    auto node_clone = std::unique_ptr<DotNode>(new DotNode(
        element_id,
        this->debug_info(),
        vertex,
        parent,
        this->implementation_type_,
        this->precision_,
        this->n_,
        this->incx_,
        this->incy_
    ));
    return std::move(node_clone);
}

nlohmann::json DotNodeSerializer::serialize(const data_flow::LibraryNode& library_node) {
    const DotNode& gemm_node = static_cast<const DotNode&>(library_node);
    nlohmann::json j;

    serializer::JSONSerializer serializer;
    j["code"] = gemm_node.code().value();
    j["precision"] = gemm_node.precision();
    j["n"] = serializer.expression(gemm_node.n());
    j["incx"] = serializer.expression(gemm_node.incx());
    j["incy"] = serializer.expression(gemm_node.incy());

    return j;
}

data_flow::LibraryNode& DotNodeSerializer::deserialize(
    const nlohmann::json& j, builder::StructuredSDFGBuilder& builder, structured_control_flow::Block& parent
) {
    // Assertions for required fields
    assert(j.contains("element_id"));
    assert(j.contains("code"));
    assert(j.contains("debug_info"));

    auto code = j["code"].get<std::string>();
    if (code != LibraryNodeType_DOT.value()) {
        throw std::runtime_error("Invalid library node code");
    }

    // Extract debug info using JSONSerializer
    sdfg::serializer::JSONSerializer serializer;
    DebugInfo debug_info = serializer.json_to_debug_info(j["debug_info"]);

    auto precision = j.at("precision").get<BLAS_Precision>();
    auto n = symbolic::parse(j.at("n"));
    auto incx = symbolic::parse(j.at("incx"));
    auto incy = symbolic::parse(j.at("incy"));

    auto implementation_type = j.at("implementation_type").get<std::string>();

    return builder.add_library_node<DotNode>(parent, debug_info, implementation_type, precision, n, incx, incy);
}

DotNodeDispatcher_BLAS::DotNodeDispatcher_BLAS(
    codegen::LanguageExtension& language_extension,
    const Function& function,
    const data_flow::DataFlowGraph& data_flow_graph,
    const DotNode& node
)
    : codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}

void DotNodeDispatcher_BLAS::dispatch_code(
    codegen::PrettyPrinter& stream,
    codegen::PrettyPrinter& globals_stream,
    codegen::CodeSnippetFactory& library_snippet_factory
) {
    stream << "{" << std::endl;
    stream.setIndent(stream.indent() + 4);

    auto& dot_node = static_cast<const DotNode&>(this->node_);

    sdfg::types::Scalar base_type(types::PrimitiveType::Void);
    BLAS_Precision precision = dot_node.precision();
    switch (precision) {
        case BLAS_Precision::h:
            base_type = types::Scalar(types::PrimitiveType::Half);
            break;
        case BLAS_Precision::s:
            base_type = types::Scalar(types::PrimitiveType::Float);
            break;
        case BLAS_Precision::d:
            base_type = types::Scalar(types::PrimitiveType::Double);
            break;
        default:
            throw std::runtime_error("Invalid BLAS_Precision value");
    }

    stream << dot_node.outputs().at(0) << " = ";
    stream << "cblas_" << BLAS_Precision_to_string(precision) << "dot(";
    stream.setIndent(stream.indent() + 4);
    stream << this->language_extension_.expression(dot_node.n());
    stream << ", ";
    stream << dot_node.inputs().at(0);
    stream << ", ";
    stream << this->language_extension_.expression(dot_node.incx());
    stream << ", ";
    stream << dot_node.inputs().at(1);
    stream << ", ";
    stream << this->language_extension_.expression(dot_node.incy());
    stream.setIndent(stream.indent() - 4);
    stream << ");" << std::endl;

    stream.setIndent(stream.indent() - 4);
    stream << "}" << std::endl;
}

DotNodeDispatcher_CUBLASWithTransfers::DotNodeDispatcher_CUBLASWithTransfers(
    codegen::LanguageExtension& language_extension,
    const Function& function,
    const data_flow::DataFlowGraph& data_flow_graph,
    const DotNode& node
)
    : codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}

void DotNodeDispatcher_CUBLASWithTransfers::dispatch_code(
    codegen::PrettyPrinter& stream,
    codegen::PrettyPrinter& globals_stream,
    codegen::CodeSnippetFactory& library_snippet_factory
) {
    auto& dot_node = static_cast<const DotNode&>(this->node_);

    globals_stream << "#include <cuda.h>" << std::endl;
    globals_stream << "#include <cublas_v2.h>" << std::endl;

    std::string type, type2;
    switch (dot_node.precision()) {
        case s:
            type = "float";
            type2 = "S";
            break;
        case d:
            type = "double";
            type2 = "D";
            break;
        default:
            throw std::runtime_error("Invalid precision for CUBLAS DOT node");
    }

    const std::string x_size =
        this->language_extension_.expression(
            symbolic::add(symbolic::mul(symbolic::sub(dot_node.n(), symbolic::one()), dot_node.incx()), symbolic::one())
        ) +
        " * sizeof(" + type + ")";
    const std::string y_size =
        this->language_extension_.expression(
            symbolic::add(symbolic::mul(symbolic::sub(dot_node.n(), symbolic::one()), dot_node.incy()), symbolic::one())
        ) +
        " * sizeof(" + type + ")";

    stream << type << " *dx, *dy;" << std::endl;
    stream << "cudaMalloc(&dx, " << x_size << ");" << std::endl;
    stream << "cudaMalloc(&dy, " << y_size << ");" << std::endl;

    stream << "cudaMemcpy(dx, __x, " << x_size << ", cudaMemcpyHostToDevice);" << std::endl;
    stream << "cudaMemcpy(dy, __y, " << y_size << ", cudaMemcpyHostToDevice);" << std::endl;

    stream << "cublasStatus_t err;" << std::endl;
    stream << "cublasHandle_t handle;" << std::endl;
    stream << "err = cublasCreate(&handle);" << std::endl;
    stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;
    stream.setIndent(stream.indent() + 4);
    stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;
    stream.setIndent(stream.indent() - 4);
    stream << "}" << std::endl;
    stream << "err = cublas" << type2 << "dot(handle, " << this->language_extension_.expression(dot_node.n())
           << ", dx, " << this->language_extension_.expression(dot_node.incx()) << ", dy, "
           << this->language_extension_.expression(dot_node.incy()) << ", &__out);" << std::endl;
    stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;
    stream.setIndent(stream.indent() + 4);
    stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;
    stream.setIndent(stream.indent() - 4);
    stream << "}" << std::endl;
    stream << "err = cublasDestroy(handle);" << std::endl;
    stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;
    stream.setIndent(stream.indent() + 4);
    stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;
    stream.setIndent(stream.indent() - 4);
    stream << "}" << std::endl;

    stream << "cudaFree(dx);" << std::endl;
    stream << "cudaFree(dy);" << std::endl;
}

DotNodeDispatcher_CUBLASWithoutTransfers::DotNodeDispatcher_CUBLASWithoutTransfers(
    codegen::LanguageExtension& language_extension,
    const Function& function,
    const data_flow::DataFlowGraph& data_flow_graph,
    const DotNode& node
)
    : codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}

void DotNodeDispatcher_CUBLASWithoutTransfers::dispatch_code(
    codegen::PrettyPrinter& stream,
    codegen::PrettyPrinter& globals_stream,
    codegen::CodeSnippetFactory& library_snippet_factory
) {
    auto& dot_node = static_cast<const DotNode&>(this->node_);

    globals_stream << "#include <cuda.h>" << std::endl;
    globals_stream << "#include <cublas_v2.h>" << std::endl;

    stream << "cublasStatus_t err;" << std::endl;
    stream << "cublasHandle_t handle;" << std::endl;
    stream << "err = cublasCreate(&handle);" << std::endl;
    stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;
    stream.setIndent(stream.indent() + 4);
    stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;
    stream.setIndent(stream.indent() - 4);
    stream << "}" << std::endl;
    stream << "err = cublas";
    switch (dot_node.precision()) {
        case s:
            stream << "S";
            break;
        case d:
            stream << "D";
            break;
        default:
            throw std::runtime_error("Invalid precision for CUBLAS DOT node");
    }
    stream << "dot(handle, " << this->language_extension_.expression(dot_node.n()) << ", __x, "
           << this->language_extension_.expression(dot_node.incx()) << ", __y, "
           << this->language_extension_.expression(dot_node.incy()) << ", &__out);" << std::endl;
    stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;
    stream.setIndent(stream.indent() + 4);
    stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;
    stream.setIndent(stream.indent() - 4);
    stream << "}" << std::endl;
    stream << "err = cublasDestroy(handle);" << std::endl;
    stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;
    stream.setIndent(stream.indent() + 4);
    stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;
    stream.setIndent(stream.indent() - 4);
    stream << "}" << std::endl;
}

} // namespace blas
} // namespace math
} // namespace sdfg

1	#include "sdfg/data_flow/library_nodes/math/blas/dot_node.h"
2	#include <stdexcept>
3	#include <string>
4
5	#include "sdfg/analysis/analysis.h"
6	#include "sdfg/builder/structured_sdfg_builder.h"
7
8	#include "sdfg/analysis/scope_analysis.h"
9	#include "sdfg/symbolic/symbolic.h"
10
11	namespace sdfg {
12	namespace math {
13	namespace blas {
14
15	DotNode::DotNode(
16	size_t element_id,
17	const DebugInfo& debug_info,
18	const graph::Vertex vertex,
19	data_flow::DataFlowGraph& parent,
20	const data_flow::ImplementationType& implementation_type,
21	const BLAS_Precision& precision,
22	symbolic::Expression n,
23	symbolic::Expression incx,
24	symbolic::Expression incy
25	)
26	: BLASNode(	1✔
27	element_id,	1✔
28	debug_info,	1✔
29	vertex,	1✔
30	parent,	1✔
31	LibraryNodeType_DOT,	1✔
32	{"__out"},	1✔
33	{"__x", "__y"},	1✔
34	implementation_type,	1✔
35	precision	1✔
36	),	1✔
37	n_(n), incx_(incx), incy_(incy) {}	1✔
38
39	symbolic::Expression DotNode::n() const { return this->n_; };	×
40
41	symbolic::Expression DotNode::incx() const { return this->incx_; };	×
42
43	symbolic::Expression DotNode::incy() const { return this->incy_; };	×
44
45	symbolic::SymbolSet DotNode::symbols() const {	×
46	symbolic::SymbolSet syms;	×
47
48	for (auto& atom : symbolic::atoms(this->n_)) {	×
49	syms.insert(atom);	×
50	}	×
51	for (auto& atom : symbolic::atoms(this->incx_)) {	×
52	syms.insert(atom);	×
53	}	×
54	for (auto& atom : symbolic::atoms(this->incy_)) {	×
55	syms.insert(atom);	×
56	}	×
57
58	return syms;	×
59	};	×
60
61	void DotNode::replace(const symbolic::Expression old_expression, const symbolic::Expression new_expression) {	×
62	this->n_ = symbolic::subs(this->n_, old_expression, new_expression);	×
63	this->incx_ = symbolic::subs(this->incx_, old_expression, new_expression);	×
64	this->incy_ = symbolic::subs(this->incy_, old_expression, new_expression);	×
65	};	×
66
67	void DotNode::validate(const Function& function) const {}	×
68
69	bool DotNode::expand(builder::StructuredSDFGBuilder& builder, analysis::AnalysisManager& analysis_manager) {	1✔
70	auto& scope_analysis = analysis_manager.get<analysis::ScopeAnalysis>();	1✔
71
72	auto& dataflow = this->get_parent();	1✔
73	auto& block = static_cast<structured_control_flow::Block&>(*dataflow.get_parent());	1✔
74	auto& parent = static_cast<structured_control_flow::Sequence&>(*scope_analysis.parent_scope(&block));	1✔
75	int index = parent.index(block);	1✔
76	auto& transition = parent.at(index).second;	1✔
77
78	const data_flow::Memlet* iedge_x = nullptr;	1✔
79	const data_flow::Memlet* iedge_y = nullptr;	1✔
80	for (const auto& iedge : dataflow.in_edges(*this)) {	2✔
81	if (iedge.dst_conn() == "__x") {	2✔
82	iedge_x = &iedge;	1✔
83	} else if (iedge.dst_conn() == "__y") {	1✔
84	iedge_y = &iedge;	1✔
85	}	1✔
86	}	2✔
87
88	const data_flow::Memlet* oedge_res = nullptr;	1✔
89	for (const auto& oedge : dataflow.out_edges(*this)) {	1✔
90	if (oedge.src_conn() == "__out") {	1✔
91	oedge_res = &oedge;	1✔
92	break;	1✔
93	}	1✔
94	}	1✔
95
96	// Check if legal
97	auto& input_node_x = static_cast<const data_flow::AccessNode&>(iedge_x->src());	1✔
98	auto& input_node_y = static_cast<const data_flow::AccessNode&>(iedge_y->src());	1✔
99	auto& output_node_res = static_cast<const data_flow::AccessNode&>(oedge_res->dst());	1✔
100	if (dataflow.in_degree(input_node_x) != 0 \|\| dataflow.in_degree(input_node_y) != 0 \|\|	1✔
101	dataflow.out_degree(output_node_res) != 0) {	1✔
102	return false;	×
103	}	×
104
105	auto& new_sequence = builder.add_sequence_before(parent, block, transition.assignments(), block.debug_info());	1✔
106
107	std::string loop_var = builder.find_new_name("_i");	1✔
108	builder.add_container(loop_var, types::Scalar(types::PrimitiveType::UInt64));	1✔
109
110	auto loop_indvar = symbolic::symbol(loop_var);	1✔
111	auto loop_init = symbolic::integer(0);	1✔
112	auto loop_condition = symbolic::Lt(loop_indvar, this->n_);	1✔
113	auto loop_update = symbolic::add(loop_indvar, symbolic::integer(1));	1✔
114
115	auto& loop =	1✔
116	builder.add_for(new_sequence, loop_indvar, loop_condition, loop_init, loop_update, {}, block.debug_info());	1✔
117	auto& body = loop.root();	1✔
118
119	auto& new_block = builder.add_block(body);	1✔
120
121	auto& res_in = builder.add_access(new_block, output_node_res.data());	1✔
122	auto& res_out = builder.add_access(new_block, output_node_res.data());	1✔
123	auto& x = builder.add_access(new_block, input_node_x.data());	1✔
124	auto& y = builder.add_access(new_block, input_node_y.data());	1✔
125
126	auto& tasklet = builder.add_tasklet(new_block, data_flow::TaskletCode::fp_fma, "__out", {"_in1", "_in2", "_in3"});	1✔
127
128	builder.add_computational_memlet(	1✔
129	new_block,	1✔
130	x,	1✔
131	tasklet,	1✔
132	"_in1",	1✔
133	{symbolic::mul(loop_indvar, this->incx_)},	1✔
134	iedge_x->base_type(),	1✔
135	iedge_x->debug_info()	1✔
136	);	1✔
137	builder.add_computational_memlet(	1✔
138	new_block,	1✔
139	y,	1✔
140	tasklet,	1✔
141	"_in2",	1✔
142	{symbolic::mul(loop_indvar, this->incy_)},	1✔
143	iedge_y->base_type(),	1✔
144	iedge_y->debug_info()	1✔
145	);	1✔
146	builder	1✔
147	.add_computational_memlet(new_block, res_in, tasklet, "_in3", {}, oedge_res->base_type(), oedge_res->debug_info());	1✔
148	builder.add_computational_memlet(	1✔
149	new_block, tasklet, "__out", res_out, {}, oedge_res->base_type(), oedge_res->debug_info()	1✔
150	);	1✔
151
152	// Clean up
153	builder.remove_memlet(block, *iedge_x);	1✔
154	builder.remove_memlet(block, *iedge_y);	1✔
155	builder.remove_memlet(block, *oedge_res);	1✔
156	builder.remove_node(block, input_node_x);	1✔
157	builder.remove_node(block, input_node_y);	1✔
158	builder.remove_node(block, output_node_res);	1✔
159	builder.remove_node(block, *this);	1✔
160	builder.remove_child(parent, index + 1);	1✔
161
162	return true;	1✔
163	}	1✔
164
165	symbolic::Expression DotNode::flop() const {	×
166	auto muls = this->n_;	×
167	auto adds = symbolic::sub(this->n_, symbolic::one());	×
168	return symbolic::add(muls, adds);	×
169	}	×
170
171	std::unique_ptr<data_flow::DataFlowNode> DotNode::
172	clone(size_t element_id, const graph::Vertex vertex, data_flow::DataFlowGraph& parent) const {	×
173	auto node_clone = std::unique_ptr<DotNode>(new DotNode(	×
174	element_id,	×
175	this->debug_info(),	×
176	vertex,	×
177	parent,	×
178	this->implementation_type_,	×
179	this->precision_,	×
180	this->n_,	×
181	this->incx_,	×
182	this->incy_	×
183	));	×
184	return std::move(node_clone);	×
185	}	×
186
187	nlohmann::json DotNodeSerializer::serialize(const data_flow::LibraryNode& library_node) {	×
188	const DotNode& gemm_node = static_cast<const DotNode&>(library_node);	×
189	nlohmann::json j;	×
190
191	serializer::JSONSerializer serializer;	×
192	j["code"] = gemm_node.code().value();	×
193	j["precision"] = gemm_node.precision();	×
194	j["n"] = serializer.expression(gemm_node.n());	×
195	j["incx"] = serializer.expression(gemm_node.incx());	×
196	j["incy"] = serializer.expression(gemm_node.incy());	×
197
198	return j;	×
199	}	×
200
201	data_flow::LibraryNode& DotNodeSerializer::deserialize(
202	const nlohmann::json& j, builder::StructuredSDFGBuilder& builder, structured_control_flow::Block& parent
203	) {	×
204	// Assertions for required fields
205	assert(j.contains("element_id"));	×
206	assert(j.contains("code"));	×
207	assert(j.contains("debug_info"));	×
208
209	auto code = j["code"].get<std::string>();	×
210	if (code != LibraryNodeType_DOT.value()) {	×
211	throw std::runtime_error("Invalid library node code");	×
212	}	×
213
214	// Extract debug info using JSONSerializer
215	sdfg::serializer::JSONSerializer serializer;	×
216	DebugInfo debug_info = serializer.json_to_debug_info(j["debug_info"]);	×
217
218	auto precision = j.at("precision").get<BLAS_Precision>();	×
219	auto n = symbolic::parse(j.at("n"));	×
220	auto incx = symbolic::parse(j.at("incx"));	×
221	auto incy = symbolic::parse(j.at("incy"));	×
222
223	auto implementation_type = j.at("implementation_type").get<std::string>();	×
224
225	return builder.add_library_node<DotNode>(parent, debug_info, implementation_type, precision, n, incx, incy);	×
226	}	×
227
228	DotNodeDispatcher_BLAS::DotNodeDispatcher_BLAS(
229	codegen::LanguageExtension& language_extension,
230	const Function& function,
231	const data_flow::DataFlowGraph& data_flow_graph,
232	const DotNode& node
233	)
234	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
235
236	void DotNodeDispatcher_BLAS::dispatch_code(
237	codegen::PrettyPrinter& stream,
238	codegen::PrettyPrinter& globals_stream,
239	codegen::CodeSnippetFactory& library_snippet_factory
240	) {	×
241	stream << "{" << std::endl;	×
242	stream.setIndent(stream.indent() + 4);	×
243
244	auto& dot_node = static_cast<const DotNode&>(this->node_);	×
245
246	sdfg::types::Scalar base_type(types::PrimitiveType::Void);	×
247	BLAS_Precision precision = dot_node.precision();	×
248	switch (precision) {	×
249	case BLAS_Precision::h:	×
250	base_type = types::Scalar(types::PrimitiveType::Half);	×
251	break;	×
252	case BLAS_Precision::s:	×
253	base_type = types::Scalar(types::PrimitiveType::Float);	×
254	break;	×
255	case BLAS_Precision::d:	×
256	base_type = types::Scalar(types::PrimitiveType::Double);	×
257	break;	×
258	default:	×
259	throw std::runtime_error("Invalid BLAS_Precision value");	×
260	}	×
261
262	stream << dot_node.outputs().at(0) << " = ";	×
263	stream << "cblas_" << BLAS_Precision_to_string(precision) << "dot(";	×
264	stream.setIndent(stream.indent() + 4);	×
265	stream << this->language_extension_.expression(dot_node.n());	×
266	stream << ", ";	×
267	stream << dot_node.inputs().at(0);	×
268	stream << ", ";	×
269	stream << this->language_extension_.expression(dot_node.incx());	×
270	stream << ", ";	×
271	stream << dot_node.inputs().at(1);	×
272	stream << ", ";	×
273	stream << this->language_extension_.expression(dot_node.incy());	×
274	stream.setIndent(stream.indent() - 4);	×
275	stream << ");" << std::endl;	×
276
277	stream.setIndent(stream.indent() - 4);	×
278	stream << "}" << std::endl;	×
279	}	×
280
281	DotNodeDispatcher_CUBLASWithTransfers::DotNodeDispatcher_CUBLASWithTransfers(
282	codegen::LanguageExtension& language_extension,
283	const Function& function,
284	const data_flow::DataFlowGraph& data_flow_graph,
285	const DotNode& node
286	)
287	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
288
289	void DotNodeDispatcher_CUBLASWithTransfers::dispatch_code(
290	codegen::PrettyPrinter& stream,
291	codegen::PrettyPrinter& globals_stream,
292	codegen::CodeSnippetFactory& library_snippet_factory
293	) {	×
294	auto& dot_node = static_cast<const DotNode&>(this->node_);	×
295
296	globals_stream << "#include <cuda.h>" << std::endl;	×
297	globals_stream << "#include <cublas_v2.h>" << std::endl;	×
298
299	std::string type, type2;	×
300	switch (dot_node.precision()) {	×
301	case s:	×
302	type = "float";	×
303	type2 = "S";	×
304	break;	×
305	case d:	×
306	type = "double";	×
307	type2 = "D";	×
308	break;	×
309	default:	×
310	throw std::runtime_error("Invalid precision for CUBLAS DOT node");	×
311	}	×
312
313	const std::string x_size =	×
314	this->language_extension_.expression(	×
315	symbolic::add(symbolic::mul(symbolic::sub(dot_node.n(), symbolic::one()), dot_node.incx()), symbolic::one())	×
316	) +	×
317	" * sizeof(" + type + ")";	×
318	const std::string y_size =	×
319	this->language_extension_.expression(	×
320	symbolic::add(symbolic::mul(symbolic::sub(dot_node.n(), symbolic::one()), dot_node.incy()), symbolic::one())	×
321	) +	×
322	" * sizeof(" + type + ")";	×
323
324	stream << type << " dx, dy;" << std::endl;	×
325	stream << "cudaMalloc(&dx, " << x_size << ");" << std::endl;	×
326	stream << "cudaMalloc(&dy, " << y_size << ");" << std::endl;	×
327
328	stream << "cudaMemcpy(dx, __x, " << x_size << ", cudaMemcpyHostToDevice);" << std::endl;	×
329	stream << "cudaMemcpy(dy, __y, " << y_size << ", cudaMemcpyHostToDevice);" << std::endl;	×
330
331	stream << "cublasStatus_t err;" << std::endl;	×
332	stream << "cublasHandle_t handle;" << std::endl;	×
333	stream << "err = cublasCreate(&handle);" << std::endl;	×
334	stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;	×
335	stream.setIndent(stream.indent() + 4);	×
336	stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;	×
337	stream.setIndent(stream.indent() - 4);	×
338	stream << "}" << std::endl;	×
339	stream << "err = cublas" << type2 << "dot(handle, " << this->language_extension_.expression(dot_node.n())	×
340	<< ", dx, " << this->language_extension_.expression(dot_node.incx()) << ", dy, "	×
341	<< this->language_extension_.expression(dot_node.incy()) << ", &__out);" << std::endl;	×
342	stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;	×
343	stream.setIndent(stream.indent() + 4);	×
344	stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;	×
345	stream.setIndent(stream.indent() - 4);	×
346	stream << "}" << std::endl;	×
347	stream << "err = cublasDestroy(handle);" << std::endl;	×
348	stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;	×
349	stream.setIndent(stream.indent() + 4);	×
350	stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;	×
351	stream.setIndent(stream.indent() - 4);	×
352	stream << "}" << std::endl;	×
353
354	stream << "cudaFree(dx);" << std::endl;	×
355	stream << "cudaFree(dy);" << std::endl;	×
356	}	×
357
358	DotNodeDispatcher_CUBLASWithoutTransfers::DotNodeDispatcher_CUBLASWithoutTransfers(
359	codegen::LanguageExtension& language_extension,
360	const Function& function,
361	const data_flow::DataFlowGraph& data_flow_graph,
362	const DotNode& node
363	)
364	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
365
366	void DotNodeDispatcher_CUBLASWithoutTransfers::dispatch_code(
367	codegen::PrettyPrinter& stream,
368	codegen::PrettyPrinter& globals_stream,
369	codegen::CodeSnippetFactory& library_snippet_factory
370	) {	×
371	auto& dot_node = static_cast<const DotNode&>(this->node_);	×
372
373	globals_stream << "#include <cuda.h>" << std::endl;	×
374	globals_stream << "#include <cublas_v2.h>" << std::endl;	×
375
376	stream << "cublasStatus_t err;" << std::endl;	×
377	stream << "cublasHandle_t handle;" << std::endl;	×
378	stream << "err = cublasCreate(&handle);" << std::endl;	×
379	stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;	×
380	stream.setIndent(stream.indent() + 4);	×
381	stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;	×
382	stream.setIndent(stream.indent() - 4);	×
383	stream << "}" << std::endl;	×
384	stream << "err = cublas";	×
385	switch (dot_node.precision()) {	×
386	case s:	×
387	stream << "S";	×
388	break;	×
389	case d:	×
390	stream << "D";	×
391	break;	×
392	default:	×
393	throw std::runtime_error("Invalid precision for CUBLAS DOT node");	×
394	}	×
395	stream << "dot(handle, " << this->language_extension_.expression(dot_node.n()) << ", __x, "	×
396	<< this->language_extension_.expression(dot_node.incx()) << ", __y, "	×
397	<< this->language_extension_.expression(dot_node.incy()) << ", &__out);" << std::endl;	×
398	stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;	×
399	stream.setIndent(stream.indent() + 4);	×
400	stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;	×
401	stream.setIndent(stream.indent() - 4);	×
402	stream << "}" << std::endl;	×
403	stream << "err = cublasDestroy(handle);" << std::endl;	×
404	stream << "if (err != CUBLAS_STATUS_SUCCESS) {" << std::endl;	×
405	stream.setIndent(stream.indent() + 4);	×
406	stream << this->language_extension_.external_prefix() << "exit(1);" << std::endl;	×
407	stream.setIndent(stream.indent() - 4);	×
408	stream << "}" << std::endl;	×
409	}	×
410
411	} // namespace blas
412	} // namespace math
413	} // namespace sdfg

daisytuner / sdfglib / 20764569418

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