27876065887

Committed 20 Jun 2026 03:46PM UTC coverage: 61.765% (-0.07%) from 61.832%

Build # 27876065887

Build Type

Pull #789

github

Committed by

web-flow

Commit Message

Merge 79a38c1d2 into b7103c21a

Pull Request Pull Request #789: Device Residency

Coverage Stats

81 of 204 new or added lines in 6 files covered. (39.71%)

2 existing lines in 2 files now uncovered.

37026 of 59947 relevant lines covered (61.76%)

1015.8 hits per line

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

49.21

/opt/src/passes/offloading/cuda_library_node_rewriter_pass.cpp

#include "sdfg/passes/offloading/cuda_library_node_rewriter_pass.h"
#include <optional>

#include "sdfg/data_flow/library_node.h"
#include "sdfg/data_flow/library_nodes/math/math.h"
#include "sdfg/data_flow/library_nodes/stdlib/memcpy.h"
#include "sdfg/data_flow/library_nodes/stdlib/memset.h"
#include "sdfg/structured_sdfg.h"
#include "sdfg/symbolic/symbolic.h"
#include "sdfg/targets/cuda/cuda.h"

namespace sdfg {
namespace cuda {

std::optional<data_flow::ImplementationType> CudaLibraryNodeRewriter::
    try_library_node_implementation(const data_flow::LibraryNode& lib_node, types::PrimitiveType data_type) {
    if (data_type == types::PrimitiveType::Float || data_type == types::PrimitiveType::Double) {
        if (lib_node.code() == math::blas::LibraryNodeType_GEMM.value()) {
            auto& gemm_node = static_cast<const math::blas::GEMMNode&>(lib_node);
            return try_cublas_gemm_node_implementation(gemm_node, data_type);
        } else if (lib_node.code() == math::blas::LibraryNodeType_BatchedGEMM.value()) {
            return cuda::ImplementationType_CUDAWithTransfers;
        } else if (lib_node.code() == math::blas::LibraryNodeType_DOT.value()) {
            return cuda::ImplementationType_CUDAWithTransfers;
        } else if (lib_node.code() == math::blas::LibraryNodeType_BatchedGEMM.value()) {
            auto& batched_gemm_node = static_cast<const math::blas::BatchedGEMMNode&>(lib_node);
            return try_cublas_batched_gemm_node_implementation(batched_gemm_node, data_type);
        } else {
            return std::nullopt;
        }
    } else {
        return std::nullopt;
    }
}

std::optional<data_flow::ImplementationType> CudaLibraryNodeRewriter::
    try_cublas_gemm_node_implementation(const math::blas::GEMMNode& gemm_node, types::PrimitiveType data_type) {
    return cuda::ImplementationType_CUDAWithTransfers;
}

std::optional<data_flow::ImplementationType> CudaLibraryNodeRewriter::
    try_memset_implementation(const ::sdfg::stdlib::MemsetNode& memset_node) {
    return cuda::ImplementationType_CUDAWithTransfers;
}

std::optional<data_flow::ImplementationType> CudaLibraryNodeRewriter::
    try_memcpy_implementation(const ::sdfg::stdlib::MemcpyNode& memcpy_node) {
    return cuda::ImplementationType_CUDAWithTransfers;
}

std::optional<data_flow::ImplementationType> CudaLibraryNodeRewriter::try_cublas_batched_gemm_node_implementation(
    const math::blas::BatchedGEMMNode& batched_gemm_node, types::PrimitiveType data_type
) {
    // Heuristic: Avoid using CUBLAS for very small matrix multiplications
    auto m = batched_gemm_node.m();
    auto n = batched_gemm_node.n();
    auto k = batched_gemm_node.k();
    auto size = symbolic::mul(symbolic::mul(m, n), k);
    if (symbolic::eq(size, symbolic::one())) {
        return std::nullopt;
    }
    return cuda::ImplementationType_CUDAWithTransfers;
}

CudaLibraryNodeRewriter::
    CudaLibraryNodeRewriter(builder::StructuredSDFGBuilder& builder, analysis::AnalysisManager& analysis_manager)
    : visitor::StructuredSDFGVisitor(builder, analysis_manager) {}

bool CudaLibraryNodeRewriter::accept(structured_control_flow::Block& node) {
    auto& dataflow = node.dataflow();
    for (auto& library_node : dataflow.nodes()) {
        if (auto lib_node = dynamic_cast<math::blas::BLASNode*>(&library_node)) {
            auto implType = try_library_node_implementation(*lib_node, lib_node->scalar_primitive());

            if (implType) {
                lib_node->implementation_type() = implType.value();
            }
        }
        if (auto memset_node = dynamic_cast<::sdfg::stdlib::MemsetNode*>(&library_node)) {
            auto implType = try_memset_implementation(*memset_node);
            if (implType) {
                memset_node->implementation_type() = implType.value();
            }
        }
        if (auto memcpy_node = dynamic_cast<::sdfg::stdlib::MemcpyNode*>(&library_node)) {
            auto implType = try_memcpy_implementation(*memcpy_node);
            if (implType) {
                memcpy_node->implementation_type() = implType.value();
            }
        }
    }
    return false;
}

} // namespace cuda
} // namespace sdfg

