15512392058

Committed 07 Jun 2025 10:47PM UTC coverage: 62.261% (+4.8%) from 57.416%

Build # 15512392058

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #63 from daisytuner/schedules

Schedules

Run Details

152 of 194 new or added lines in 24 files covered. (78.35%)

39 existing lines in 7 files now uncovered.

7467 of 11993 relevant lines covered (62.26%)

127.08 hits per line

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

14.49

/src/transformations/loop_to_kernel_dim.cpp

#include "sdfg/transformations/loop_to_kernel_dim.h"

#include <utility>

#include "sdfg/analysis/assumptions_analysis.h"
#include "sdfg/analysis/data_parallelism_analysis.h"
#include "sdfg/builder/structured_sdfg_builder.h"
#include "sdfg/passes/pipeline.h"
#include "sdfg/passes/structured_control_flow/dead_cfg_elimination.h"
#include "sdfg/passes/structured_control_flow/sequence_fusion.h"
#include "sdfg/structured_control_flow/if_else.h"
#include "sdfg/structured_control_flow/sequence.h"
#include "sdfg/structured_control_flow/structured_loop.h"
#include "sdfg/symbolic/symbolic.h"
#include "sdfg/transformations/utils.h"
#include "symengine/symengine_rcp.h"

namespace sdfg {
namespace transformations {

LoopToKernelDim::LoopToKernelDim(structured_control_flow::Sequence& parent,
                                 structured_control_flow::StructuredLoop& loop)
    : parent_(parent), loop_(loop) {};

std::string LoopToKernelDim::name() { return "LoopToKernelDim"; };

bool LoopToKernelDim::can_be_applied(builder::StructuredSDFGBuilder& builder,
                                     analysis::AnalysisManager& analysis_manager) {
    sdfg::passes::Pipeline expression_combine = sdfg::passes::Pipeline::expression_combine();
    sdfg::passes::Pipeline memlet_combine = sdfg::passes::Pipeline::memlet_combine();
    memlet_combine.run(builder, analysis_manager);

    auto& sdfg = builder.subject();
    if (sdfg.type() != FunctionType_NV_GLOBAL) {
        return false;
    }

    // Criterion: Iteration count is known and an Integer
    symbolic::Integer iteration_count = get_iteration_count(loop_);
    if (iteration_count == SymEngine::null) {
        return false;
    }

    // Criterion: Indvar is only used to access arrays
    auto& users = analysis_manager.get<analysis::Users>();
    analysis::UsersView body_users(users, loop_.root());
    for (auto user : body_users.reads(loop_.indvar()->get_name())) {
        if (dynamic_cast<data_flow::AccessNode*>(user->element())) {
            return false;
        }
    }

    // Criterion: Kernel dimensions are known and an Integer
    auto& assumptions_analysis = analysis_manager.get<analysis::AssumptionsAnalysis>();
    auto assumptions = assumptions_analysis.get(loop_.root());
    auto x_dim_size = assumptions[symbolic::blockDim_x()].integer_value();
    if (x_dim_size == SymEngine::null) {
        return false;
    }
    auto y_dim_size = assumptions[symbolic::blockDim_y()].integer_value();
    if (y_dim_size == SymEngine::null) {
        return false;
    }
    auto z_dim_size = assumptions[symbolic::blockDim_z()].integer_value();
    if (z_dim_size == SymEngine::null) {
        return false;
    }

    // Criterion: Available kernel dimensions is free
    bool x_dim_available = false;
    bool y_dim_available = false;
    bool z_dim_available = false;

    if (!symbolic::eq(x_dim_size, symbolic::integer(1))) {
        if (body_users.reads("threadIdx.x").empty()) {
            x_dim_available = true;
        }
    }
    if (!symbolic::eq(y_dim_size, symbolic::integer(1))) {
        if (body_users.reads("threadIdx.y").empty()) {
            y_dim_available = true;
        }
    }
    if (!symbolic::eq(z_dim_size, symbolic::integer(1))) {
        if (body_users.reads("threadIdx.z").empty()) {
            z_dim_available = true;
        }
    }
    if (!x_dim_available && !y_dim_available && !z_dim_available) {
        return false;
    }

    // Criterion: Unused kernel dimension is bigger or equal to loop iteration count
    bool x_match = false;
    bool y_match = false;
    bool z_match = false;
    if (x_dim_available) {
        auto cond = symbolic::Ge(x_dim_size, iteration_count);
        if (symbolic::is_true(cond)) {
            x_match = true;
        }
    }
    if (y_dim_available) {
        auto cond = symbolic::Ge(y_dim_size, iteration_count);
        if (symbolic::is_true(cond)) {
            y_match = true;
        }
    }
    if (z_dim_available) {
        auto cond = symbolic::Ge(z_dim_size, iteration_count);
        if (symbolic::is_true(cond)) {
            z_match = true;
        }
    }
    if (!x_match && !y_match && !z_match) {
        return false;
    }

    // Criterion: Loop is a map
    auto& analysis = analysis_manager.get<analysis::DataParallelismAnalysis>();
    auto& data_dependencies = analysis.get(this->loop_);
    for (auto& dep : data_dependencies) {
        auto& dep_type = dep.second;
        if (dep_type < analysis::Parallelism::PARALLEL) {
            return false;
        }
    }

    return true;
};

void LoopToKernelDim::apply(builder::StructuredSDFGBuilder& builder,
                            analysis::AnalysisManager& analysis_manager) {
    auto& assumptions_analysis = analysis_manager.get<analysis::AssumptionsAnalysis>();
    auto assumptions = assumptions_analysis.get(loop_.root());
    auto x_dim_size = assumptions[symbolic::blockDim_x()].integer_value();
    auto y_dim_size = assumptions[symbolic::blockDim_y()].integer_value();
    auto z_dim_size = assumptions[symbolic::blockDim_z()].integer_value();

    bool x_dim_available = false;
    bool y_dim_available = false;
    auto& users = analysis_manager.get<analysis::Users>();
    analysis::UsersView body_users(users, loop_.root());

    if (!symbolic::eq(x_dim_size, symbolic::integer(1))) {
        if (body_users.reads("threadIdx.x").empty()) {
            x_dim_available = true;
        }
    }
    if (!symbolic::eq(y_dim_size, symbolic::integer(1))) {
        if (body_users.reads("threadIdx.y").empty()) {
            y_dim_available = true;
        }
    }

    bool x_match = false;
    bool y_match = false;
    symbolic::Integer iteration_count = get_iteration_count(loop_);
    if (x_dim_available) {
        auto cond = symbolic::Ge(x_dim_size, iteration_count);
        if (symbolic::is_true(cond)) {
            x_match = true;
        }
    }
    if (y_dim_available) {
        auto cond = symbolic::Ge(y_dim_size, iteration_count);
        if (symbolic::is_true(cond)) {
            y_match = true;
        }
    }

    auto target_dim = symbolic::threadIdx_z();
    if (x_match) {
        target_dim = symbolic::threadIdx_x();
    } else if (y_match) {
        target_dim = symbolic::threadIdx_y();
    }

    auto& parent = builder.parent(loop_);
    auto& if_else = builder.add_if_else_before(parent, loop_).first;
    auto condition =
        symbolic::subs(loop_.condition(), loop_.indvar(), symbolic::add(target_dim, loop_.init()));
    auto& branch = builder.add_case(if_else, condition);

    builder.insert_children(branch, loop_.root(), 0);
    branch.replace(loop_.indvar(), symbolic::add(target_dim, loop_.init()));
    builder.remove_child(parent, loop_);

    analysis_manager.invalidate_all();
    passes::SequenceFusion sf_pass;
    passes::DeadCFGElimination dce_pass;
    bool applies = false;
    do {
        applies = false;
        applies |= dce_pass.run(builder, analysis_manager);
        applies |= sf_pass.run(builder, analysis_manager);
    } while (applies);
};

}  // namespace transformations
}  // namespace sdfg

