15494289007

Committed 06 Jun 2025 03:36PM UTC coverage: 57.304% (-0.4%) from 57.704%

Build # 15494289007

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #60 from daisytuner/kernels

removes kernel node in favor of function types

Run Details

78 of 99 new or added lines in 11 files covered. (78.79%)

91 existing lines in 14 files now uncovered.

7583 of 13233 relevant lines covered (57.3%)

116.04 hits per line

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

15.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/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::For& loop)
    : parent_(parent), loop_(loop) {};

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

bool LoopToKernelDim::can_be_applied(Schedule& schedule) {
    auto& analysis_manager = schedule.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(schedule.builder(), analysis_manager);
    auto& builder = schedule.builder();

    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(Schedule& schedule) {
    auto& analysis_manager = schedule.analysis_manager();
    auto& builder = schedule.builder();

    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(schedule.builder(), analysis_manager);
        applies |= sf_pass.run(schedule.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/symbolic/symbolic.h"
14	#include "sdfg/transformations/utils.h"
15	#include "symengine/symengine_rcp.h"
16
17	namespace sdfg {
18	namespace transformations {
19
20	LoopToKernelDim::LoopToKernelDim(structured_control_flow::Sequence& parent,	3✔
21	structured_control_flow::For& loop)
22	: parent_(parent), loop_(loop) {};	3✔
23
24	std::string LoopToKernelDim::name() { return "LoopToKernelDim"; };	×
25
26	bool LoopToKernelDim::can_be_applied(Schedule& schedule) {	3✔
27	auto& analysis_manager = schedule.analysis_manager();	3✔
28	sdfg::passes::Pipeline expression_combine = sdfg::passes::Pipeline::expression_combine();	3✔
29	sdfg::passes::Pipeline memlet_combine = sdfg::passes::Pipeline::memlet_combine();	3✔
30	memlet_combine.run(schedule.builder(), analysis_manager);	3✔
31	auto& builder = schedule.builder();	3✔
32
33	auto& sdfg = builder.subject();	3✔
34	if (sdfg.type() != FunctionType::NV_GLOBAL) {	3✔
UNCOV 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	}
NEW 60	auto y_dim_size = assumptions[symbolic::blockDim_y()].integer_value();	×
UNCOV 61	if (y_dim_size == SymEngine::null) {	×
62	return false;	×
63	}
NEW 64	auto z_dim_size = assumptions[symbolic::blockDim_z()].integer_value();	×
UNCOV 65	if (z_dim_size == SymEngine::null) {	×
66	return false;	×
67	}
68
69	// Criterion: Available kernel dimensions is free
UNCOV 70	bool x_dim_available = false;	×
UNCOV 71	bool y_dim_available = false;	×
UNCOV 72	bool z_dim_available = false;	×
73
UNCOV 74	if (!symbolic::eq(x_dim_size, symbolic::integer(1))) {	×
NEW 75	if (body_users.reads("threadIdx.x").empty()) {	×
76	x_dim_available = true;	×
77	}	×
UNCOV 78	}	×
UNCOV 79	if (!symbolic::eq(y_dim_size, symbolic::integer(1))) {	×
NEW 80	if (body_users.reads("threadIdx.y").empty()) {	×
UNCOV 81	y_dim_available = true;	×
UNCOV 82	}	×
UNCOV 83	}	×
UNCOV 84	if (!symbolic::eq(z_dim_size, symbolic::integer(1))) {	×
NEW 85	if (body_users.reads("threadIdx.z").empty()) {	×
86	z_dim_available = true;	×
87	}	×
88	}	×
UNCOV 89	if (!x_dim_available && !y_dim_available && !z_dim_available) {	×
UNCOV 90	return false;	×
91	}
92
93	// Criterion: Unused kernel dimension is bigger or equal to loop iteration count
UNCOV 94	bool x_match = false;	×
UNCOV 95	bool y_match = false;	×
UNCOV 96	bool z_match = false;	×
UNCOV 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	}	×
UNCOV 103	if (y_dim_available) {	×
UNCOV 104	auto cond = symbolic::Ge(y_dim_size, iteration_count);	×
UNCOV 105	if (symbolic::is_true(cond)) {	×
UNCOV 106	y_match = true;	×
UNCOV 107	}	×
UNCOV 108	}	×
UNCOV 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	}	×
UNCOV 115	if (!x_match && !y_match && !z_match) {	×
UNCOV 116	return false;	×
117	}
118
119	// Criterion: Loop is a map
UNCOV 120	auto& analysis = analysis_manager.get<analysis::DataParallelismAnalysis>();	×
UNCOV 121	auto& data_dependencies = analysis.get(this->loop_);	×
UNCOV 122	for (auto& dep : data_dependencies) {	×
UNCOV 123	auto& dep_type = dep.second;	×
UNCOV 124	if (dep_type < analysis::Parallelism::PARALLEL) {	×
UNCOV 125	return false;	×
126	}
127	}
128
129	return true;	×
130	};	3✔
131
132	void LoopToKernelDim::apply(Schedule& schedule) {	×
133	auto& analysis_manager = schedule.analysis_manager();	×
134	auto& builder = schedule.builder();	×
135
136	auto& assumptions_analysis = analysis_manager.get<analysis::AssumptionsAnalysis>();	×
137	auto assumptions = assumptions_analysis.get(loop_.root());	×
NEW 138	auto x_dim_size = assumptions[symbolic::blockDim_x()].integer_value();	×
NEW 139	auto y_dim_size = assumptions[symbolic::blockDim_y()].integer_value();	×
NEW 140	auto z_dim_size = assumptions[symbolic::blockDim_z()].integer_value();	×
141
142	bool x_dim_available = false;	×
143	bool y_dim_available = false;	×
144	auto& users = analysis_manager.get<analysis::Users>();	×
145	analysis::UsersView body_users(users, loop_.root());	×
146
147	if (!symbolic::eq(x_dim_size, symbolic::integer(1))) {	×
NEW 148	if (body_users.reads("threadIdx.x").empty()) {	×
149	x_dim_available = true;	×
150	}	×
151	}	×
152	if (!symbolic::eq(y_dim_size, symbolic::integer(1))) {	×
NEW 153	if (body_users.reads("threadIdx.y").empty()) {	×
154	y_dim_available = true;	×
155	}	×
156	}	×
157
158	bool x_match = false;	×
159	bool y_match = false;	×
160	symbolic::Integer iteration_count = get_iteration_count(loop_);	×
161	if (x_dim_available) {	×
162	auto cond = symbolic::Ge(x_dim_size, iteration_count);	×
163	if (symbolic::is_true(cond)) {	×
164	x_match = true;	×
165	}	×
166	}	×
167	if (y_dim_available) {	×
168	auto cond = symbolic::Ge(y_dim_size, iteration_count);	×
169	if (symbolic::is_true(cond)) {	×
170	y_match = true;	×
171	}	×
172	}	×
173
NEW 174	auto target_dim = symbolic::threadIdx_z();	×
175	if (x_match) {	×
NEW 176	target_dim = symbolic::threadIdx_x();	×
177	} else if (y_match) {	×
NEW 178	target_dim = symbolic::threadIdx_y();	×
179	}	×
180
181	auto& parent = builder.parent(loop_);	×
182	auto& if_else = builder.add_if_else_before(parent, loop_).first;	×
183	auto condition =
184	symbolic::subs(loop_.condition(), loop_.indvar(), symbolic::add(target_dim, loop_.init()));	×
185	auto& branch = builder.add_case(if_else, condition);	×
186
187	builder.insert_children(branch, loop_.root(), 0);	×
188	branch.replace(loop_.indvar(), symbolic::add(target_dim, loop_.init()));	×
189	builder.remove_child(parent, loop_);	×
190
191	analysis_manager.invalidate_all();	×
192	passes::SequenceFusion sf_pass;	×
193	passes::DeadCFGElimination dce_pass;	×
194	bool applies = false;	×
195	do {	×
196	applies = false;	×
197	applies \|= dce_pass.run(schedule.builder(), analysis_manager);	×
198	applies \|= sf_pass.run(schedule.builder(), analysis_manager);	×
199	} while (applies);	×
200	};	×
201
202	} // namespace transformations
203	} // namespace sdfg

daisytuner / sdfglib / 15494289007

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