24408034210

Committed 14 Apr 2026 03:32PM UTC coverage: 64.413% (-0.06%) from 64.469%

Build # 24408034210

Build Type

Pull #679

github

Committed by

web-flow

Commit Message

Merge 7f86463c3 into 505b640ce

Pull Request Pull Request #679: Single GPU BLAS Handle Instantiation

Coverage Stats

9 of 92 new or added lines in 15 files covered. (9.78%)

14 existing lines in 7 files now uncovered.

30559 of 47442 relevant lines covered (64.41%)

583.53 hits per line

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

0.0

/opt/src/targets/rocm/stdlib/memset.cpp

#include "sdfg/targets/rocm/stdlib/memset.h"
#include "sdfg/targets/rocm/rocm.h"

namespace sdfg::rocm::stdlib {

MemsetNodeDispatcher_ROCMWithTransfers::MemsetNodeDispatcher_ROCMWithTransfers(
    codegen::LanguageExtension& language_extension,
    const Function& function,
    const data_flow::DataFlowGraph& data_flow_graph,
    const sdfg::stdlib::MemsetNode& node
)
    : codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}

void MemsetNodeDispatcher_ROCMWithTransfers::dispatch_code(
    codegen::PrettyPrinter& stream,
    codegen::PrettyPrinter& globals_stream,
    codegen::CodeSnippetFactory& library_snippet_factory
) {
    auto& node = static_cast<const sdfg::stdlib::MemsetNode&>(node_);

    library_snippet_factory.add_global("#include <hip/hip_runtime.h>");

    stream << "hipError_t err_hip;" << std::endl;

    std::string num_expr = language_extension_.expression(node.num());

    stream << "void *d_ptr;" << std::endl;
    stream << "err_hip = hipMalloc(&d_ptr, " << num_expr << ");" << std::endl;
    rocm_error_checking(stream, language_extension_, "err_hip");

    stream << "err_hip = hipMemset(d_ptr, " << language_extension_.expression(node.value()) << ", " << num_expr << ");"
           << std::endl;
    rocm_error_checking(stream, language_extension_, "err_hip");

    stream << "err_hip = hipMemcpy(" << node.outputs().at(0) << ", d_ptr, " << num_expr << ", hipMemcpyDeviceToHost);"
           << std::endl;
    rocm_error_checking(stream, language_extension_, "err_hip");

    stream << "err_hip = hipFree(d_ptr);" << std::endl;
    rocm_error_checking(stream, language_extension_, "err_hip");
}

MemsetNodeDispatcher_ROCMWithoutTransfers::MemsetNodeDispatcher_ROCMWithoutTransfers(
    codegen::LanguageExtension& language_extension,
    const Function& function,
    const data_flow::DataFlowGraph& data_flow_graph,
    const sdfg::stdlib::MemsetNode& node
)
    : codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}

void MemsetNodeDispatcher_ROCMWithoutTransfers::dispatch_code(
    codegen::PrettyPrinter& stream,
    codegen::PrettyPrinter& globals_stream,
    codegen::CodeSnippetFactory& library_snippet_factory
) {
    auto& node = static_cast<const sdfg::stdlib::MemsetNode&>(node_);

    library_snippet_factory.add_global("#include <hip/hip_runtime.h>");

    stream << "hipError_t err_hip;" << std::endl;
    stream << "err_hip = hipMemset(" << node.outputs().at(0) << ", " << language_extension_.expression(node.value())
           << ", " << language_extension_.expression(node.num()) << ");" << std::endl;
    rocm_error_checking(stream, language_extension_, "err_hip");
}

} // namespace sdfg::rocm::stdlib

1	#include "sdfg/targets/rocm/stdlib/memset.h"
2	#include "sdfg/targets/rocm/rocm.h"
3
4	namespace sdfg::rocm::stdlib {
5
6	MemsetNodeDispatcher_ROCMWithTransfers::MemsetNodeDispatcher_ROCMWithTransfers(
7	codegen::LanguageExtension& language_extension,
8	const Function& function,
9	const data_flow::DataFlowGraph& data_flow_graph,
10	const sdfg::stdlib::MemsetNode& node
11	)
12	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
13
14	void MemsetNodeDispatcher_ROCMWithTransfers::dispatch_code(
15	codegen::PrettyPrinter& stream,
16	codegen::PrettyPrinter& globals_stream,
17	codegen::CodeSnippetFactory& library_snippet_factory
18	) {	×
19	auto& node = static_cast<const sdfg::stdlib::MemsetNode&>(node_);	×
20
NEW 21	library_snippet_factory.add_global("#include <hip/hip_runtime.h>");	×
22
23	stream << "hipError_t err_hip;" << std::endl;	×
24
25	std::string num_expr = language_extension_.expression(node.num());	×
26
27	stream << "void *d_ptr;" << std::endl;	×
28	stream << "err_hip = hipMalloc(&d_ptr, " << num_expr << ");" << std::endl;	×
29	rocm_error_checking(stream, language_extension_, "err_hip");	×
30
31	stream << "err_hip = hipMemset(d_ptr, " << language_extension_.expression(node.value()) << ", " << num_expr << ");"	×
32	<< std::endl;	×
33	rocm_error_checking(stream, language_extension_, "err_hip");	×
34
35	stream << "err_hip = hipMemcpy(" << node.outputs().at(0) << ", d_ptr, " << num_expr << ", hipMemcpyDeviceToHost);"	×
36	<< std::endl;	×
37	rocm_error_checking(stream, language_extension_, "err_hip");	×
38
39	stream << "err_hip = hipFree(d_ptr);" << std::endl;	×
40	rocm_error_checking(stream, language_extension_, "err_hip");	×
41	}	×
42
43	MemsetNodeDispatcher_ROCMWithoutTransfers::MemsetNodeDispatcher_ROCMWithoutTransfers(
44	codegen::LanguageExtension& language_extension,
45	const Function& function,
46	const data_flow::DataFlowGraph& data_flow_graph,
47	const sdfg::stdlib::MemsetNode& node
48	)
49	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
50
51	void MemsetNodeDispatcher_ROCMWithoutTransfers::dispatch_code(
52	codegen::PrettyPrinter& stream,
53	codegen::PrettyPrinter& globals_stream,
54	codegen::CodeSnippetFactory& library_snippet_factory
55	) {	×
56	auto& node = static_cast<const sdfg::stdlib::MemsetNode&>(node_);	×
57
NEW 58	library_snippet_factory.add_global("#include <hip/hip_runtime.h>");	×
59
60	stream << "hipError_t err_hip;" << std::endl;	×
61	stream << "err_hip = hipMemset(" << node.outputs().at(0) << ", " << language_extension_.expression(node.value())	×
62	<< ", " << language_extension_.expression(node.num()) << ");" << std::endl;	×
63	rocm_error_checking(stream, language_extension_, "err_hip");	×
64	}	×
65
66	} // namespace sdfg::rocm::stdlib

daisytuner / docc / 24408034210

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