Skip to content

GitLab

Unverified Commit 3e6d18a4 authored by Paul Fultz II's avatar Paul Fultz II Committed by GitHub
Browse files

Merge pull request #48 from ROCmSoftwarePlatform/revert-46-memory_coloring 

Revert "Memory coloring"
parents 7a2edaf4 ef99eda3
Hide whitespace changes
Inline Side-by-side
Showing with 9 additions and 575 deletions
src/CMakeLists.txt
View file @ 3e6d18a4
......@@ -9,8 +9,6 @@ add_library(migraph
program.cpp
shape.cpp
simplify_reshapes.cpp
opt/memory_coloring.cpp
opt/memory_coloring_impl.cpp
)
rocm_clang_tidy_check(migraph)
target_include_directories(migraph PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)
......@@ -20,9 +18,3 @@ add_subdirectory(targets/cpu)
if(MIGRAPH_ENABLE_GPU)
add_subdirectory(targets/gpu)
endif()
#install (TARGETS migraph
# LIBRARY DESTINATION /opt/rocm/lib)
#install (DIRECTORY include/migraph DESTINATION /opt/rocm/include)
src/include/migraph/memory_coloring.hpp deleted 100644 → 0
View file @ 7a2edaf4
#ifndef MIGRAPH_GUARD_RTGLIB_MEMORY_COLORING_HPP
#define MIGRAPH_GUARD_RTGLIB_MEMORY_COLORING_HPP
#include <string>
#include <migraph/instruction_ref.hpp>
namespace migraph {
struct program;
struct memory_coloring
{
std::string name() const { return "memory coloring"; }
void apply(program& p) const;
};
} // namespace migraph
#endif
src/include/migraph/operators.hpp
View file @ 3e6d18a4
......@@ -534,27 +534,6 @@ struct div : binary
std::string name() const { return "div"; }
};
struct get_mem_ptr
{
std::string name() const { return "get_mem_ptr:" + std::to_string(offset); }
shape compute_shape(std::vector<shape> inputs) const { return inputs.at(1); }
argument compute(context&, const shape& output_shape, const std::vector<argument>& args) const
{
return {std::move(output_shape), args.at(0).data() + offset};
}
std::size_t offset = 0;
};
struct write_literal
{
std::string name() const { return "write_literal"; }
shape compute_shape(std::vector<shape> inputs) const { return inputs.at(2); }
argument compute(context&, const shape&, const std::vector<argument>&) const
{
MIGRAPH_THROW("not computable");
}
};
struct outline
{
shape s;
......@@ -569,6 +548,7 @@ struct outline
return {s, nullptr};
}
};
} // namespace migraph
#endif
src/include/migraph/program.hpp
View file @ 3e6d18a4
......@@ -100,6 +100,7 @@ struct program
private:
std::unique_ptr<program_impl> impl;
};
} // namespace migraph
#endif
src/opt/common_header.hpp deleted 100644 → 0
View file @ 7a2edaf4
#ifndef MIGRAPH_GUARD_RTGLIB_COMMON_HEADER_HPP
#define MIGRAPH_GUARD_RTGLIB_COMMON_HEADER_HPP
#include <migraph/program.hpp>
#include <migraph/stringutils.hpp>
#include <migraph/instruction.hpp>
#include <migraph/operators.hpp>
#include <migraph/iterator_for.hpp>
#include <set>
#include <list>
#include <vector>
#include <queue>
//#define MIGRAPH_DEBUG_OPT
#ifdef MIGRAPH_DEBUG_OPT
#define MIGRAPH_DEBUG(s) s
#else
#define MIGRAPH_DEBUG(s)
#endif // MIGRAPH_DEBUG_OPT
#endif // MIGRAPH_GUARD_RTGLIB_COMMON_HEADER_HPP
src/opt/memory_coloring.cpp deleted 100644 → 0
View file @ 7a2edaf4
#include <migraph/memory_coloring.hpp>
#include "memory_coloring_impl.hpp"
namespace migraph {
void memory_coloring::apply(program& p) const
{
memory_coloring_impl opt(&p);
opt.run();
}
} // namespace migraph
src/opt/memory_coloring_impl.cpp deleted 100644 → 0
View file @ 7a2edaf4
#include "memory_coloring_impl.hpp"
namespace migraph {
void memory_coloring_impl::run()
{
build();
if(num_of_lives != 0)
{
MIGRAPH_DEBUG(dump("---Before memory coloring---"));
MIGRAPH_DEBUG(dump_program());
MIGRAPH_DEBUG(dump_intervals());
// Coloring
while(!alloc_queue.empty())
{
interval_ptr interval = alloc_queue.top();
allocate(interval);
alloc_queue.pop();
}
rewrite();
MIGRAPH_DEBUG(verify());
}
}
bool memory_coloring_impl::allocate(interval_ptr interval)
{
shape s = interval->result;
std::size_t size = s.bytes();
if(size == 0)
return false;
std::size_t element_size = size / s.elements();
live_range& segment = interval->segment;
int vn = segment.vn;
std::priority_queue<live_range*, std::vector<live_range*>, ordering> conflict_queue;
std::unordered_map<long long, live_range*> offset2_live;
offset2_live.clear();
if(conflict_table.find(vn) != conflict_table.end())
{
std::set<int>& vn_set = conflict_table[vn];
for(auto& iter : vn_set)
{
live_range* range = live_ranges[iter];
long long offset = range->offset;
if(offset != InvalidOffset)
{
conflict_queue.push(range);
if(offset2_live.find(offset) == offset2_live.end())
{
offset2_live[offset] = range;
}
else
{
live_range* prev = offset2_live[offset];
assert(prev->offset == offset);
if(prev->size < range->size)
offset2_live[offset] = range;
}
}
}
}
long long offset = 0;
while(!conflict_queue.empty())
{
live_range* range = conflict_queue.top();
long long cur_offset = range->offset;
if(offset2_live[cur_offset] == range)
{
if((cur_offset > offset) && (cur_offset - offset) >= size)
{
break;
}
offset = cur_offset + range->size;
if((offset % element_size) != 0)
offset += (element_size - (offset % element_size));
}
conflict_queue.pop();
}
segment.offset = offset;
MIGRAPH_DEBUG(segment.dump());
required_bytes = std::max(required_bytes, offset + segment.size);
return true;
}
void memory_coloring_impl::build()
{
std::size_t num_of_instrs = p_program->size();
if(num_of_instrs == 0)
return;
int cur_points = num_of_instrs * 2;
instruction_ref iter = std::prev(p_program->end());
instruction_ref begin = p_program->begin();
std::vector<instruction_ref> dead_instrs;
std::set<int> live_set;
// Build live intervals.
live_intervals.resize(num_of_instrs);
do
{
const instruction* p_iter = &(*iter);
interval_ptr def_interval = nullptr;
bool is_dead = false;
if(instr2_live.find(p_iter) != instr2_live.end())
{
def_interval = instr2_live[p_iter];
bool is_lit = is_literal(iter);
if(is_allocate(iter) || is_lit)
{
live_range& range = def_interval->segment;
def_interval->result = iter->result;
def_interval->is_literal = is_lit;
alloc_queue.push(def_interval);
range.begin = cur_points;
range.size = (iter->result).bytes();
live_set.erase(range.vn);
}
}
else if(!is_param(iter) && !is_outline(iter) && !is_check_context(iter))
{
is_dead = true;
}
int tie_ndx = get_input_tie_ndx(iter);
int cnt = -1;
for(auto&& arg : iter->arguments)
{
cnt++;
if(is_param(arg) || is_outline(arg))
{
if(is_output_param(arg))
is_dead = false;
continue;
}
const instruction* p_arg = &(*arg);
if(cnt == tie_ndx)
{
// input memory is used as this instruction's output.
// def is considered as use. Coalesce the live intervals.
assert(def_interval != nullptr);
def_interval->add_use(cur_points);
instr2_live[p_arg] = def_interval;
}
else if(instr2_live.find(p_arg) == instr2_live.end())
{
// First time see a use, create a live interval.
int id = num_of_lives++;
interval_ptr interval = &(live_intervals[id]);
interval->id = id;
interval->segment.end = cur_points;
interval->segment.vn = ++max_value_number;
interval->add_use(cur_points);
instr2_live[p_arg] = interval;
add_conflicts(live_set, max_value_number);
live_set.insert(max_value_number);
live_ranges[max_value_number] = &(interval->segment);
}
else
{
interval_ptr interval = instr2_live[p_arg];
interval->add_use(cur_points);
assert(live_set.find(interval->id) != live_set.end());
}
}
if(is_dead)
dead_instrs.push_back(iter);
cur_points -= 2;
iter = std::prev(iter);
} while(iter != begin);
}
void memory_coloring_impl::rewrite()
{
instruction_ref end = p_program->end();
instruction_ref scratch_param = end;
std::vector<std::size_t> dims;
dims.push_back(required_bytes / sizeof(float));
shape s = {shape::float_type, dims};
scratch_param = p_program->add_parameter("scratch", s);
for(auto ins : iterator_for(*p_program))
{
const instruction* p_iter = &(*ins);
if(instr2_live.find(p_iter) != instr2_live.end())
{
interval_ptr interval = instr2_live[p_iter];
if(interval->get_offset() == InvalidOffset)
{
assert((interval->get_begin() == InvalidOffset) || interval->result.bytes() == 0);
continue;
}
std::size_t offset = interval->get_offset();
if(is_allocate(ins))
{
p_program->replace_instruction(
ins, get_mem_ptr{offset}, scratch_param, ins->arguments.at(0));
}
else if(is_literal(ins))
{
auto pre = p_program->add_literal(ins->lit);
auto index = p_program->add_literal(offset);
p_program->replace_instruction(ins, write_literal{}, scratch_param, index, pre);
}
}
}
MIGRAPH_DEBUG(dump("---After rewrite---"));
MIGRAPH_DEBUG(dump_program());
}
#ifdef MIGRAPH_DEBUG_OPT
void memory_coloring_impl::dump(const std::string str) { std::cout << str << std::endl; }
void memory_coloring_impl::dump_program() { std::cout << *p_program << std::endl; }
void memory_coloring_impl::dump_intervals()
{
if(num_of_lives > 0)
{
std::cout << "---live intervals ---" << std::endl;
for(int i = 0; i < num_of_lives; ++i)
{
live_interval& interval = live_intervals[i];
interval.dump();
}
std::cout << "---conflict table---" << std::endl;
for(int i = 0; i <= max_value_number; ++i)
{
std::cout << " segment:" << i;
std::cout << " =>";
std::set<int>& table = conflict_table[i];
for(auto& iter : table)
{
std::cout << (iter) << ",";
}
}
std::cout << std::endl;
}
}
void memory_coloring_impl::verify()
{
if(num_of_lives > 0)
{
for(int i = 0; i < num_of_lives; ++i)
{
live_interval& interval = live_intervals[i];
live_range& segment = interval.segment;
if(segment.offset == InvalidOffset)
continue;
int vn = segment.vn;
if(conflict_table.find(vn) != conflict_table.end())
{
std::set<int>& vn_set = conflict_table[vn];
for(auto& iter : vn_set)
{
live_range* range = live_ranges[iter];
if(range->offset == InvalidOffset)
continue;
if(!is_disjoin(*range, segment))
assert(false);
}
}
}
}
}
#define GET_INS_ENUM(x) (((x) > 0) ? (((x) >> 1) - 1) : InvalidOffset)
void live_range::dump()
{
std::cout << " segment:" << vn;
std::cout << " [" << GET_INS_ENUM(begin) << ", " << GET_INS_ENUM(end) << "]";
if(offset != InvalidOffset)
{
std::cout << " mem:";
std::cout << " [" << offset << "," << offset + size - 1 << "]";
}
std::cout << std::endl;
}
void live_interval::dump()
{
std::cout << "id:" << id;
segment.dump();
std::cout << " uses:";
for(auto& iter : use_points)
{
std::cout << " " << GET_INS_ENUM(iter) << ",";
}
if(is_literal)
std::cout << " literal";
std::cout << " " << result;
std::cout << std::endl;
}
#endif
} // namespace migraph
src/opt/memory_coloring_impl.hpp deleted 100644 → 0
View file @ 7a2edaf4
#ifndef MIGRAPH_GUARD_RTGLIB_MEMORY_COLORING_IMPL_HPP
#define MIGRAPH_GUARD_RTGLIB_MEMORY_COLORING_IMPL_HPP
#include "common_header.hpp"
namespace migraph {
static const int InvalidOffset = -1;
struct live_range
{
int begin; // begin point in the instruction stream.
int end; // end point in the instruction stream.
long long offset; // offset to base pointer of allocated memory trunk.
int vn; // value number that identifies this live_range.
long long size; // size of required memory in bytes
#ifdef MIGRAPH_DEBUG_OPT
void dump();
#endif
};
struct live_interval
{
live_interval() : segment({-1, -1, InvalidOffset, -1, 0})
{
id = -1;
is_literal = false;
}
void add_use(int use) { use_points.push_front(use); }
int get_begin() const { return segment.begin; }
int get_end() const { return segment.end; }
long long get_offset() const { return segment.offset; }
#ifdef MIGRAPH_DEBUG_OPT
void dump();
#endif
live_range segment;
int id;
std::list<int> use_points;
shape result;
bool is_literal;
};
typedef live_interval* interval_ptr;
struct memory_coloring_impl
{
memory_coloring_impl(program* p) : p_program(p)
{
instr2_live.clear();
live_ranges.clear();
conflict_table.clear();
num_of_lives = 0;
max_value_number = -1;
required_bytes = 0;
}
bool allocate(interval_ptr);
void add_conflicts(std::set<int>& live_set, int val)
{
for(auto& iter : live_set)
{
conflict_table[iter].insert(val);
conflict_table[val].insert(iter);
}
}
void build();
void run();
void rewrite();
private:
static bool is_param(const instruction_ref ins) { return ins->op.name() == "@param"; }
static bool is_output_param(const instruction_ref ins)
{
return is_param(ins) && any_cast<builtin::param>(ins->op).parameter == "output";
}
static bool is_allocate(const instruction_ref ins) { return ins->op.name() == "hip::allocate"; }
static bool is_outline(const instruction_ref ins) { return ins->op.name() == "@outline"; }
static bool is_literal(const instruction_ref ins) { return ins->op.name() == "@literal"; }
static bool is_check_context(const instruction_ref ins)
{
return ins->op.name() == "check_context";
}
static bool is_transpose(const instruction_ref ins) { return ins->op.name() == "transpose"; }
int get_input_tie_ndx(const instruction_ref ins)
{
if(is_transpose(ins))
return 0;
int cnt = -1;
int last_allocate = -1;
for(auto&& arg : ins->arguments)
{
cnt++;
if(is_allocate(arg))
last_allocate = cnt;
}
return last_allocate;
}
#ifdef MIGRAPH_DEBUG_OPT
static bool is_disjoin(live_range& range1, live_range& range2)
{
long long end1 = range1.offset + range1.size - 1;
long long end2 = range2.offset + range2.size - 1;
return ((end1 < range2.offset) || (end2 < range1.offset));
}
void dump(const std::string);
void dump_program();
void dump_intervals();
void verify();
#endif
struct ordering
{
bool operator()(const interval_ptr i1, const interval_ptr i2) const
{
int len1 = i1->get_end() - i1->get_begin();
int len2 = i2->get_end() - i2->get_begin();
if(len1 != len2)
{
return (len1 < len2) ? true : false;
}
else if(i1->result.bytes() != i2->result.bytes())
{
return (i1->result.bytes() < i2->result.bytes()) ? true : false;
}
else
{
return i1->id > i2->id;
}
}
bool operator()(const live_range* i1, const live_range* i2) const
{
return (i1->offset > i2->offset);
}
};
program* p_program;
std::unordered_map<const instruction*, interval_ptr> instr2_live;
// universe of live intervals.
std::vector<live_interval> live_intervals;
// Map live range value number to live range.
std::unordered_map<int, live_range*> live_ranges;
// Map live range value number to a set of conflicting live ranges' value numbers.
std::unordered_map<int, std::set<int>> conflict_table;
// Priority queue for coloring.
std::priority_queue<interval_ptr, std::vector<interval_ptr>, ordering> alloc_queue;
int num_of_lives;
int max_value_number;
long long required_bytes;
};
} // namespace migraph
#endif
src/program.cpp
View file @ 3e6d18a4
......@@ -425,4 +425,5 @@ std::ostream& operator<<(std::ostream& os, const program& p)
print_program(os, p, [](auto&&...) {});
return os;
}
} // namespace migraph
src/targets/cpu/CMakeLists.txt
View file @ 3e6d18a4
......@@ -13,8 +13,3 @@ target_link_libraries(migraph_cpu migraph Threads::Threads)
target_include_directories(migraph_cpu PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)
target_include_directories(migraph_cpu PRIVATE ${BLAZE_INCLUDE})
target_compile_definitions(migraph_cpu PRIVATE -DBLAZE_USE_CPP_THREADS)
#install (TARGETS migraph_cpu
# LIBRARY DESTINATION /opt/rocm/lib)
#install (DIRECTORY include/migraph DESTINATION /opt/rocm/include)
src/targets/gpu/CMakeLists.txt
View file @ 3e6d18a4
......@@ -33,11 +33,3 @@ add_library(migraph_gpu
rocm_clang_tidy_check(migraph_gpu)
target_link_libraries(migraph_gpu migraph MIOpen migraph_device roc::rocblas)
target_include_directories(migraph_gpu PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)
#install (TARGETS migraph_gpu
# LIBRARY DESTINATION /opt/rocm/lib)
#install (DIRECTORY include/migraph DESTINATION /opt/rocm/include)
#install (TARGETS migraph_device
# LIBRARY DESTINATION /opt/rocm/lib)
#install (DIRECTORY include/migraph DESTINATION /opt/rocm/include)
src/targets/gpu/hip.cpp
View file @ 3e6d18a4
......@@ -58,9 +58,7 @@ std::vector<T> read_from_gpu(const void* x, std::size_t sz)
hip_ptr write_to_gpu(const void* x, std::size_t sz, bool host = false)
{
auto result = allocate_gpu(sz, host);
// gpu_sync();
auto status = hipMemcpy(result.get(), x, sz, hipMemcpyHostToDevice);
if(status != hipSuccess)
MIGRAPH_THROW("Copy to gpu failed: " + hip_error(status));
......@@ -92,9 +90,6 @@ argument from_gpu(argument arg)
void gpu_sync() { hipDeviceSynchronize(); }
void copy_to_gpu(char* dst, const char* src, std::size_t size)
{
hipMemcpy(dst, src, size, hipMemcpyHostToDevice);
}
} // namespace gpu
} // namespace migraph
src/targets/gpu/include/migraph/gpu/hip.hpp
View file @ 3e6d18a4
......@@ -14,7 +14,6 @@ migraph::argument to_gpu(migraph::argument arg, bool host = false);
migraph::argument from_gpu(migraph::argument arg);
void gpu_sync();
void copy_to_gpu(char* dst, const char* src, std::size_t size);
struct hip_allocate
{
......@@ -82,21 +81,8 @@ struct hip_write
}
};
struct hip_memcpy
{
std::string name() const { return "hip_memcpy"; }
shape compute_shape(std::vector<shape> inputs) const { return inputs.at(2); }
argument compute(context&, shape output_shape, std::vector<argument> args) const
{
std::size_t* p_data = reinterpret_cast<std::size_t*>(args.at(1).data());
char* dst = args.at(0).data() + p_data[0];
const char* src = args.at(2).data();
std::size_t size = args.at(2).get_shape().bytes();
copy_to_gpu(dst, src, size);
return {output_shape, dst};
}
};
} // namespace gpu
} // namespace migraph
#endif
src/targets/gpu/lowering.cpp
View file @ 3e6d18a4
......@@ -108,9 +108,7 @@ struct miopen_convolution
workspace_shape = shape{shape::int8_type, {workspace_size}};
auto x = to_gpu(generate_argument(inputs[0]->get_shape()));
gpu_sync();
auto w = to_gpu(generate_argument(inputs[1]->get_shape()));
gpu_sync();
auto y = to_gpu(generate_argument(output_shape));
auto workspace = allocate_gpu(workspace_shape);
......
src/targets/gpu/target.cpp
View file @ 3e6d18a4
#include <migraph/gpu/target.hpp>
#include <migraph/gpu/lowering.hpp>
#include <migraph/memory_coloring.hpp>
#include <migraph/gpu/write_literals.hpp>
#include <migraph/gpu/context.hpp>
#include <migraph/gpu/eliminate_workspace.hpp>
......@@ -29,16 +28,13 @@ std::vector<pass> target::get_passes(migraph::context& gctx) const
simplify_reshapes{},
dead_code_elimination{},
lowering{ctx},
memory_coloring{},
fuse_ops{},
dead_code_elimination{},
// eliminate_workspace{},
eliminate_workspace{},
eliminate_contiguous{},
dead_code_elimination{},
write_literals{&ctx},
// eliminate_allocation{},
eliminate_allocation{},
check_context<context>{},
dead_code_elimination{}
};
......
src/targets/gpu/write_literals.cpp
View file @ 3e6d18a4
......@@ -28,7 +28,6 @@ void write_literals::apply(program& p) const
assert(ctx != nullptr);
for(auto ins : iterator_for(p))
{
#if 0
if(ins->op.name() == "@literal")
{
argument a = to_gpu(ins->lit.get_argument());
......@@ -36,13 +35,9 @@ void write_literals::apply(program& p) const
ctx->literals.push_back(a);
p.replace_instruction(ins, hip_load_literal{a.get_shape(), n});
}
#else
if(ins->op.name() == "write_literal")
{
p.replace_instruction(ins, hip_memcpy{}, ins->arguments);
}
#endif
}
}
} // namespace gpu
} // namespace migraph
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment