27325952678

Committed 10 Jun 2026 03:49PM UTC coverage: 61.108% (-0.08%) from 61.191%

Build # 27325952678

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #747 from daisytuner/memcpyOnDevice

Translate CPU memcpy to GPU memcpy

Coverage Stats

41 of 149 new or added lines in 6 files covered. (27.52%)

2 existing lines in 2 files now uncovered.

36048 of 58991 relevant lines covered (61.11%)

747.87 hits per line

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

0.0

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

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

namespace sdfg::cuda::stdlib {

MemsetNodeDispatcher_CUDAWithTransfers::MemsetNodeDispatcher_CUDAWithTransfers(
    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_CUDAWithTransfers::dispatch_code_with_edges(
    codegen::CodegenOutput& out,
    std::vector<codegen::DispatchInput>& inputs,
    std::vector<codegen::DispatchOutput>& outputs
) {
    auto& node = static_cast<const sdfg::stdlib::MemsetNode&>(node_);

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

    out.stream << "cudaError_t err_cuda;" << std::endl;

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

    out.stream << "void *d_ptr;" << std::endl;
    out.stream << "err_cuda = cudaMalloc(&d_ptr, " << num_expr << ");" << std::endl;
    cuda_error_checking(out.stream, language_extension_, "err_cuda");

    out.stream << "err_cuda = cudaMemset(d_ptr, " << language_extension_.expression(node.value()) << ", " << num_expr
               << ");" << std::endl;
    cuda_error_checking(out.stream, language_extension_, "err_cuda");

    out.stream << "err_cuda = cudaMemcpy(" << inputs.at(0).expr << ", d_ptr, " << num_expr
               << ", cudaMemcpyDeviceToHost);" << std::endl;
    cuda_error_checking(out.stream, language_extension_, "err_cuda");

    out.stream << "err_cuda = cudaFree(d_ptr);" << std::endl;
    cuda_error_checking(out.stream, language_extension_, "err_cuda");
}

MemsetNodeDispatcher_CUDAWithoutTransfers::MemsetNodeDispatcher_CUDAWithoutTransfers(
    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_CUDAWithoutTransfers::dispatch_code_with_edges(
    codegen::CodegenOutput& out,
    std::vector<codegen::DispatchInput>& inputs,
    std::vector<codegen::DispatchOutput>& outputs
) {
    auto& node = static_cast<const sdfg::stdlib::MemsetNode&>(node_);

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

    out.stream << "cudaError_t err_cuda;" << std::endl;
    out.stream << "err_cuda = cudaMemset(" << inputs.at(0).expr << ", " << language_extension_.expression(node.value())
               << ", " << language_extension_.expression(node.num()) << ");" << std::endl;
    cuda_error_checking(out.stream, language_extension_, "err_cuda");
}

} // namespace sdfg::cuda::stdlib

1	#include "sdfg/targets/cuda/stdlib/memset.h"
2	#include "sdfg/targets/cuda/cuda.h"
3
4	namespace sdfg::cuda::stdlib {
5
6	MemsetNodeDispatcher_CUDAWithTransfers::MemsetNodeDispatcher_CUDAWithTransfers(
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_CUDAWithTransfers::dispatch_code_with_edges(
15	codegen::CodegenOutput& out,
16	std::vector<codegen::DispatchInput>& inputs,
17	std::vector<codegen::DispatchOutput>& outputs
18	) {	×
19	auto& node = static_cast<const sdfg::stdlib::MemsetNode&>(node_);	×
20
NEW 21	out.library_snippet_factory.add_global("#include <cuda.h>");	×
22
NEW 23	out.stream << "cudaError_t err_cuda;" << std::endl;	×
24
25	std::string num_expr = language_extension_.expression(node.num());	×
26
NEW 27	out.stream << "void *d_ptr;" << std::endl;	×
NEW 28	out.stream << "err_cuda = cudaMalloc(&d_ptr, " << num_expr << ");" << std::endl;	×
NEW 29	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
30
NEW 31	out.stream << "err_cuda = cudaMemset(d_ptr, " << language_extension_.expression(node.value()) << ", " << num_expr	×
NEW 32	<< ");" << std::endl;	×
NEW 33	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
34
NEW 35	out.stream << "err_cuda = cudaMemcpy(" << inputs.at(0).expr << ", d_ptr, " << num_expr	×
NEW 36	<< ", cudaMemcpyDeviceToHost);" << std::endl;	×
NEW 37	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
38
NEW 39	out.stream << "err_cuda = cudaFree(d_ptr);" << std::endl;	×
NEW 40	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
UNCOV 41	}	×
42
43	MemsetNodeDispatcher_CUDAWithoutTransfers::MemsetNodeDispatcher_CUDAWithoutTransfers(
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_CUDAWithoutTransfers::dispatch_code_with_edges(
52	codegen::CodegenOutput& out,
53	std::vector<codegen::DispatchInput>& inputs,
54	std::vector<codegen::DispatchOutput>& outputs
55	) {	×
56	auto& node = static_cast<const sdfg::stdlib::MemsetNode&>(node_);	×
57
58	out.library_snippet_factory.add_global("#include <cuda.h>");	×
59
60	out.stream << "cudaError_t err_cuda;" << std::endl;	×
61	out.stream << "err_cuda = cudaMemset(" << inputs.at(0).expr << ", " << language_extension_.expression(node.value())	×
62	<< ", " << language_extension_.expression(node.num()) << ");" << std::endl;	×
63	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
64	}	×
65
66	} // namespace sdfg::cuda::stdlib

daisytuner / docc / 27325952678

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