24237992383

Committed 10 Apr 2026 10:12AM UTC coverage: 64.377% (-0.02%) from 64.394%

Build # 24237992383

Build Type

push

github

Committed by

web-flow

Commit Message

Offload Memset to GPU (#668)

* Rename implementation type

* add memset offloading and expansion

* rename offload expansion passes and transformations

* Add unit tests

* rename offloading expansion to data transfer extraction

* fix counters

Coverage Stats

249 of 381 new or added lines in 18 files covered. (65.35%)

29966 of 46548 relevant lines covered (64.38%)

584.15 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

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	)
NEW 12	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
13
14	void MemsetNodeDispatcher_CUDAWithTransfers::dispatch_code(
15	codegen::PrettyPrinter& stream,
16	codegen::PrettyPrinter& globals_stream,
17	codegen::CodeSnippetFactory& library_snippet_factory
NEW 18	) {	×
NEW 19	auto& node = static_cast<const sdfg::stdlib::MemsetNode&>(node_);	×
20
NEW 21	globals_stream << "#include <cuda.h>" << std::endl;	×
22
NEW 23	stream << "cudaError_t err_cuda;" << std::endl;	×
24
NEW 25	std::string num_expr = language_extension_.expression(node.num());	×
26
NEW 27	stream << "void *d_ptr;" << std::endl;	×
NEW 28	stream << "err_cuda = cudaMalloc(&d_ptr, " << num_expr << ");" << std::endl;	×
NEW 29	cuda_error_checking(stream, language_extension_, "err_cuda");	×
30
NEW 31	stream << "err_cuda = cudaMemset(d_ptr, " << language_extension_.expression(node.value()) << ", " << num_expr	×
NEW 32	<< ");" << std::endl;	×
NEW 33	cuda_error_checking(stream, language_extension_, "err_cuda");	×
34
NEW 35	stream << "err_cuda = cudaMemcpy(" << node.outputs().at(0) << ", d_ptr, " << num_expr	×
NEW 36	<< ", cudaMemcpyDeviceToHost);" << std::endl;	×
NEW 37	cuda_error_checking(stream, language_extension_, "err_cuda");	×
38
NEW 39	stream << "err_cuda = cudaFree(d_ptr);" << std::endl;	×
NEW 40	cuda_error_checking(stream, language_extension_, "err_cuda");	×
NEW 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	)
NEW 49	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
50
51	void MemsetNodeDispatcher_CUDAWithoutTransfers::dispatch_code(
52	codegen::PrettyPrinter& stream,
53	codegen::PrettyPrinter& globals_stream,
54	codegen::CodeSnippetFactory& library_snippet_factory
NEW 55	) {	×
NEW 56	auto& node = static_cast<const sdfg::stdlib::MemsetNode&>(node_);	×
57
NEW 58	globals_stream << "#include <cuda.h>" << std::endl;	×
59
NEW 60	stream << "cudaError_t err_cuda;" << std::endl;	×
NEW 61	stream << "err_cuda = cudaMemset(" << node.outputs().at(0) << ", " << language_extension_.expression(node.value())	×
NEW 62	<< ", " << language_extension_.expression(node.num()) << ");" << std::endl;	×
NEW 63	cuda_error_checking(stream, language_extension_, "err_cuda");	×
NEW 64	}	×
65
66	} // namespace sdfg::cuda::stdlib

daisytuner / docc / 24237992383

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