1	#include "sdfg/transformations/loop_to_kernel_dim.h"
2
3	#include <utility>
4
5	#include "sdfg/analysis/assumptions_analysis.h"
6	#include "sdfg/analysis/data_parallelism_analysis.h"
7	#include "sdfg/builder/structured_sdfg_builder.h"
8	#include "sdfg/passes/pipeline.h"
9	#include "sdfg/passes/structured_control_flow/dead_cfg_elimination.h"
10	#include "sdfg/passes/structured_control_flow/sequence_fusion.h"
11	#include "sdfg/structured_control_flow/if_else.h"
12	#include "sdfg/structured_control_flow/sequence.h"
13	#include "sdfg/structured_control_flow/structured_loop.h"
14	#include "sdfg/symbolic/symbolic.h"
15	#include "sdfg/transformations/utils.h"
16	#include "symengine/symengine_rcp.h"
17
18	namespace sdfg {
19	namespace transformations {
20
21	LoopToKernelDim::LoopToKernelDim(structured_control_flow::Sequence& parent,	3✔
22	structured_control_flow::StructuredLoop& loop)
23	: parent_(parent), loop_(loop) {};	3✔
24
25	std::string LoopToKernelDim::name() { return "LoopToKernelDim"; };	×
26
27	bool LoopToKernelDim::can_be_applied(builder::StructuredSDFGBuilder& builder,	3✔
28	analysis::AnalysisManager& analysis_manager) {
29	sdfg::passes::Pipeline expression_combine = sdfg::passes::Pipeline::expression_combine();	3✔
30	sdfg::passes::Pipeline memlet_combine = sdfg::passes::Pipeline::memlet_combine();	3✔
31	memlet_combine.run(builder, analysis_manager);	3✔
32
33	auto& sdfg = builder.subject();	3✔
34	if (sdfg.type() != FunctionType_NV_GLOBAL) {	3✔
35	return false;	×
36	}
37
38	// Criterion: Iteration count is known and an Integer
39	symbolic::Integer iteration_count = get_iteration_count(loop_);	3✔
40	if (iteration_count == SymEngine::null) {	3✔
41	return false;	×
42	}
43
44	// Criterion: Indvar is only used to access arrays
45	auto& users = analysis_manager.get<analysis::Users>();	3✔
46	analysis::UsersView body_users(users, loop_.root());	3✔
47	for (auto user : body_users.reads(loop_.indvar()->get_name())) {	8✔
48	if (dynamic_cast<data_flow::AccessNode*>(user->element())) {	5✔
49	return false;	×
50	}
51	}
52
53	// Criterion: Kernel dimensions are known and an Integer
54	auto& assumptions_analysis = analysis_manager.get<analysis::AssumptionsAnalysis>();	3✔
55	auto assumptions = assumptions_analysis.get(loop_.root());	3✔
56	auto x_dim_size = assumptions[symbolic::blockDim_x()].integer_value();	3✔
57	if (x_dim_size == SymEngine::null) {	3✔
58	return false;	3✔
59	}
60	auto y_dim_size = assumptions[symbolic::blockDim_y()].integer_value();	×
61	if (y_dim_size == SymEngine::null) {	×
62	return false;	×
63	}
64	auto z_dim_size = assumptions[symbolic::blockDim_z()].integer_value();	×
65	if (z_dim_size == SymEngine::null) {	×
66	return false;	×
67	}
68
69	// Criterion: Available kernel dimensions is free
70	bool x_dim_available = false;	×
71	bool y_dim_available = false;	×
72	bool z_dim_available = false;	×
73
74	if (!symbolic::eq(x_dim_size, symbolic::integer(1))) {	×
75	if (body_users.reads("threadIdx.x").empty()) {	×
76	x_dim_available = true;	×
77	}	×
78	}	×
79	if (!symbolic::eq(y_dim_size, symbolic::integer(1))) {	×
80	if (body_users.reads("threadIdx.y").empty()) {	×
81	y_dim_available = true;	×
82	}	×
83	}	×
84	if (!symbolic::eq(z_dim_size, symbolic::integer(1))) {	×
85	if (body_users.reads("threadIdx.z").empty()) {	×
86	z_dim_available = true;	×
87	}	×
88	}	×
89	if (!x_dim_available && !y_dim_available && !z_dim_available) {	×
90	return false;	×
91	}
92
93	// Criterion: Unused kernel dimension is bigger or equal to loop iteration count
94	bool x_match = false;	×
95	bool y_match = false;	×
96	bool z_match = false;	×
97	if (x_dim_available) {	×
98	auto cond = symbolic::Ge(x_dim_size, iteration_count);	×
99	if (symbolic::is_true(cond)) {	×
100	x_match = true;	×
101	}	×
102	}	×
103	if (y_dim_available) {	×
104	auto cond = symbolic::Ge(y_dim_size, iteration_count);	×
105	if (symbolic::is_true(cond)) {	×
106	y_match = true;	×
107	}	×
108	}	×
109	if (z_dim_available) {	×
110	auto cond = symbolic::Ge(z_dim_size, iteration_count);	×
111	if (symbolic::is_true(cond)) {	×
112	z_match = true;	×
113	}	×
114	}	×
115	if (!x_match && !y_match && !z_match) {	×
116	return false;	×
117	}
118
119	// Criterion: Loop is a map
120	auto& analysis = analysis_manager.get<analysis::DataParallelismAnalysis>();	×
121	auto& data_dependencies = analysis.get(this->loop_);	×
122	for (auto& dep : data_dependencies) {	×
123	auto& dep_type = dep.second;	×
124	if (dep_type < analysis::Parallelism::PARALLEL) {	×
125	return false;	×
126	}
127	}
128
129	return true;	×
130	};	3✔
131
NEW 132	void LoopToKernelDim::apply(builder::StructuredSDFGBuilder& builder,	×
133	analysis::AnalysisManager& analysis_manager) {
134	auto& assumptions_analysis = analysis_manager.get<analysis::AssumptionsAnalysis>();	×
135	auto assumptions = assumptions_analysis.get(loop_.root());	×
136	auto x_dim_size = assumptions[symbolic::blockDim_x()].integer_value();	×
137	auto y_dim_size = assumptions[symbolic::blockDim_y()].integer_value();	×
138	auto z_dim_size = assumptions[symbolic::blockDim_z()].integer_value();	×
139
140	bool x_dim_available = false;	×
141	bool y_dim_available = false;	×
142	auto& users = analysis_manager.get<analysis::Users>();	×
143	analysis::UsersView body_users(users, loop_.root());	×
144
145	if (!symbolic::eq(x_dim_size, symbolic::integer(1))) {	×
146	if (body_users.reads("threadIdx.x").empty()) {	×
147	x_dim_available = true;	×
148	}	×
149	}	×
150	if (!symbolic::eq(y_dim_size, symbolic::integer(1))) {	×
151	if (body_users.reads("threadIdx.y").empty()) {	×
152	y_dim_available = true;	×
153	}	×
154	}	×
155
156	bool x_match = false;	×
157	bool y_match = false;	×
158	symbolic::Integer iteration_count = get_iteration_count(loop_);	×
159	if (x_dim_available) {	×
160	auto cond = symbolic::Ge(x_dim_size, iteration_count);	×
161	if (symbolic::is_true(cond)) {	×
162	x_match = true;	×
163	}	×
164	}	×
165	if (y_dim_available) {	×
166	auto cond = symbolic::Ge(y_dim_size, iteration_count);	×
167	if (symbolic::is_true(cond)) {	×
168	y_match = true;	×
169	}	×
170	}	×
171
172	auto target_dim = symbolic::threadIdx_z();	×
173	if (x_match) {	×
174	target_dim = symbolic::threadIdx_x();	×
175	} else if (y_match) {	×
176	target_dim = symbolic::threadIdx_y();	×
177	}	×
178
179	auto& parent = builder.parent(loop_);	×
180	auto& if_else = builder.add_if_else_before(parent, loop_).first;	×
181	auto condition =
182	symbolic::subs(loop_.condition(), loop_.indvar(), symbolic::add(target_dim, loop_.init()));	×
183	auto& branch = builder.add_case(if_else, condition);	×
184
185	builder.insert_children(branch, loop_.root(), 0);	×
186	branch.replace(loop_.indvar(), symbolic::add(target_dim, loop_.init()));	×
187	builder.remove_child(parent, loop_);	×
188
189	analysis_manager.invalidate_all();	×
190	passes::SequenceFusion sf_pass;	×
191	passes::DeadCFGElimination dce_pass;	×
192	bool applies = false;	×
193	do {	×
194	applies = false;	×
NEW 195	applies \|= dce_pass.run(builder, analysis_manager);	×
NEW 196	applies \|= sf_pass.run(builder, analysis_manager);	×
197	} while (applies);	×
198	};	×
199
200	} // namespace transformations
201	} // namespace sdfg

daisytuner / sdfglib / 15512392058

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