30543265435

Committed 30 Jul 2026 12:36PM UTC coverage: 64.549% (-0.2%) from 64.738%

Build # 30543265435

Build Type

Pull #913

github

Committed by

web-flow

Commit Message

Merge 6970d8cbe into 4a519a416

Pull Request Pull Request #913: delegates instrumentation event type decision to specific node dispat…

Coverage Stats

17 of 236 new or added lines in 25 files covered. (7.2%)

7 existing lines in 7 files now uncovered.

44347 of 68703 relevant lines covered (64.55%)

714.95 hits per line

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

0.0

/opt/src/targets/cuda/blas/dot.cpp

#include "sdfg/targets/cuda/blas/dot.h"
#include <sdfg/data_flow/library_nodes/math/blas/dot_node.h>
#include "sdfg/codegen/dispatchers/block_dispatcher.h"
#include "sdfg/data_flow/library_nodes/math/blas/blas_node.h"
#include "sdfg/targets/cuda/blas/utils.h"
#include "sdfg/targets/cuda/cuda.h"

namespace sdfg::cuda::blas {

DotNodeDispatcher_CUBLASWithTransfers::DotNodeDispatcher_CUBLASWithTransfers(
    codegen::LanguageExtension& language_extension,
    const Function& function,
    const data_flow::DataFlowGraph& data_flow_graph,
    const math::blas::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 math::blas::DotNode&>(this->node_);

    library_snippet_factory.add_global("#include <cuda.h>");
    library_snippet_factory.add_global("#include <cublas_v2.h>");

    std::string type, type2;
    switch (dot_node.precision()) {
        case sdfg::math::blas::BLAS_Precision::s:
            type = "float";
            type2 = "S";
            break;
        case sdfg::math::blas::BLAS_Precision::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 << "cudaError_t err_cuda;" << std::endl;
    stream << type << " *dx, *dy;" << std::endl;
    stream << "err_cuda = cudaMalloc((void**) &dx, " << x_size << ");" << std::endl;
    cuda_error_checking(stream, this->language_extension_, "err_cuda");
    stream << "err_cuda = cudaMalloc((void**) &dy, " << y_size << ");" << std::endl;
    cuda_error_checking(stream, this->language_extension_, "err_cuda");

    stream << "err_cuda = cudaMemcpy(dx, __x, " << x_size << ", cudaMemcpyHostToDevice);" << std::endl;
    cuda_error_checking(stream, this->language_extension_, "err_cuda");
    stream << "err_cuda = cudaMemcpy(dy, __y, " << y_size << ", cudaMemcpyHostToDevice);" << std::endl;
    cuda_error_checking(stream, this->language_extension_, "err_cuda");

    setup_blas_handle(library_snippet_factory, this->language_extension_);
    stream << "cublasStatus_t err;" << 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;
    cublas_error_checking(stream, this->language_extension_, "err");
    check_cuda_kernel_launch_errors(stream, this->language_extension_, false);

    stream << "err_cuda = cudaFree(dx);" << std::endl;
    cuda_error_checking(stream, this->language_extension_, "err_cuda");
    stream << "err_cuda = cudaFree(dy);" << std::endl;
    cuda_error_checking(stream, this->language_extension_, "err_cuda");
}

codegen::InstrumentationInfo DotNodeDispatcher_CUBLASWithTransfers::instrumentation_info() const {
    return {
        node_.element_id(),
        std::string(node_.element_type()) + ":::" + node_.code().value(),
        TargetType_CUDA,
        codegen::InstrumentationEventType::CUDA,
        analysis::LoopInfo{},
        {}
    };
}

DotNodeDispatcher_CUBLASWithoutTransfers::DotNodeDispatcher_CUBLASWithoutTransfers(
    codegen::LanguageExtension& language_extension,
    const Function& function,
    const data_flow::DataFlowGraph& data_flow_graph,
    const math::blas::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 math::blas::DotNode&>(this->node_);
    library_snippet_factory.add_global("#include <cuda.h>");
    library_snippet_factory.add_global("#include <cublas_v2.h>");

    setup_blas_handle(library_snippet_factory, this->language_extension_);

    stream << "cudaError_t err_cuda;" << std::endl;
    stream << "cublasStatus_t err;" << std::endl;

    stream << "err = cublas";
    switch (dot_node.precision()) {
        case sdfg::math::blas::BLAS_Precision::s:
            stream << "S";
            break;
        case sdfg::math::blas::BLAS_Precision::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;

    cublas_error_checking(stream, this->language_extension_, "err");
    check_cuda_kernel_launch_errors(stream, this->language_extension_, false);
}

codegen::InstrumentationInfo DotNodeDispatcher_CUBLASWithoutTransfers::instrumentation_info() const {
    return {
        node_.element_id(),
        std::string(node_.element_type()) + ":::" + node_.code().value(),
        TargetType_CUDA,
        codegen::InstrumentationEventType::CUDA,
        analysis::LoopInfo{},
        {}
    };
}

} // namespace sdfg::cuda::blas

1	#include "sdfg/targets/cuda/blas/dot.h"
2	#include <sdfg/data_flow/library_nodes/math/blas/dot_node.h>
3	#include "sdfg/codegen/dispatchers/block_dispatcher.h"
4	#include "sdfg/data_flow/library_nodes/math/blas/blas_node.h"
5	#include "sdfg/targets/cuda/blas/utils.h"
6	#include "sdfg/targets/cuda/cuda.h"
7
8	namespace sdfg::cuda::blas {
9
10	DotNodeDispatcher_CUBLASWithTransfers::DotNodeDispatcher_CUBLASWithTransfers(
11	codegen::LanguageExtension& language_extension,
12	const Function& function,
13	const data_flow::DataFlowGraph& data_flow_graph,
14	const math::blas::DotNode& node
15	)
16	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
17
18	void DotNodeDispatcher_CUBLASWithTransfers::dispatch_code(
19	codegen::PrettyPrinter& stream,
20	codegen::PrettyPrinter& globals_stream,
21	codegen::CodeSnippetFactory& library_snippet_factory
22	) {	×
23	auto& dot_node = static_cast<const math::blas::DotNode&>(this->node_);	×
24
25	library_snippet_factory.add_global("#include <cuda.h>");	×
26	library_snippet_factory.add_global("#include <cublas_v2.h>");	×
27
28	std::string type, type2;	×
29	switch (dot_node.precision()) {	×
30	case sdfg::math::blas::BLAS_Precision::s:	×
31	type = "float";	×
32	type2 = "S";	×
33	break;	×
34	case sdfg::math::blas::BLAS_Precision::d:	×
35	type = "double";	×
36	type2 = "D";	×
37	break;	×
38	default:	×
39	throw std::runtime_error("Invalid precision for CUBLAS DOT node");	×
40	}	×
41
42	const std::string x_size =	×
43	this->language_extension_.expression(	×
44	symbolic::add(symbolic::mul(symbolic::sub(dot_node.n(), symbolic::one()), dot_node.incx()), symbolic::one())	×
45	) +	×
46	" * sizeof(" + type + ")";	×
47	const std::string y_size =	×
48	this->language_extension_.expression(	×
49	symbolic::add(symbolic::mul(symbolic::sub(dot_node.n(), symbolic::one()), dot_node.incy()), symbolic::one())	×
50	) +	×
51	" * sizeof(" + type + ")";	×
52
53	stream << "cudaError_t err_cuda;" << std::endl;	×
54	stream << type << " dx, dy;" << std::endl;	×
55	stream << "err_cuda = cudaMalloc((void**) &dx, " << x_size << ");" << std::endl;	×
56	cuda_error_checking(stream, this->language_extension_, "err_cuda");	×
57	stream << "err_cuda = cudaMalloc((void**) &dy, " << y_size << ");" << std::endl;	×
58	cuda_error_checking(stream, this->language_extension_, "err_cuda");	×
59
60	stream << "err_cuda = cudaMemcpy(dx, __x, " << x_size << ", cudaMemcpyHostToDevice);" << std::endl;	×
61	cuda_error_checking(stream, this->language_extension_, "err_cuda");	×
62	stream << "err_cuda = cudaMemcpy(dy, __y, " << y_size << ", cudaMemcpyHostToDevice);" << std::endl;	×
63	cuda_error_checking(stream, this->language_extension_, "err_cuda");	×
64
65	setup_blas_handle(library_snippet_factory, this->language_extension_);	×
66	stream << "cublasStatus_t err;" << std::endl;	×
67
68	stream << "err = cublas" << type2 << "dot(handle, " << this->language_extension_.expression(dot_node.n())	×
69	<< ", dx, " << this->language_extension_.expression(dot_node.incx()) << ", dy, "	×
70	<< this->language_extension_.expression(dot_node.incy()) << ", &__out);" << std::endl;	×
71	cublas_error_checking(stream, this->language_extension_, "err");	×
72	check_cuda_kernel_launch_errors(stream, this->language_extension_, false);	×
73
74	stream << "err_cuda = cudaFree(dx);" << std::endl;	×
75	cuda_error_checking(stream, this->language_extension_, "err_cuda");	×
76	stream << "err_cuda = cudaFree(dy);" << std::endl;	×
77	cuda_error_checking(stream, this->language_extension_, "err_cuda");	×
78	}	×
79
NEW 80	codegen::InstrumentationInfo DotNodeDispatcher_CUBLASWithTransfers::instrumentation_info() const {	×
NEW 81	return {	×
NEW 82	node_.element_id(),	×
NEW 83	std::string(node_.element_type()) + ":::" + node_.code().value(),	×
NEW 84	TargetType_CUDA,	×
NEW 85	codegen::InstrumentationEventType::CUDA,	×
NEW 86	analysis::LoopInfo{},	×
NEW 87	{}	×
NEW 88	};	×
NEW 89	}	×
90
91	DotNodeDispatcher_CUBLASWithoutTransfers::DotNodeDispatcher_CUBLASWithoutTransfers(
92	codegen::LanguageExtension& language_extension,
93	const Function& function,
94	const data_flow::DataFlowGraph& data_flow_graph,
95	const math::blas::DotNode& node
96	)
97	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
98
99	void DotNodeDispatcher_CUBLASWithoutTransfers::dispatch_code(
100	codegen::PrettyPrinter& stream,
101	codegen::PrettyPrinter& globals_stream,
102	codegen::CodeSnippetFactory& library_snippet_factory
103	) {	×
104	auto& dot_node = static_cast<const math::blas::DotNode&>(this->node_);	×
105	library_snippet_factory.add_global("#include <cuda.h>");	×
106	library_snippet_factory.add_global("#include <cublas_v2.h>");	×
107
108	setup_blas_handle(library_snippet_factory, this->language_extension_);	×
109
110	stream << "cudaError_t err_cuda;" << std::endl;	×
111	stream << "cublasStatus_t err;" << std::endl;	×
112
113	stream << "err = cublas";	×
114	switch (dot_node.precision()) {	×
115	case sdfg::math::blas::BLAS_Precision::s:	×
116	stream << "S";	×
117	break;	×
118	case sdfg::math::blas::BLAS_Precision::d:	×
119	stream << "D";	×
120	break;	×
121	default:	×
122	throw std::runtime_error("Invalid precision for CUBLAS DOT node");	×
123	}	×
124	stream << "dot(handle, " << this->language_extension_.expression(dot_node.n()) << ", __x, "	×
125	<< this->language_extension_.expression(dot_node.incx()) << ", __y, "	×
126	<< this->language_extension_.expression(dot_node.incy()) << ", &__out);" << std::endl;	×
127
128	cublas_error_checking(stream, this->language_extension_, "err");	×
129	check_cuda_kernel_launch_errors(stream, this->language_extension_, false);	×
130	}	×
131
NEW 132	codegen::InstrumentationInfo DotNodeDispatcher_CUBLASWithoutTransfers::instrumentation_info() const {	×
NEW 133	return {	×
NEW 134	node_.element_id(),	×
NEW 135	std::string(node_.element_type()) + ":::" + node_.code().value(),	×
NEW 136	TargetType_CUDA,	×
NEW 137	codegen::InstrumentationEventType::CUDA,	×
NEW 138	analysis::LoopInfo{},	×
NEW 139	{}	×
NEW 140	};	×
NEW 141	}	×
142
143	} // namespace sdfg::cuda::blas

daisytuner / docc / 30543265435

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