27288215780

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

Build # 27288215780

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/memcpy.cpp

1	#include "sdfg/targets/cuda/stdlib/memcpy.h"
2	#include "sdfg/targets/cuda/cuda.h"
3
4	namespace sdfg::cuda::stdlib {
5
6	MemcpyNodeDispatcher_CUDAWithTransfers::MemcpyNodeDispatcher_CUDAWithTransfers(
7	codegen::LanguageExtension& language_extension,
8	const Function& function,
9	const data_flow::DataFlowGraph& data_flow_graph,
10	const sdfg::stdlib::MemcpyNode& node
11	)
NEW 12	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
13
14	void MemcpyNodeDispatcher_CUDAWithTransfers::dispatch_code_with_edges(
15	codegen::CodegenOutput& out,
16	std::vector<codegen::DispatchInput>& inputs,
17	std::vector<codegen::DispatchOutput>& outputs
NEW 18	) {	×
NEW 19	auto& node = static_cast<const sdfg::stdlib::MemcpyNode&>(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
NEW 25	std::string num_expr = language_extension_.expression(node.count());	×
26
NEW 27	out.stream << "void *d_ptr_in;" << std::endl;	×
NEW 28	out.stream << "void *d_ptr_out;" << std::endl;	×
29
NEW 30	out.stream << "err_cuda = cudaMalloc(&d_ptr_in, " << num_expr << ");" << std::endl;	×
NEW 31	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
NEW 32	out.stream << "err_cuda = cudaMalloc(&d_ptr_out, " << num_expr << ");" << std::endl;	×
NEW 33	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
34
NEW 35	out.stream << "err_cuda = cudaMemcpy(d_ptr_in, " << inputs.at(0).expr << ", " << num_expr	×
NEW 36	<< ", cudaMemcpyHostToDevice);" << std::endl;	×
NEW 37	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
38
NEW 39	out.stream << "err_cuda = cudaMemcpy(d_ptr_out, d_ptr_in, " << num_expr << ", cudaMemcpyDeviceToDevice);"	×
NEW 40	<< std::endl;	×
NEW 41	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
42
NEW 43	out.stream << "err_cuda = cudaMemcpy(" << inputs.at(1).expr << ", d_ptr_out, " << num_expr	×
NEW 44	<< ", cudaMemcpyDeviceToHost);" << std::endl;	×
NEW 45	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
46
NEW 47	out.stream << "err_cuda = cudaFree(d_ptr_in);" << std::endl;	×
NEW 48	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
NEW 49	out.stream << "err_cuda = cudaFree(d_ptr_out);" << std::endl;	×
NEW 50	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
NEW 51	}	×
52
53	MemcpyNodeDispatcher_CUDAWithoutTransfers::MemcpyNodeDispatcher_CUDAWithoutTransfers(
54	codegen::LanguageExtension& language_extension,
55	const Function& function,
56	const data_flow::DataFlowGraph& data_flow_graph,
57	const sdfg::stdlib::MemcpyNode& node
58	)
NEW 59	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
60
61	void MemcpyNodeDispatcher_CUDAWithoutTransfers::dispatch_code_with_edges(
62	codegen::CodegenOutput& out,
63	std::vector<codegen::DispatchInput>& inputs,
64	std::vector<codegen::DispatchOutput>& outputs
NEW 65	) {	×
NEW 66	auto& node = static_cast<const sdfg::stdlib::MemcpyNode&>(node_);	×
67
NEW 68	out.library_snippet_factory.add_global("#include <cuda.h>");	×
69
NEW 70	out.stream << "cudaError_t err_cuda;" << std::endl;	×
NEW 71	out.stream << "err_cuda = cudaMemcpy(" << inputs.at(1).expr << ", " << inputs.at(0).expr << ", "	×
NEW 72	<< language_extension_.expression(node.count()) << ", cudaMemcpyDeviceToDevice);" << std::endl;	×
NEW 73	cuda_error_checking(out.stream, language_extension_, "err_cuda");	×
NEW 74	}	×
75
76	} // namespace sdfg::cuda::stdlib

daisytuner / docc / 27288215780

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