1	#include "sdfg/passes/offloading/cuda_library_node_rewriter_pass.h"
2	#include <optional>
3
4	#include "sdfg/data_flow/library_node.h"
5	#include "sdfg/data_flow/library_nodes/math/math.h"
6	#include "sdfg/data_flow/library_nodes/stdlib/memcpy.h"
7	#include "sdfg/data_flow/library_nodes/stdlib/memset.h"
8	#include "sdfg/structured_sdfg.h"
9	#include "sdfg/symbolic/symbolic.h"
10	#include "sdfg/targets/cuda/cuda.h"
11
12	namespace sdfg {
13	namespace cuda {
14
15	std::optional<data_flow::ImplementationType> CudaLibraryNodeRewriter::
16	try_library_node_implementation(const data_flow::LibraryNode& lib_node, types::PrimitiveType data_type) {	1✔
17	if (data_type == types::PrimitiveType::Float \|\| data_type == types::PrimitiveType::Double) {	1✔
18	if (lib_node.code() == math::blas::LibraryNodeType_GEMM.value()) {	1✔
19	auto& gemm_node = static_cast<const math::blas::GEMMNode&>(lib_node);	×
20	return try_cublas_gemm_node_implementation(gemm_node, data_type);	×
21	} else if (lib_node.code() == math::blas::LibraryNodeType_BatchedGEMM.value()) {	1✔
22	return cuda::ImplementationType_CUDAWithTransfers;	1✔
23	} else if (lib_node.code() == math::blas::LibraryNodeType_DOT.value()) {	1✔
24	return cuda::ImplementationType_CUDAWithTransfers;	×
25	} else if (lib_node.code() == math::blas::LibraryNodeType_BatchedGEMM.value()) {	×
26	auto& batched_gemm_node = static_cast<const math::blas::BatchedGEMMNode&>(lib_node);	×
27	return try_cublas_batched_gemm_node_implementation(batched_gemm_node, data_type);	×
28	} else {	×
29	return std::nullopt;	×
30	}	×
31	} else {	1✔
32	return std::nullopt;	×
33	}	×
34	}	1✔
35
36	std::optional<data_flow::ImplementationType> CudaLibraryNodeRewriter::
37	try_cublas_gemm_node_implementation(const math::blas::GEMMNode& gemm_node, types::PrimitiveType data_type) {	×
UNCOV 38	return cuda::ImplementationType_CUDAWithTransfers;	×
39	}	×
40
41	std::optional<data_flow::ImplementationType> CudaLibraryNodeRewriter::
42	try_memset_implementation(const ::sdfg::stdlib::MemsetNode& memset_node) {	2✔
43	return cuda::ImplementationType_CUDAWithTransfers;	2✔
44	}	2✔
45
46	std::optional<data_flow::ImplementationType> CudaLibraryNodeRewriter::
47	try_memcpy_implementation(const ::sdfg::stdlib::MemcpyNode& memcpy_node) {	×
48	return cuda::ImplementationType_CUDAWithTransfers;	×
49	}	×
50
51	std::optional<data_flow::ImplementationType> CudaLibraryNodeRewriter::try_cublas_batched_gemm_node_implementation(
52	const math::blas::BatchedGEMMNode& batched_gemm_node, types::PrimitiveType data_type
53	) {	×
54	// Heuristic: Avoid using CUBLAS for very small matrix multiplications
55	auto m = batched_gemm_node.m();	×
56	auto n = batched_gemm_node.n();	×
57	auto k = batched_gemm_node.k();	×
58	auto size = symbolic::mul(symbolic::mul(m, n), k);	×
59	if (symbolic::eq(size, symbolic::one())) {	×
60	return std::nullopt;	×
61	}	×
62	return cuda::ImplementationType_CUDAWithTransfers;	×
63	}	×
64
65	CudaLibraryNodeRewriter::
66	CudaLibraryNodeRewriter(builder::StructuredSDFGBuilder& builder, analysis::AnalysisManager& analysis_manager)
67	: visitor::StructuredSDFGVisitor(builder, analysis_manager) {}	3✔
68
69	bool CudaLibraryNodeRewriter::accept(structured_control_flow::Block& node) {	3✔
70	auto& dataflow = node.dataflow();	3✔
71	for (auto& library_node : dataflow.nodes()) {	10✔
72	if (auto lib_node = dynamic_cast<math::blas::BLASNode*>(&library_node)) {	10✔
73	auto implType = try_library_node_implementation(*lib_node, lib_node->scalar_primitive());	1✔
74
75	if (implType) {	1✔
76	lib_node->implementation_type() = implType.value();	1✔
77	}	1✔
78	}	1✔
79	if (auto memset_node = dynamic_cast<::sdfg::stdlib::MemsetNode*>(&library_node)) {	10✔
80	auto implType = try_memset_implementation(*memset_node);	2✔
81	if (implType) {	2✔
82	memset_node->implementation_type() = implType.value();	2✔
83	}	2✔
84	}	2✔
85	if (auto memcpy_node = dynamic_cast<::sdfg::stdlib::MemcpyNode*>(&library_node)) {	10✔
86	auto implType = try_memcpy_implementation(*memcpy_node);	×
87	if (implType) {	×
88	memcpy_node->implementation_type() = implType.value();	×
89	}	×
90	}	×
91	}	10✔
92	return false;	3✔
93	}	3✔
94
95	} // namespace cuda
96	} // namespace sdfg

daisytuner / docc / 27876065887

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