diff --git a/llvm/integration/rodinia_openmp_test.py b/llvm/integration/rodinia_openmp_test.py index 850a2f587..43d61a135 100644 --- a/llvm/integration/rodinia_openmp_test.py +++ b/llvm/integration/rodinia_openmp_test.py @@ -354,9 +354,9 @@ def test_cfd(compiler="clang++-19"): verification={ "sdfgs": 6, "MAP": 14, - "CPU_PARALLEL": 5, + "CPU_PARALLEL": 12, "FOR": 13, - "SEQUENTIAL": 22, + "SEQUENTIAL": 15, } ) runner = TestRunner( diff --git a/opt/CMakeLists.txt b/opt/CMakeLists.txt index eb97545be..fe91d2487 100644 --- a/opt/CMakeLists.txt +++ b/opt/CMakeLists.txt @@ -53,6 +53,7 @@ set(SOURCE_FILES src/passes/map_fusion_by_domain_pass.cpp src/passes/opt_pipeline.cpp src/passes/redundant_load_elimination_pass.cpp + src/passes/zero_fill_to_memset_pass.cpp src/passes/scheduler/cuda_scheduler.cpp src/passes/scheduler/rocm_scheduler.cpp src/passes/scheduler/vectorize_scheduler.cpp diff --git a/opt/include/sdfg/passes/zero_fill_to_memset_pass.h b/opt/include/sdfg/passes/zero_fill_to_memset_pass.h new file mode 100644 index 000000000..f453c3600 --- /dev/null +++ b/opt/include/sdfg/passes/zero_fill_to_memset_pass.h @@ -0,0 +1,78 @@ +#pragma once + +#include +#include +#include + +#include "sdfg/analysis/loop_analysis.h" +#include "sdfg/analysis/memory_layout_analysis.h" +#include "sdfg/passes/pass.h" +#include "sdfg/structured_control_flow/map.h" +#include "sdfg/symbolic/symbolic.h" +#include "sdfg/types/type.h" +#include "sdfg/visitor/structured_sdfg_visitor.h" + +namespace sdfg::passes { + +/** + * @brief Replaces perfectly-nested Map loops that zero out an entire array with a single memset. + * + * The pass detects a perfect nest of Map loops whose only effect is to overwrite the + * complete contents of a (multi-dimensional) array container with the literal value 0 + * (and nothing else). Such a nest is semantically equivalent to a `memset(A, 0, sizeof(A))` + * and is rewritten to a single stdlib::MemsetNode. + * + * A nest qualifies when: + * - Each loop level is a Map with `init == 0`, unit stride and an extractable exclusive + * upper bound; the nest is perfect (each Map body contains exactly the next level). + * - The innermost body is a single Block whose only computation assigns the constant 0 + * to one element of the array, indexed exactly by the induction variables of the nest. + * - The MemoryLayoutAnalysis tile for the target container at the outermost Map covers the + * whole array contiguously from offset 0 (bounded dimensions match the declared extent, + * and the covered elements form a dense block), guaranteeing the whole array is covered. + * Relying on the layout analysis also transparently handles linearized accesses such as + * `A[i*M + j]`, which are delinearized to the underlying multi-dimensional shape. + */ +class ZeroFillToMemsetPass : public sdfg::passes::Pass { +public: + struct State { + size_t applied = 0; + }; + + std::string name() override { return "ZeroFillToMemset"; } + + bool run_pass(builder::StructuredSDFGBuilder& builder, analysis::AnalysisManager& analysis_manager) override; +}; + +class ZeroFillToMemsetVisitor : public sdfg::visitor::ActualStructuredSDFGVisitor { +public: + struct Candidate { + structured_control_flow::Map* map; + std::string array; + symbolic::Expression num; + std::unique_ptr ptr_type; + }; + +private: + builder::StructuredSDFGBuilder& builder_; + ZeroFillToMemsetPass::State& state_; + analysis::MemoryLayoutAnalysis& memory_layout_analysis_; + analysis::LoopAnalysis& loop_analysis_; + std::vector candidates_; + + bool match(structured_control_flow::Map& node, Candidate& candidate); + +public: + ZeroFillToMemsetVisitor( + builder::StructuredSDFGBuilder& builder, + ZeroFillToMemsetPass::State& state, + analysis::MemoryLayoutAnalysis& memory_layout_analysis, + analysis::LoopAnalysis& loop_analysis + ); + + bool visit(structured_control_flow::Map& node) override; + + void apply(); +}; + +} // namespace sdfg::passes diff --git a/opt/include/sdfg/targets/rocm/stdlib/memset.h b/opt/include/sdfg/targets/rocm/stdlib/memset.h index 277b36c66..ee188135e 100644 --- a/opt/include/sdfg/targets/rocm/stdlib/memset.h +++ b/opt/include/sdfg/targets/rocm/stdlib/memset.h @@ -14,10 +14,10 @@ class MemsetNodeDispatcher_ROCMWithTransfers : public codegen::LibraryNodeDispat const sdfg::stdlib::MemsetNode& node ); - void dispatch_code( - codegen::PrettyPrinter& stream, - codegen::PrettyPrinter& globals_stream, - codegen::CodeSnippetFactory& library_snippet_factory + void dispatch_code_with_edges( + codegen::CodegenOutput& out, + std::vector& inputs, + std::vector& outputs ) override; }; diff --git a/opt/src/passes/scheduler/omp_scheduler.cpp b/opt/src/passes/scheduler/omp_scheduler.cpp index 48c440162..6303f7f3e 100644 --- a/opt/src/passes/scheduler/omp_scheduler.cpp +++ b/opt/src/passes/scheduler/omp_scheduler.cpp @@ -23,8 +23,7 @@ SchedulerAction OMPScheduler::find( auto& loop_analysis = analysis_manager.get(); auto loop_info = loop_analysis.loop_info(&loop); - if (loop_info.loopnest_index == -1 || loop_info.num_maps <= 1 || loop_info.is_perfectly_nested || - loop_info.has_side_effects) { + if (loop_info.loopnest_index == -1 || loop_info.num_maps <= 1 || loop_info.is_perfectly_nested) { return NEXT; } else { return CHILDREN; diff --git a/opt/src/passes/zero_fill_to_memset_pass.cpp b/opt/src/passes/zero_fill_to_memset_pass.cpp new file mode 100644 index 000000000..69cb369c5 --- /dev/null +++ b/opt/src/passes/zero_fill_to_memset_pass.cpp @@ -0,0 +1,311 @@ +#include "sdfg/passes/zero_fill_to_memset_pass.h" + +#include +#include +#include +#include + +#include "sdfg/analysis/loop_analysis.h" +#include "sdfg/analysis/memory_layout_analysis.h" +#include "sdfg/data_flow/access_node.h" +#include "sdfg/data_flow/library_nodes/stdlib/memset.h" +#include "sdfg/data_flow/memlet.h" +#include "sdfg/data_flow/tasklet.h" +#include "sdfg/structured_control_flow/block.h" +#include "sdfg/structured_control_flow/sequence.h" +#include "sdfg/structured_control_flow/structured_loop.h" +#include "sdfg/symbolic/symbolic.h" +#include "sdfg/types/pointer.h" +#include "sdfg/types/scalar.h" +#include "sdfg/types/utils.h" + +namespace sdfg::passes { + +// Returns true if the literal string represents a numeric zero (e.g. "0", "0.0", "-0.0f", "false"). +static bool is_zero_literal(const std::string& value) { + std::string t; + for (char c : value) { + if (!std::isspace(static_cast(c))) { + t += c; + } + } + if (t.empty()) { + return false; + } + if (t == "false") { + return true; + } + + size_t begin = 0; + if (t[begin] == '+' || t[begin] == '-') { + begin++; + } + size_t end = t.size(); + while (end > begin && (t[end - 1] == 'f' || t[end - 1] == 'F' || t[end - 1] == 'l' || t[end - 1] == 'L')) { + end--; + } + + bool saw_digit = false; + for (size_t k = begin; k < end; k++) { + char c = t[k]; + if (c == '0') { + saw_digit = true; + } else if (c == '.' || c == 'x' || c == 'X') { + continue; + } else { + return false; + } + } + return saw_digit; +} + +// Decides whether the tile of accessed elements covers the entire container contiguously +// starting at offset 0, so that the loop nest is equivalent to a single memset. +// +// `bounds` are the exclusive upper bounds of the loop nest (outermost first); their product +// is the number of written elements. The check succeeds when +// - every dimension starts at index 0, +// - every dimension with a known (bounded) extent is fully covered (extent == declared +// shape), which rejects partial fills of statically sized arrays, and +// - the accessed elements form a dense, contiguous block [0, count) where count equals +// the number of loop iterations, which rejects strided / gapped accesses. +static bool covers_full_container(const analysis::MemoryTile& tile, const std::vector& bounds) { + size_t ndims = tile.min_subset.size(); + if (ndims == 0) { + return false; + } + const auto& shape = tile.layout.shape(); + if (shape.size() != ndims) { + return false; + } + + auto extents = tile.extents(); + if (extents.size() != ndims) { + return false; + } + + for (size_t d = 0; d < ndims; d++) { + if (!symbolic::eq(tile.min_subset[d], symbolic::zero())) { + return false; + } + if (extents[d].is_null()) { + return false; + } + // Only the leading dimension can be unbounded (raw pointer base); its extent is + // not encoded in the type and is therefore trusted to match the loop. Every other + // dimension must fully cover its declared extent. + bool bounded = (d != 0) || tile.first_dim_bounded; + if (bounded && !symbolic::eq(extents[d], shape[d])) { + return false; + } + } + + // The written elements must form a contiguous block [0, trip) with no gaps. + auto [first, last] = tile.contiguous_range(); + if (first.is_null() || last.is_null()) { + return false; + } + if (!symbolic::eq(first, symbolic::zero())) { + return false; + } + + symbolic::Expression trip = symbolic::one(); + for (const auto& bound : bounds) { + trip = symbolic::mul(trip, bound); + } + return symbolic::eq(symbolic::add(last, symbolic::one()), trip); +} + +ZeroFillToMemsetVisitor::ZeroFillToMemsetVisitor( + builder::StructuredSDFGBuilder& builder, + ZeroFillToMemsetPass::State& state, + analysis::MemoryLayoutAnalysis& memory_layout_analysis, + analysis::LoopAnalysis& loop_analysis +) + : builder_(builder), state_(state), memory_layout_analysis_(memory_layout_analysis), loop_analysis_(loop_analysis) { +} + +bool ZeroFillToMemsetVisitor::match(structured_control_flow::Map& node, Candidate& candidate) { + // The nest must be a perfect stack consisting solely of Maps. LoopAnalysis decides this: + // `is_perfectly_nested` guarantees every level holds exactly one child loop (a single + // linear chain), and `num_maps == num_loops` guarantees each of those loops is a Map. + // This gives a finite nesting depth to iterate over instead of an open-ended walk. + const analysis::LoopInfo info = loop_analysis_.loop_info(&node); + if (!info.is_perfectly_nested || info.num_loops == 0 || info.num_loops != info.num_maps) { + return false; + } + const size_t depth = info.num_loops; + + // Walk the (now bounded) perfect nest of Maps, collecting the exclusive upper bounds. + std::vector bounds; + bounds.reserve(depth); + + structured_control_flow::StructuredLoop* loop = &node; + structured_control_flow::Block* terminal = nullptr; + for (size_t level = 0; level < depth; level++) { + // Each loop must be a normalized [0, bound) unit-stride loop. + if (!symbolic::eq(loop->init(), symbolic::zero()) || !loop->is_contiguous()) { + return false; + } + auto bound = loop->canonical_bound_upper(); + if (bound.is_null()) { + return false; + } + bounds.push_back(bound); + + // Perfect nesting guarantees a single child; additionally reject transitions that + // carry symbol assignments, which the loop analysis does not consider. + auto& body = loop->root(); + if (body.size() != 1) { + return false; + } + auto entry = body.at(0); + if (!entry.second.assignments().empty()) { + return false; + } + + auto& child = entry.first; + if (level + 1 < depth) { + // Intermediate level: the single child is the next Map in the chain. + auto* inner_map = dynamic_cast(&child); + if (inner_map == nullptr) { + return false; + } + loop = inner_map; + } else { + // Innermost level: the single child is the terminal computation block. + terminal = dynamic_cast(&child); + if (terminal == nullptr) { + return false; + } + } + } + if (terminal == nullptr) { + return false; + } + + // The innermost body must be a single "A[...] = 0" assignment and nothing else. + auto& dfg = terminal->dataflow(); + if (dfg.tasklets().size() != 1 || !dfg.library_nodes().empty() || dfg.data_nodes().size() != 2) { + return false; + } + + auto* tasklet = *dfg.tasklets().begin(); + if (tasklet->code() != data_flow::TaskletCode::assign || tasklet->inputs().size() != 1) { + return false; + } + + const data_flow::Memlet* in_edge = nullptr; + for (auto& edge : dfg.in_edges(*tasklet)) { + if (in_edge != nullptr) { + return false; + } + in_edge = &edge; + } + const data_flow::Memlet* out_edge = nullptr; + for (auto& edge : dfg.out_edges(*tasklet)) { + if (out_edge != nullptr) { + return false; + } + out_edge = &edge; + } + if (in_edge == nullptr || out_edge == nullptr) { + return false; + } + + // Input: the constant literal 0. + auto* constant = dynamic_cast(&in_edge->src()); + if (constant == nullptr || !is_zero_literal(constant->data())) { + return false; + } + + // Output: a plain array element write. + auto* access = dynamic_cast(&out_edge->dst()); + if (access == nullptr || dynamic_cast(access) != nullptr) { + return false; + } + const std::string& container = access->data(); + + // Coverage is decided by the memory-layout analysis: the tile for this container at the + // outermost Map must cover the whole array contiguously. This handles native arrays, + // pointer bases and linearized accesses (e.g. `A[i*M + j]`) uniformly. + const analysis::MemoryTile* tile = memory_layout_analysis_.tile(node, container); + if (tile == nullptr || !covers_full_container(*tile, bounds)) { + return false; + } + + // The container's element must be a scalar so a byte-wise memset is well-defined. + const auto& container_type = builder_.subject().type(container); + const auto& element_type = types::peel_to_innermost_element(container_type); + const auto* scalar = dynamic_cast(&element_type); + if (scalar == nullptr) { + return false; + } + auto element_size = types::get_contiguous_element_size(container_type); + if (element_size.is_null()) { + return false; + } + + // Total size in bytes = element size times the number of written elements. + symbolic::Expression num = element_size; + for (const auto& bound : bounds) { + num = symbolic::mul(num, bound); + } + + candidate.map = &node; + candidate.array = container; + candidate.num = num; + candidate.ptr_type = std::make_unique(*scalar); + return true; +} + +bool ZeroFillToMemsetVisitor::visit(structured_control_flow::Map& node) { + Candidate candidate; + if (match(node, candidate)) { + candidates_.push_back(std::move(candidate)); + // Do not descend into a matched nest. + return false; + } + return ActualStructuredSDFGVisitor::visit(node); +} + +void ZeroFillToMemsetVisitor::apply() { + for (auto& candidate : candidates_) { + auto* parent = dynamic_cast(candidate.map->get_parent()); + if (parent == nullptr) { + continue; + } + int index = parent->index(*candidate.map); + if (index < 0) { + continue; + } + + // Preserve any symbol assignments attached to the replaced nest. + auto assignments = parent->at(index).second.assignments(); + + auto& block = builder_.add_block_before(*parent, *candidate.map, assignments); + stdlib::add_memset_node(builder_, block, candidate.array, symbolic::integer(0), candidate.num, *candidate.ptr_type); + + int map_index = parent->index(*candidate.map); + if (map_index < 0) { + continue; + } + builder_.remove_child(*parent, map_index); + state_.applied++; + } +} + +bool ZeroFillToMemsetPass::run_pass(builder::StructuredSDFGBuilder& builder, analysis::AnalysisManager& analysis_manager) { + State state; + + auto& memory_layout_analysis = analysis_manager.get(); + auto& loop_analysis = analysis_manager.get(); + + ZeroFillToMemsetVisitor visitor(builder, state, memory_layout_analysis, loop_analysis); + visitor.dispatch(builder.subject().root()); + visitor.apply(); + + return state.applied > 0; +} + +} // namespace sdfg::passes diff --git a/opt/src/targets/rocm/stdlib/memset.cpp b/opt/src/targets/rocm/stdlib/memset.cpp index 0fe83e41e..45deda10d 100644 --- a/opt/src/targets/rocm/stdlib/memset.cpp +++ b/opt/src/targets/rocm/stdlib/memset.cpp @@ -11,33 +11,33 @@ MemsetNodeDispatcher_ROCMWithTransfers::MemsetNodeDispatcher_ROCMWithTransfers( ) : 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 +void MemsetNodeDispatcher_ROCMWithTransfers::dispatch_code_with_edges( + codegen::CodegenOutput& out, + std::vector& inputs, + std::vector& outputs ) { auto& node = static_cast(node_); - library_snippet_factory.add_global("#include "); + out.library_snippet_factory.add_global("#include "); - stream << "hipError_t err_hip;" << std::endl; + out.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"); + out.stream << "void *d_ptr;" << std::endl; + out.stream << "err_hip = hipMalloc(&d_ptr, " << num_expr << ");" << std::endl; + rocm_error_checking(out.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"); + out.stream << "err_hip = hipMemset(d_ptr, " << language_extension_.expression(node.value()) << ", " << num_expr + << ");" << std::endl; + rocm_error_checking(out.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"); + out.stream << "err_hip = hipMemcpy(" << inputs.at(0).expr << ", d_ptr, " << num_expr << ", hipMemcpyDeviceToHost);" + << std::endl; + rocm_error_checking(out.stream, language_extension_, "err_hip"); - stream << "err_hip = hipFree(d_ptr);" << std::endl; - rocm_error_checking(stream, language_extension_, "err_hip"); + out.stream << "err_hip = hipFree(d_ptr);" << std::endl; + rocm_error_checking(out.stream, language_extension_, "err_hip"); } MemsetNodeDispatcher_ROCMWithoutTransfers::MemsetNodeDispatcher_ROCMWithoutTransfers( diff --git a/opt/tests/CMakeLists.txt b/opt/tests/CMakeLists.txt index b2d285372..92f7ed75b 100644 --- a/opt/tests/CMakeLists.txt +++ b/opt/tests/CMakeLists.txt @@ -38,6 +38,7 @@ set(TEST_FILES optimizations/gpu_kernels_test.cpp passes/map_fusion_by_domain_pass_test.cpp passes/redundant_load_elimination_pass_test.cpp + passes/zero_fill_to_memset_pass_test.cpp passes/offloading/cuda_library_node_transfer_extraction_pass_test.cpp passes/offloading/cuda_library_node_expand_tests.cpp passes/offloading/rocm_library_node_expand_tests.cpp diff --git a/opt/tests/passes/zero_fill_to_memset_pass_test.cpp b/opt/tests/passes/zero_fill_to_memset_pass_test.cpp new file mode 100644 index 000000000..e1fc86b71 --- /dev/null +++ b/opt/tests/passes/zero_fill_to_memset_pass_test.cpp @@ -0,0 +1,619 @@ +#include "sdfg/passes/zero_fill_to_memset_pass.h" + +#include + +#include +#include +#include + +#include "sdfg/builder/structured_sdfg_builder.h" +#include "sdfg/data_flow/library_nodes/stdlib/memset.h" +#include "sdfg/structured_control_flow/map.h" +#include "sdfg/structured_control_flow/structured_loop.h" +#include "sdfg/symbolic/symbolic.h" +#include "sdfg/types/array.h" +#include "sdfg/types/pointer.h" +#include "sdfg/types/scalar.h" +#include "sdfg_debug_dump.h" + +using namespace sdfg; + +namespace { + +// Element type used throughout: 8-byte double. +static const types::Scalar kElement(types::PrimitiveType::Double); +static constexpr int kElementBytes = 8; + +// Helper wrapper that builds normalized [0, bound) unit-stride Map nests. +class NestFixture { +public: + builder::StructuredSDFGBuilder& builder; + structured_control_flow::ScheduleType sched = structured_control_flow::ScheduleType_Sequential::create(); + + explicit NestFixture(builder::StructuredSDFGBuilder& builder) : builder(builder) {} + + // Registers a size symbol as an integer container (idempotent per name). + symbolic::Symbol size(const std::string& name) { + if (!builder.subject().exists(name)) { + builder.add_container(name, types::Scalar(types::PrimitiveType::Int64), true); + // Array/loop extents are positive; make that explicit for the layout analysis. + builder.subject().assumption(symbolic::symbol(name)).add_lower_bound(symbolic::one()); + } + return symbolic::symbol(name); + } + + structured_control_flow::Map& + add_map(structured_control_flow::Sequence& parent, const std::string& iv, const symbolic::Expression& bound) { + builder.add_container(iv, types::Scalar(types::PrimitiveType::Int64)); + auto sym = symbolic::symbol(iv); + return builder + .add_map(parent, sym, symbolic::Lt(sym, bound), symbolic::zero(), symbolic::add(sym, symbolic::one()), sched); + } + + // Builds a perfect nest of maps [iv_0, bound_0], ... and returns the innermost body. + structured_control_flow::Sequence& build_nest( + structured_control_flow::Sequence& root, + const std::vector& ivs, + const std::vector& bounds + ) { + structured_control_flow::Sequence* cur = &root; + for (size_t i = 0; i < ivs.size(); i++) { + auto& map = add_map(*cur, ivs[i], bounds[i]); + cur = &map.root(); + } + return *cur; + } + + // Emits a single "array[subset] = " assignment into a fresh block. + void write_scalar( + structured_control_flow::Sequence& body, + const std::string& array, + const data_flow::Subset& subset, + const types::IType& base_type, + const std::string& value + ) { + auto& block = builder.add_block(body); + auto& constant = builder.add_constant(block, value, kElement); + auto& access = builder.add_access(block, array); + auto& tasklet = builder.add_tasklet(block, data_flow::TaskletCode::assign, "_out", {"_in"}); + builder.add_computational_memlet(block, constant, tasklet, "_in", {}); + builder.add_computational_memlet(block, tasklet, "_out", access, subset, base_type); + } + + // Emits a single "dst[subset] = src[subset]" load-and-store into a fresh block. + void copy_scalar( + structured_control_flow::Sequence& body, + const std::string& src, + const std::string& dst, + const data_flow::Subset& subset, + const types::IType& base_type + ) { + auto& block = builder.add_block(body); + auto& src_access = builder.add_access(block, src); + auto& dst_access = builder.add_access(block, dst); + auto& tasklet = builder.add_tasklet(block, data_flow::TaskletCode::assign, "_out", {"_in"}); + builder.add_computational_memlet(block, src_access, tasklet, "_in", subset, base_type); + builder.add_computational_memlet(block, tasklet, "_out", dst_access, subset, base_type); + } +}; + +// Collects every MemsetNode reachable at the top level of the SDFG root. +std::vector collect_memsets(structured_control_flow::Sequence& root) { + std::vector result; + for (size_t i = 0; i < root.size(); i++) { + auto* block = dynamic_cast(&root.at(i).first); + if (block == nullptr) { + continue; + } + for (auto* node : block->dataflow().library_nodes()) { + if (auto* memset = dynamic_cast(node)) { + result.push_back(memset); + } + } + } + return result; +} + +// Returns the name of the container the memset writes to (its "_ptr" input). +std::string memset_target(structured_control_flow::Sequence& root, const stdlib::MemsetNode& memset) { + for (size_t i = 0; i < root.size(); i++) { + auto* block = dynamic_cast(&root.at(i).first); + if (block == nullptr) { + continue; + } + auto& dfg = block->dataflow(); + // Only inspect the block that actually owns this memset node; calling + // in_edges() for a node from a different graph is undefined. + bool owns = false; + for (auto* node : dfg.library_nodes()) { + if (static_cast(node) == static_cast(&memset)) { + owns = true; + break; + } + } + if (!owns) { + continue; + } + for (auto& edge : dfg.in_edges(memset)) { + if (auto* access = dynamic_cast(&edge.src())) { + return access->data(); + } + } + } + return ""; +} + +bool run(builder::StructuredSDFGBuilder& builder) { + analysis::AnalysisManager analysis_manager(builder.subject()); + passes::ZeroFillToMemsetPass pass; + return pass.run_pass(builder, analysis_manager); +} + +} // namespace + +// --------------------------------------------------------------------------- +// Positive cases: a full zero-fill nest is rewritten into a single memset. +// --------------------------------------------------------------------------- + +TEST(ZeroFillToMemsetPassTest, FullAccess_1DArray) { + builder::StructuredSDFGBuilder builder("sdfg_1d_array", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + types::Array array(kElement, n); + builder.add_container("A", array); + + auto& body = fx.build_nest(builder.subject().root(), {"i"}, {n}); + fx.write_scalar(body, "A", {symbolic::symbol("i")}, array, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_TRUE(run(builder)); + + auto& root = builder.subject().root(); + auto memsets = collect_memsets(root); + ASSERT_EQ(memsets.size(), 1); + EXPECT_EQ(memset_target(root, *memsets[0]), "A"); + EXPECT_TRUE(symbolic::eq(memsets[0]->value(), symbolic::zero())); + EXPECT_TRUE(symbolic::eq(memsets[0]->num(), symbolic::mul(symbolic::integer(kElementBytes), n))); +} + +TEST(ZeroFillToMemsetPassTest, FullAccess_2DArray) { + builder::StructuredSDFGBuilder builder("sdfg_2d_array", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + types::Array inner(kElement, m); + types::Array array(inner, n); + builder.add_container("A", array); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j"}, {n, m}); + fx.write_scalar(body, "A", {symbolic::symbol("i"), symbolic::symbol("j")}, array, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_TRUE(run(builder)); + + auto& root = builder.subject().root(); + auto memsets = collect_memsets(root); + ASSERT_EQ(memsets.size(), 1); + EXPECT_EQ(memset_target(root, *memsets[0]), "A"); + EXPECT_TRUE(symbolic::eq(memsets[0]->num(), symbolic::mul(symbolic::mul(symbolic::integer(kElementBytes), n), m))); +} + +TEST(ZeroFillToMemsetPassTest, FullAccess_1DPointer) { + builder::StructuredSDFGBuilder builder("sdfg_1d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + types::Pointer ptr(kElement); + builder.add_container("A", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i"}, {n}); + fx.write_scalar(body, "A", {symbolic::symbol("i")}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_TRUE(run(builder)); + + auto& root = builder.subject().root(); + auto memsets = collect_memsets(root); + ASSERT_EQ(memsets.size(), 1); + EXPECT_EQ(memset_target(root, *memsets[0]), "A"); + EXPECT_TRUE(symbolic::eq(memsets[0]->num(), symbolic::mul(symbolic::integer(kElementBytes), n))); +} + +TEST(ZeroFillToMemsetPassTest, FullAccess_2DPointer) { + builder::StructuredSDFGBuilder builder("sdfg_2d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + types::Array inner(kElement, m); + types::Pointer ptr(inner); + builder.add_container("A", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j"}, {n, m}); + fx.write_scalar(body, "A", {symbolic::symbol("i"), symbolic::symbol("j")}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_TRUE(run(builder)); + + auto& root = builder.subject().root(); + auto memsets = collect_memsets(root); + ASSERT_EQ(memsets.size(), 1); + EXPECT_EQ(memset_target(root, *memsets[0]), "A"); + EXPECT_TRUE(symbolic::eq(memsets[0]->num(), symbolic::mul(symbolic::mul(symbolic::integer(kElementBytes), n), m))); +} + +TEST(ZeroFillToMemsetPassTest, FullAccess_3DPointer) { + builder::StructuredSDFGBuilder builder("sdfg_3d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + auto k = fx.size("K"); + types::Array inner(kElement, k); + types::Array middle(inner, m); + types::Pointer ptr(middle); + builder.add_container("A", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j", "l"}, {n, m, k}); + fx.write_scalar(body, "A", {symbolic::symbol("i"), symbolic::symbol("j"), symbolic::symbol("l")}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_TRUE(run(builder)); + + auto& root = builder.subject().root(); + auto memsets = collect_memsets(root); + ASSERT_EQ(memsets.size(), 1); + EXPECT_EQ(memset_target(root, *memsets[0]), "A"); + EXPECT_TRUE(symbolic:: + eq(memsets[0]->num(), + symbolic::mul(symbolic::mul(symbolic::mul(symbolic::integer(kElementBytes), n), m), k))); +} + +TEST(ZeroFillToMemsetPassTest, FullAccess_TwoIndependent1DPointers) { + builder::StructuredSDFGBuilder builder("sdfg_two_1d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + types::Pointer ptr(kElement); + builder.add_container("A", ptr, true); + builder.add_container("B", ptr, true); + + auto& root = builder.subject().root(); + auto& body_a = fx.build_nest(root, {"i"}, {n}); + fx.write_scalar(body_a, "A", {symbolic::symbol("i")}, ptr, "0"); + auto& body_b = fx.build_nest(root, {"j"}, {n}); + fx.write_scalar(body_b, "B", {symbolic::symbol("j")}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_TRUE(run(builder)); + + auto memsets = collect_memsets(root); + ASSERT_EQ(memsets.size(), 2); + std::vector targets{memset_target(root, *memsets[0]), memset_target(root, *memsets[1])}; + EXPECT_NE(std::find(targets.begin(), targets.end(), "A"), targets.end()); + EXPECT_NE(std::find(targets.begin(), targets.end(), "B"), targets.end()); +} + +TEST(ZeroFillToMemsetPassTest, FullAccess_TwoIndependent2DPointers) { + builder::StructuredSDFGBuilder builder("sdfg_two_2d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + types::Array inner(kElement, m); + types::Pointer ptr(inner); + builder.add_container("A", ptr, true); + builder.add_container("B", ptr, true); + + auto& root = builder.subject().root(); + auto& body_a = fx.build_nest(root, {"i", "j"}, {n, m}); + fx.write_scalar(body_a, "A", {symbolic::symbol("i"), symbolic::symbol("j")}, ptr, "0"); + auto& body_b = fx.build_nest(root, {"p", "q"}, {n, m}); + fx.write_scalar(body_b, "B", {symbolic::symbol("p"), symbolic::symbol("q")}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_TRUE(run(builder)); + + auto memsets = collect_memsets(root); + ASSERT_EQ(memsets.size(), 2); + std::vector targets{memset_target(root, *memsets[0]), memset_target(root, *memsets[1])}; + EXPECT_NE(std::find(targets.begin(), targets.end(), "A"), targets.end()); + EXPECT_NE(std::find(targets.begin(), targets.end(), "B"), targets.end()); +} + +// --------------------------------------------------------------------------- +// Positive cases: linearized accesses. The nest indexes a flat pointer buffer with a +// single delinearizable expression (e.g. `A[i*M + j]`) instead of one index per loop. +// These reach the memory-layout analysis' delinearization path, which the plain +// per-dimension accesses above never exercise. +// +// Note: linearization is only modeled for flat pointer buffers. A native fixed-size +// array carries its own shape and is never delinearized, so there is no +// "linearized native array" case here. +// --------------------------------------------------------------------------- + +TEST(ZeroFillToMemsetPassTest, Linearized_FullAccess_2DPointer) { + builder::StructuredSDFGBuilder builder("sdfg_lin_2d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + types::Pointer ptr(kElement); // flat pointer, layout inferred from A[i*M + j] + builder.add_container("A", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j"}, {n, m}); + auto linearized = symbolic::add(symbolic::mul(symbolic::symbol("i"), m), symbolic::symbol("j")); + fx.write_scalar(body, "A", {linearized}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_TRUE(run(builder)); + + auto& root = builder.subject().root(); + auto memsets = collect_memsets(root); + ASSERT_EQ(memsets.size(), 1); + EXPECT_EQ(memset_target(root, *memsets[0]), "A"); + EXPECT_TRUE(symbolic::eq(memsets[0]->num(), symbolic::mul(symbolic::mul(symbolic::integer(kElementBytes), n), m))); +} + +TEST(ZeroFillToMemsetPassTest, Linearized_FullAccess_3DPointer) { + builder::StructuredSDFGBuilder builder("sdfg_lin_3d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + auto k = fx.size("K"); + types::Pointer ptr(kElement); + builder.add_container("A", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j", "l"}, {n, m, k}); + // A[i*M*K + j*K + l] + auto i = symbolic::symbol("i"); + auto j = symbolic::symbol("j"); + auto l = symbolic::symbol("l"); + auto linearized = symbolic::add(symbolic::add(symbolic::mul(symbolic::mul(i, m), k), symbolic::mul(j, k)), l); + fx.write_scalar(body, "A", {linearized}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_TRUE(run(builder)); + + auto& root = builder.subject().root(); + auto memsets = collect_memsets(root); + ASSERT_EQ(memsets.size(), 1); + EXPECT_EQ(memset_target(root, *memsets[0]), "A"); + EXPECT_TRUE(symbolic:: + eq(memsets[0]->num(), + symbolic::mul(symbolic::mul(symbolic::mul(symbolic::integer(kElementBytes), n), m), k))); +} + +TEST(ZeroFillToMemsetPassTest, Linearized_FullAccess_TwoIndependent2DPointers) { + builder::StructuredSDFGBuilder builder("sdfg_lin_two_2d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + types::Pointer ptr(kElement); + builder.add_container("A", ptr, true); + builder.add_container("B", ptr, true); + + auto& root = builder.subject().root(); + auto& body_a = fx.build_nest(root, {"i", "j"}, {n, m}); + fx.write_scalar( + body_a, "A", {symbolic::add(symbolic::mul(symbolic::symbol("i"), m), symbolic::symbol("j"))}, ptr, "0" + ); + auto& body_b = fx.build_nest(root, {"p", "q"}, {n, m}); + fx.write_scalar( + body_b, "B", {symbolic::add(symbolic::mul(symbolic::symbol("p"), m), symbolic::symbol("q"))}, ptr, "0" + ); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_TRUE(run(builder)); + + auto memsets = collect_memsets(root); + ASSERT_EQ(memsets.size(), 2); + std::vector targets{memset_target(root, *memsets[0]), memset_target(root, *memsets[1])}; + EXPECT_NE(std::find(targets.begin(), targets.end(), "A"), targets.end()); + EXPECT_NE(std::find(targets.begin(), targets.end(), "B"), targets.end()); +} + +// --------------------------------------------------------------------------- +// Negative cases: the nest is left untouched. +// --------------------------------------------------------------------------- + +// The loop bound is smaller than the declared array extent, so the fill is partial. +TEST(ZeroFillToMemsetPassTest, PartialAccess_1DArray) { + builder::StructuredSDFGBuilder builder("sdfg_partial_1d_array", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + types::Array array(kElement, n); // container has N elements ... + builder.add_container("A", array); + + auto& body = fx.build_nest(builder.subject().root(), {"i"}, {m}); // ... but only M are written + fx.write_scalar(body, "A", {symbolic::symbol("i")}, array, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_FALSE(run(builder)); + EXPECT_TRUE(collect_memsets(builder.subject().root()).empty()); +} + +// The inner dimension is only partially covered. +TEST(ZeroFillToMemsetPassTest, PartialAccess_2DArray) { + builder::StructuredSDFGBuilder builder("sdfg_partial_2d_array", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + auto m2 = fx.size("M2"); + types::Array inner(kElement, m); // inner extent is M ... + types::Array array(inner, n); + builder.add_container("A", array); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j"}, {n, m2}); // ... but only M2 written + fx.write_scalar(body, "A", {symbolic::symbol("i"), symbolic::symbol("j")}, array, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_FALSE(run(builder)); + EXPECT_TRUE(collect_memsets(builder.subject().root()).empty()); +} + +// The pointee array is only partially covered by the inner loop. +TEST(ZeroFillToMemsetPassTest, PartialAccess_2DPointer) { + builder::StructuredSDFGBuilder builder("sdfg_partial_2d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + auto m2 = fx.size("M2"); + types::Array inner(kElement, m); + types::Pointer ptr(inner); + builder.add_container("A", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j"}, {n, m2}); + fx.write_scalar(body, "A", {symbolic::symbol("i"), symbolic::symbol("j")}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_FALSE(run(builder)); + EXPECT_TRUE(collect_memsets(builder.subject().root()).empty()); +} + +// Linearized negatives: same linearized shapes as the positive cases above, but the inner +// loop stops short of the linearization stride, so the delinearized inner extent (M / K) +// does not match the loop bound (M2 / K2) -> partial fill, left untouched. + +// A[i*M + j] on a flat pointer, but the inner loop only covers M2 < M of the M-wide stride. +TEST(ZeroFillToMemsetPassTest, Linearized_PartialAccess_2DPointer) { + builder::StructuredSDFGBuilder builder("sdfg_lin_partial_2d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + auto m2 = fx.size("M2"); + types::Pointer ptr(kElement); + builder.add_container("A", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j"}, {n, m2}); + auto linearized = symbolic::add(symbolic::mul(symbolic::symbol("i"), m), symbolic::symbol("j")); + fx.write_scalar(body, "A", {linearized}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_FALSE(run(builder)); + EXPECT_TRUE(collect_memsets(builder.subject().root()).empty()); +} + +// A[i*M*K + j*K + l] on a flat pointer, but the innermost loop only covers K2 < K. +TEST(ZeroFillToMemsetPassTest, Linearized_PartialAccess_3DPointer) { + builder::StructuredSDFGBuilder builder("sdfg_lin_partial_3d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + auto k = fx.size("K"); + auto k2 = fx.size("K2"); + types::Pointer ptr(kElement); + builder.add_container("A", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j", "l"}, {n, m, k2}); + auto i = symbolic::symbol("i"); + auto j = symbolic::symbol("j"); + auto l = symbolic::symbol("l"); + auto linearized = symbolic::add(symbolic::add(symbolic::mul(symbolic::mul(i, m), k), symbolic::mul(j, k)), l); + fx.write_scalar(body, "A", {linearized}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_FALSE(run(builder)); + EXPECT_TRUE(collect_memsets(builder.subject().root()).empty()); +} + +// A full linearized fill, but the stored constant is non-zero. +TEST(ZeroFillToMemsetPassTest, Linearized_NonZeroConstant) { + builder::StructuredSDFGBuilder builder("sdfg_lin_nonzero_const", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + types::Pointer ptr(kElement); + builder.add_container("A", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j"}, {n, m}); + auto linearized = symbolic::add(symbolic::mul(symbolic::symbol("i"), m), symbolic::symbol("j")); + fx.write_scalar(body, "A", {linearized}, ptr, "1.0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_FALSE(run(builder)); + EXPECT_TRUE(collect_memsets(builder.subject().root()).empty()); +} + +// A full linearized nest, but the stored value is a variable load, not a constant. +TEST(ZeroFillToMemsetPassTest, Linearized_VariableInput) { + builder::StructuredSDFGBuilder builder("sdfg_lin_variable_input", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + auto m = fx.size("M"); + types::Pointer ptr(kElement); + builder.add_container("A", ptr, true); + builder.add_container("B", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i", "j"}, {n, m}); + auto linearized = symbolic::add(symbolic::mul(symbolic::symbol("i"), m), symbolic::symbol("j")); + fx.copy_scalar(body, "B", "A", {linearized}, ptr); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_FALSE(run(builder)); + EXPECT_TRUE(collect_memsets(builder.subject().root()).empty()); +} + +// A strided write (A[2*i]) skips half of the elements: not a full fill. +TEST(ZeroFillToMemsetPassTest, PartialAccess_1DPointerStridedSubset) { + builder::StructuredSDFGBuilder builder("sdfg_partial_1d_ptr", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + types::Pointer ptr(kElement); + builder.add_container("A", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i"}, {n}); + fx.write_scalar(body, "A", {symbolic::mul(symbolic::integer(2), symbolic::symbol("i"))}, ptr, "0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_FALSE(run(builder)); + EXPECT_TRUE(collect_memsets(builder.subject().root()).empty()); +} + +// The stored constant is non-zero. +TEST(ZeroFillToMemsetPassTest, NonZeroConstant) { + builder::StructuredSDFGBuilder builder("sdfg_nonzero_const", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + types::Array array(kElement, n); + builder.add_container("A", array); + + auto& body = fx.build_nest(builder.subject().root(), {"i"}, {n}); + fx.write_scalar(body, "A", {symbolic::symbol("i")}, array, "1.0"); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_FALSE(run(builder)); + EXPECT_TRUE(collect_memsets(builder.subject().root()).empty()); +} + +// The stored value is a variable load, not a constant. +TEST(ZeroFillToMemsetPassTest, VariableInput) { + builder::StructuredSDFGBuilder builder("sdfg_variable_input", FunctionType_CPU); + NestFixture fx(builder); + + auto n = fx.size("N"); + types::Pointer ptr(kElement); + builder.add_container("A", ptr, true); + builder.add_container("B", ptr, true); + + auto& body = fx.build_nest(builder.subject().root(), {"i"}, {n}); + fx.copy_scalar(body, "B", "A", {symbolic::symbol("i")}, ptr); + + dump_sdfg(builder.subject(), "0.init"); + EXPECT_FALSE(run(builder)); + EXPECT_TRUE(collect_memsets(builder.subject().root()).empty()); +} diff --git a/python/benchmarks/npbench/cavity_flow/test_cavity_flow.py b/python/benchmarks/npbench/cavity_flow/test_cavity_flow.py index e85414ba8..fd368ba23 100644 --- a/python/benchmarks/npbench/cavity_flow/test_cavity_flow.py +++ b/python/benchmarks/npbench/cavity_flow/test_cavity_flow.py @@ -141,14 +141,14 @@ def test_cavity_flow(target): ) elif target == "sequential": verifier = SDFGVerification( - verification={"VECTORIZE": 35, "MAP": 61, "SEQUENTIAL": 30, "FOR": 4} + verification={"VECTORIZE": 34, "MAP": 60, "SEQUENTIAL": 30, "FOR": 4} ) elif target == "openmp": verifier = SDFGVerification( verification={ "VECTORIZE": 9, - "CPU_PARALLEL": 26, - "MAP": 42, + "CPU_PARALLEL": 25, + "MAP": 41, "SEQUENTIAL": 11, "FOR": 4, } diff --git a/python/benchmarks/npbench/polybench/test_adi.py b/python/benchmarks/npbench/polybench/test_adi.py index 4b530b891..a5c9f4bad 100644 --- a/python/benchmarks/npbench/polybench/test_adi.py +++ b/python/benchmarks/npbench/polybench/test_adi.py @@ -94,8 +94,8 @@ def test_adi(target): elif target == "openmp": verifier = SDFGVerification( verification={ - "CPU_PARALLEL": 1, - "VECTORIZE": 19, + "CPU_PARALLEL": 6, + "VECTORIZE": 14, "MAP": 20, "SEQUENTIAL": 5, "FOR": 5, diff --git a/python/benchmarks/npbench/polybench/test_correlation.py b/python/benchmarks/npbench/polybench/test_correlation.py index a02bb3d37..6a0051812 100644 --- a/python/benchmarks/npbench/polybench/test_correlation.py +++ b/python/benchmarks/npbench/polybench/test_correlation.py @@ -69,10 +69,10 @@ def test_correlation(target): verifier = SDFGVerification( verification={ "GEMM": 1, - "VECTORIZE": 5, + "VECTORIZE": 3, "REDUCE": 3, "CMath": 2, - "CPU_PARALLEL": 14, + "CPU_PARALLEL": 16, "MAP": 16, "SEQUENTIAL": 1, "FOR": 1, diff --git a/python/benchmarks/npbench/polybench/test_covariance.py b/python/benchmarks/npbench/polybench/test_covariance.py index 021d288b6..01af1448d 100644 --- a/python/benchmarks/npbench/polybench/test_covariance.py +++ b/python/benchmarks/npbench/polybench/test_covariance.py @@ -58,9 +58,9 @@ def test_covariance(target): verifier = SDFGVerification( verification={ "GEMM": 1, - "VECTORIZE": 3, + "VECTORIZE": 1, "REDUCE": 1, - "CPU_PARALLEL": 5, + "CPU_PARALLEL": 7, "MAP": 7, "SEQUENTIAL": 1, "FOR": 1, diff --git a/python/benchmarks/npbench/polybench/test_durbin.py b/python/benchmarks/npbench/polybench/test_durbin.py index 73c48388c..bd46023d9 100644 --- a/python/benchmarks/npbench/polybench/test_durbin.py +++ b/python/benchmarks/npbench/polybench/test_durbin.py @@ -58,9 +58,9 @@ def test_durbin(target): elif target == "openmp": verifier = SDFGVerification( verification={ - "CPU_PARALLEL": 1, + "CPU_PARALLEL": 4, "REDUCE": 1, - "VECTORIZE": 4, + "VECTORIZE": 1, "MAP": 4, "SEQUENTIAL": 1, "FOR": 1, diff --git a/python/benchmarks/npbench/polybench/test_symm.py b/python/benchmarks/npbench/polybench/test_symm.py index d2a2a1a65..0768ece3f 100644 --- a/python/benchmarks/npbench/polybench/test_symm.py +++ b/python/benchmarks/npbench/polybench/test_symm.py @@ -65,10 +65,10 @@ def test_symm(target): verifier = SDFGVerification( verification={ "GEMM": 1, - "VECTORIZE": 5, + "VECTORIZE": 2, "SEQUENTIAL": 2, "FOR": 2, - "CPU_PARALLEL": 2, + "CPU_PARALLEL": 5, "MAP": 7, }, ) diff --git a/python/benchmarks/npbench/polybench/test_syr2k.py b/python/benchmarks/npbench/polybench/test_syr2k.py index caf138c1b..54b0d28fa 100644 --- a/python/benchmarks/npbench/polybench/test_syr2k.py +++ b/python/benchmarks/npbench/polybench/test_syr2k.py @@ -41,7 +41,13 @@ def test_syr2k(target): ) elif target == "openmp": verifier = SDFGVerification( - verification={"VECTORIZE": 6, "MAP": 6, "SEQUENTIAL": 2, "FOR": 2}, + verification={ + "VECTORIZE": 4, + "MAP": 6, + "SEQUENTIAL": 2, + "FOR": 2, + "CPU_PARALLEL": 2, + }, ) elif target == "cuda": verifier = SDFGVerification( diff --git a/python/benchmarks/npbench/polybench/test_syrk.py b/python/benchmarks/npbench/polybench/test_syrk.py index d6f713e3a..3f757ce58 100644 --- a/python/benchmarks/npbench/polybench/test_syrk.py +++ b/python/benchmarks/npbench/polybench/test_syrk.py @@ -36,7 +36,13 @@ def test_syrk(target): ) elif target == "openmp": verifier = SDFGVerification( - verification={"VECTORIZE": 4, "MAP": 4, "SEQUENTIAL": 2, "FOR": 2}, + verification={ + "VECTORIZE": 2, + "MAP": 4, + "SEQUENTIAL": 2, + "FOR": 2, + "CPU_PARALLEL": 2, + }, ) elif target == "cuda": verifier = SDFGVerification( diff --git a/python/benchmarks/npbench/spmv/test_spmv.py b/python/benchmarks/npbench/spmv/test_spmv.py index cfd991c41..0d43ac31a 100644 --- a/python/benchmarks/npbench/spmv/test_spmv.py +++ b/python/benchmarks/npbench/spmv/test_spmv.py @@ -71,9 +71,9 @@ def test_spmv(target): elif target == "openmp": verifier = SDFGVerification( verification={ - "CPU_PARALLEL": 1, + "CPU_PARALLEL": 4, "REDUCE": 1, - "VECTORIZE": 4, + "VECTORIZE": 1, "MAP": 4, "SEQUENTIAL": 1, "FOR": 1, diff --git a/python/benchmarks/npbench/weather_stencils/test_vadv.py b/python/benchmarks/npbench/weather_stencils/test_vadv.py index fc866c3ce..8a81d19e8 100644 --- a/python/benchmarks/npbench/weather_stencils/test_vadv.py +++ b/python/benchmarks/npbench/weather_stencils/test_vadv.py @@ -135,7 +135,13 @@ def test_vadv(target): ) elif target == "openmp": verifier = SDFGVerification( - verification={"VECTORIZE": 34, "MAP": 66, "SEQUENTIAL": 39, "FOR": 7} + verification={ + "VECTORIZE": 2, + "MAP": 34, + "SEQUENTIAL": 7, + "FOR": 7, + "CPU_PARALLEL": 32, + } ) elif target == "cuda": verifier = SDFGVerification( diff --git a/python/bindings/py_structured_sdfg.cpp b/python/bindings/py_structured_sdfg.cpp index bb7269ae8..10cc04ab9 100644 --- a/python/bindings/py_structured_sdfg.cpp +++ b/python/bindings/py_structured_sdfg.cpp @@ -44,6 +44,7 @@ #include #include #include +#include #include #include @@ -451,6 +452,9 @@ void PyStructuredSDFG::normalize() { sdfg::passes::TaskletFusionPass task_fuse_pass; task_fuse_pass.run(builder, analysis_manager); } + + auto zero_fill_to_memset_pass = sdfg::passes::ZeroFillToMemsetPass(); + zero_fill_to_memset_pass.run(builder, analysis_manager); } void PyStructuredSDFG::schedule(const std::string& target, const std::string& category, bool remote_tuning) { diff --git a/sdfg/src/analysis/memory_layout_analysis.cpp b/sdfg/src/analysis/memory_layout_analysis.cpp index 3278f7fad..c81230ab1 100644 --- a/sdfg/src/analysis/memory_layout_analysis.cpp +++ b/sdfg/src/analysis/memory_layout_analysis.cpp @@ -164,16 +164,22 @@ void MemoryLayoutAnalysis:: continue; } case types::TypeID::Array: { - // Arrays are c-like stack array, so we can infer a simple row-major layout without needing - // delinearization + // Arrays are c-like stack arrays, so we can infer a simple row-major layout without needing + // delinearization. Collect the extent of every (nested) array dimension so that + // multi-dimensional arrays get a full shape instead of just their leading dimension. auto* array_type = dynamic_cast(&memlet.base_type()); - symbolic::MultiExpression shape = {array_type->num_elements()}; + symbolic::MultiExpression shape; + shape.push_back(array_type->num_elements()); while (array_type->element_type().type_id() == types::TypeID::Array) { array_type = dynamic_cast(&array_type->element_type()); + shape.push_back(array_type->num_elements()); } if (array_type->element_type().type_id() != types::TypeID::Scalar) { continue; // Skip non-scalar arrays } + if (subset.size() != shape.size()) { + continue; // Require one index per (native) array dimension + } MemoryLayout layout(shape); MemoryAccess layout_info{container_name, subset, layout, true};