28098855020

Committed 24 Jun 2026 12:35PM UTC coverage: 61.615% (-0.3%) from 61.875%

Build # 28098855020

Build Type

Pull #808

github

Committed by

web-flow

Commit Message

Merge 96a21bb28 into b991bdf8b

Pull Request Pull Request #808: enable batched gemm and memcpy on AMD GPUs analogous to the existing CUDA implementation

Coverage Stats

58 of 329 new or added lines in 6 files covered. (17.63%)

5 existing lines in 2 files now uncovered.

37222 of 60411 relevant lines covered (61.61%)

1009.57 hits per line

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

0.0

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

1	#include "sdfg/targets/rocm/stdlib/memcpy.h"
2	#include "sdfg/targets/rocm/rocm.h"
3
4	namespace sdfg::rocm::stdlib {
5
6	MemcpyNodeDispatcher_ROCMWithTransfers::MemcpyNodeDispatcher_ROCMWithTransfers(
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_ROCMWithTransfers::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 <hip/hip_runtime.h>");	×
22
NEW 23	out.stream << "hipError_t err_hip;" << 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_hip = hipMalloc(&d_ptr_in, " << num_expr << ");" << std::endl;	×
NEW 31	rocm_error_checking(out.stream, language_extension_, "err_hip");	×
NEW 32	out.stream << "err_hip = hipMalloc(&d_ptr_out, " << num_expr << ");" << std::endl;	×
NEW 33	rocm_error_checking(out.stream, language_extension_, "err_hip");	×
34
NEW 35	out.stream << "err_hip = hipMemcpy(d_ptr_in, " << inputs.at(1).expr << ", " << num_expr	×
NEW 36	<< ", hipMemcpyHostToDevice);" << std::endl;	×
NEW 37	rocm_error_checking(out.stream, language_extension_, "err_hip");	×
38
NEW 39	out.stream << "err_hip = hipMemcpy(d_ptr_out, d_ptr_in, " << num_expr << ", hipMemcpyDeviceToDevice);" << std::endl;	×
NEW 40	rocm_error_checking(out.stream, language_extension_, "err_hip");	×
41
NEW 42	out.stream << "err_hip = hipMemcpy(" << inputs.at(0).expr << ", d_ptr_out, " << num_expr	×
NEW 43	<< ", hipMemcpyDeviceToHost);" << std::endl;	×
NEW 44	rocm_error_checking(out.stream, language_extension_, "err_hip");	×
45
NEW 46	out.stream << "err_hip = hipFree(d_ptr_in);" << std::endl;	×
NEW 47	rocm_error_checking(out.stream, language_extension_, "err_hip");	×
NEW 48	out.stream << "err_hip = hipFree(d_ptr_out);" << std::endl;	×
NEW 49	rocm_error_checking(out.stream, language_extension_, "err_hip");	×
NEW 50	}	×
51
52	MemcpyNodeDispatcher_ROCMWithoutTransfers::MemcpyNodeDispatcher_ROCMWithoutTransfers(
53	codegen::LanguageExtension& language_extension,
54	const Function& function,
55	const data_flow::DataFlowGraph& data_flow_graph,
56	const sdfg::stdlib::MemcpyNode& node
57	)
NEW 58	: codegen::LibraryNodeDispatcher(language_extension, function, data_flow_graph, node) {}	×
59
60	void MemcpyNodeDispatcher_ROCMWithoutTransfers::dispatch_code_with_edges(
61	codegen::CodegenOutput& out,
62	std::vector<codegen::DispatchInput>& inputs,
63	std::vector<codegen::DispatchOutput>& outputs
NEW 64	) {	×
NEW 65	auto& node = static_cast<const sdfg::stdlib::MemcpyNode&>(node_);	×
66
NEW 67	out.library_snippet_factory.add_global("#include <hip/hip_runtime.h>");	×
68
NEW 69	out.stream << "hipError_t err_hip;" << std::endl;	×
NEW 70	out.stream << "err_hip = hipMemcpy(" << inputs.at(0).expr << ", " << inputs.at(1).expr << ", "	×
NEW 71	<< language_extension_.expression(node.count()) << ", hipMemcpyDeviceToDevice);" << std::endl;	×
NEW 72	rocm_error_checking(out.stream, language_extension_, "err_hip");	×
NEW 73	}	×
74
75	} // namespace sdfg::rocm::stdlib

daisytuner / docc / 28098855020

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