AI/rocm-systems
2
0
포크 0
코드 이슈 풀 리퀘스트 액션 패키지 프로젝트 릴리즈 위키 활동

Merge pull request #1140 from scchan/program_state_stage_2-rebase-20190524

소스 검색 
migrate more program_state logic from header into shared library (phase II)

[ROCm/clr commit: 1d5d923d36]
이 커밋은 다음에 포함됨:
Maneesh Gupta
2019-06-05 16:09:01 +05:30
커밋한 사람 GitHub
부모 932d76e45b cf2e2c5077
커밋 b4fb2b0ab4
8개의 변경된 파일과 309개의 추가작업 그리고 269개의 파일을 삭제
패치 파일 다운로드 Diff 파일 다운로드 모든 파일 펼침 모든 파일 접기
projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
+6 -7
파일 보기
@@ -22,7 +22,6 @@ THE SOFTWARE.
#pragma once
#include "code_object_bundle.hpp"
#include "concepts.hpp"
#include "helpers.hpp"
#include "program_state.hpp"
@@ -57,10 +56,10 @@ template <
std::size_t n,
typename... Ts,
typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
inline std::vector<std::uint8_t> make_kernarg(
inline hip_impl::kernarg make_kernarg(
const std::tuple<Ts...>&,
const kernargs_size_align&,
std::vector<std::uint8_t> kernarg) {
hip_impl::kernarg kernarg) {
return kernarg;
}
@@ -68,10 +67,10 @@ template <
std::size_t n,
typename... Ts,
typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
inline std::vector<std::uint8_t> make_kernarg(
inline hip_impl::kernarg make_kernarg(
const std::tuple<Ts...>& formals,
const kernargs_size_align& size_align,
std::vector<std::uint8_t> kernarg) {
hip_impl::kernarg kernarg) {
using T = typename std::tuple_element<n, std::tuple<Ts...>>::type;
static_assert(
@@ -96,7 +95,7 @@ inline std::vector<std::uint8_t> make_kernarg(
}
template <typename... Formals, typename... Actuals>
inline std::vector<std::uint8_t> make_kernarg(
inline hip_impl::kernarg make_kernarg(
void (*kernel)(Formals...), std::tuple<Actuals...> actuals) {
static_assert(sizeof...(Formals) == sizeof...(Actuals),
"The count of formal arguments must match the count of actuals.");
@@ -104,7 +103,7 @@ inline std::vector<std::uint8_t> make_kernarg(
if (sizeof...(Formals) == 0) return {};
std::tuple<Formals...> to_formals{std::move(actuals)};
std::vector<std::uint8_t> kernarg;
hip_impl::kernarg kernarg;
kernarg.reserve(sizeof(to_formals));
auto& ps = hip_impl::get_program_state();
projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp
+10 -1
파일 보기
@@ -84,12 +84,21 @@ struct is_callable_impl<F(Ts...), 4u, void_t_<decltype(std::declval<F>()(std::de
// Not callable.
template <FunctionalProcedure F>
struct is_callable_impl<F, 5u> : std::false_type {};
#else
#elif (__cplusplus < 201703L)
template <typename, typename = void>
struct is_callable_impl : std::false_type {};
template <FunctionalProcedure F, typename... Ts>
struct is_callable_impl<F(Ts...), void_t_<std::result_of_t<F(Ts...)> > > : std::true_type {};
#else
// C++17
template <typename, typename = void>
struct is_callable_impl : std::false_type {};
template <FunctionalProcedure F, typename... Ts>
struct is_callable_impl<F(Ts...), void_t_<std::invoke_result<F(Ts...)> > > : std::true_type {};
#endif
template <typename Call>
struct is_callable : is_callable_impl<Call> {};
projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
+23 -162
파일 보기
@@ -2599,180 +2599,41 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con
hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func);
struct Agent_global {
Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {}
Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt)
: name(nullptr), address(address), byte_cnt(byte_cnt) {
if (name)
this->name = strdup(name);
}
Agent_global& operator=(Agent_global&& t) {
if (this == &t) return *this;
if (name) free(name);
name = t.name;
address = t.address;
byte_cnt = t.byte_cnt;
t.name = nullptr;
t.address = nullptr;
t.byte_cnt = 0;
return *this;
}
Agent_global(Agent_global&& t)
: name(nullptr), address(nullptr), byte_cnt(0) {
*this = std::move(t);
}
// not needed, delete them to prevent bugs
Agent_global(const Agent_global&) = delete;
Agent_global& operator=(Agent_global& t) = delete;
~Agent_global() { if (name) free(name); }
char* name;
hipDeviceptr_t address;
uint32_t byte_cnt;
};
#if !__HIP_VDI__
#if defined(__cplusplus)
} // extern "C"
#endif
namespace hip_impl {
hsa_executable_t executable_for(hipModule_t);
const char* hash_for(hipModule_t);
class agent_globals_impl;
class agent_globals {
public:
agent_globals();
~agent_globals();
agent_globals(const agent_globals&) = delete;
template<typename ForwardIterator>
std::pair<hipDeviceptr_t, std::size_t> read_global_description(
ForwardIterator f, ForwardIterator l, const char* name) {
const auto it = std::find_if(f, l, [=](const Agent_global& x) {
return strcmp(x.name, name) == 0;
});
hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
hipModule_t hmod, const char* name);
hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
const char* name);
private:
agent_globals_impl* impl;
};
return it == l ?
std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt);
}
std::vector<Agent_global> read_agent_globals(hsa_agent_t agent,
hsa_executable_t executable);
hsa_agent_t this_agent();
class agent_globals_impl {
private:
std::pair<
std::mutex,
std::unordered_map<
std::string, std::vector<Agent_global>>> globals_from_module;
std::unordered_map<
hsa_agent_t,
std::pair<
std::once_flag,
std::vector<Agent_global>>> globals_from_process;
public:
hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
hipModule_t hmod, const char* name) {
// the key of the map would the hash of code object associated with the
// hipModule_t instance
std::string key(hash_for(hmod));
if (globals_from_module.second.count(key) == 0) {
std::lock_guard<std::mutex> lck{globals_from_module.first};
if (globals_from_module.second.count(key) == 0) {
globals_from_module.second.emplace(
key, read_agent_globals(this_agent(), executable_for(hmod)));
}
}
const auto it0 = globals_from_module.second.find(key);
if (it0 == globals_from_module.second.cend()) {
hip_throw(
std::runtime_error{"agent_globals data structure corrupted."});
}
std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(),
it0->second.cend(), name);
return *dptr ? hipSuccess : hipErrorNotFound;
inline
__attribute__((visibility("hidden")))
agent_globals& get_agent_globals() {
static agent_globals ag;
return ag;
}
extern "C"
inline
__attribute__((visibility("hidden")))
hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
const char* name) {
auto agent = this_agent();
std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) {
std::vector<Agent_global> tmp0;
for (auto&& executable : hip_impl::get_program_state().executables(aa)) {
auto tmp1 = read_agent_globals(aa, executable);
tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
make_move_iterator(tmp1.end()));
}
globals_from_process[aa].second = move(move(tmp0));
}, agent);
const auto it = globals_from_process.find(agent);
if (it == globals_from_process.cend()) return hipErrorNotInitialized;
std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(),
it->second.second.cend(), name);
return *dptr ? hipSuccess : hipErrorNotFound;
const char* name) {
return get_agent_globals().read_agent_global_from_process(dptr, bytes, name);
}
};
class agent_globals {
public:
agent_globals() : impl(new agent_globals_impl()) {
if (!impl)
hip_throw(
std::runtime_error{"Error when constructing agent global data structures."});
}
~agent_globals() { delete impl; }
hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
hipModule_t hmod, const char* name) {
return impl->read_agent_global_from_module(dptr, bytes, hmod, name);
}
hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
const char* name) {
return impl->read_agent_global_from_process(dptr, bytes, name);
}
private:
agent_globals_impl* impl;
};
inline
__attribute__((visibility("hidden")))
agent_globals& get_agent_globals() {
static agent_globals ag;
return ag;
}
extern "C"
inline
__attribute__((visibility("hidden")))
hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
const char* name) {
return get_agent_globals().read_agent_global_from_process(dptr, bytes, name);
}
} // Namespace hip_impl.
#if defined(__cplusplus)
projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp
+15 -34
파일 보기
@@ -30,41 +30,25 @@ THE SOFTWARE.
#include <cstddef>
#include <cstdint>
#include <cstdlib>
#include <stdexcept>
#include <string>
#include <unordered_map>
#include <vector>
struct ihipModuleSymbol_t;
using hipFunction_t = ihipModuleSymbol_t*;
namespace std {
template<>
struct hash<hsa_agent_t> {
size_t operator()(hsa_agent_t x) const {
return hash<decltype(x.handle)>{}(x.handle);
}
};
template<>
struct hash<hsa_isa_t> {
size_t operator()(hsa_isa_t x) const {
return hash<decltype(x.handle)>{}(x.handle);
}
};
} // namespace std
inline constexpr bool operator==(hsa_agent_t x, hsa_agent_t y) {
return x.handle == y.handle;
}
inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) {
return x.handle == y.handle;
}
namespace hip_impl {
[[noreturn]]
void hip_throw(const std::exception&);
struct kernarg_impl;
class kernarg {
public:
kernarg();
kernarg(kernarg&&);
~kernarg();
std::uint8_t* data();
std::size_t size();
void reserve(std::size_t);
void resize(std::size_t);
private:
kernarg_impl* impl;
};
class kernargs_size_align;
class program_state_impl;
@@ -72,6 +56,7 @@ class program_state {
public:
program_state();
~program_state();
program_state(const program_state&) = delete;
hipFunction_t kernel_descriptor(std::uintptr_t,
hsa_agent_t);
@@ -83,12 +68,8 @@ public:
void* global_addr_by_name(const char* name);
// to fix later
const std::vector<hsa_executable_t>& executables(hsa_agent_t agent);
program_state(const program_state&) = delete;
private:
friend class agent_globals_impl;
program_state_impl* impl;
};
projects/clr/hipamd/src/hip_hcc.cpp
-9
파일 보기
@@ -2492,13 +2492,4 @@ namespace hip_impl {
#endif
}
std::mutex executables_cache_mutex;
std::vector<hsa_executable_t>& executables_cache(
std::string elf, hsa_isa_t isa, hsa_agent_t agent) {
static std::unordered_map<std::string,
std::unordered_map<hsa_isa_t,
std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>>> cache;
return cache[elf][isa][agent];
}
} // Namespace hip_impl.
projects/clr/hipamd/src/hip_module.cpp
+152 -13
파일 보기
@@ -313,16 +313,6 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0));
}
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
hipModule_t hmod, const char* name) {
HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
if (!dptr || !bytes || !hmod) return hipErrorInvalidValue;
if (!name) return hipErrorNotInitialized;
return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name);
}
namespace hip_impl {
hsa_executable_t executable_for(hipModule_t hmod) {
return hmod->executable;
@@ -347,10 +337,159 @@ namespace hip_impl {
return currentDevice->_hsaAgent;
}
struct Agent_global {
Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {}
Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt)
: name(nullptr), address(address), byte_cnt(byte_cnt) {
if (name)
this->name = strdup(name);
}
Agent_global& operator=(Agent_global&& t) {
if (this == &t) return *this;
if (name) free(name);
name = t.name;
address = t.address;
byte_cnt = t.byte_cnt;
t.name = nullptr;
t.address = nullptr;
t.byte_cnt = 0;
return *this;
}
Agent_global(Agent_global&& t)
: name(nullptr), address(nullptr), byte_cnt(0) {
*this = std::move(t);
}
// not needed, delete them to prevent bugs
Agent_global(const Agent_global&) = delete;
Agent_global& operator=(Agent_global& t) = delete;
~Agent_global() { if (name) free(name); }
char* name;
hipDeviceptr_t address;
uint32_t byte_cnt;
};
template<typename ForwardIterator>
std::pair<hipDeviceptr_t, std::size_t> read_global_description(
ForwardIterator f, ForwardIterator l, const char* name) {
const auto it = std::find_if(f, l, [=](const Agent_global& x) {
return strcmp(x.name, name) == 0;
});
return it == l ?
std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt);
}
std::vector<Agent_global> read_agent_globals(hsa_agent_t agent,
hsa_executable_t executable);
class agent_globals_impl {
private:
std::pair<
std::mutex,
std::unordered_map<
std::string, std::vector<Agent_global>>> globals_from_module;
std::unordered_map<
hsa_agent_t,
std::pair<
std::once_flag,
std::vector<Agent_global>>> globals_from_process;
public:
hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
hipModule_t hmod, const char* name) {
// the key of the map would the hash of code object associated with the
// hipModule_t instance
std::string key(hash_for(hmod));
if (globals_from_module.second.count(key) == 0) {
std::lock_guard<std::mutex> lck{globals_from_module.first};
if (globals_from_module.second.count(key) == 0) {
globals_from_module.second.emplace(
key, read_agent_globals(this_agent(), executable_for(hmod)));
}
}
const auto it0 = globals_from_module.second.find(key);
if (it0 == globals_from_module.second.cend()) {
hip_throw(
std::runtime_error{"agent_globals data structure corrupted."});
}
std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(),
it0->second.cend(), name);
return *dptr ? hipSuccess : hipErrorNotFound;
}
hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
const char* name) {
auto agent = this_agent();
std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) {
std::vector<Agent_global> tmp0;
for (auto&& executable : hip_impl::get_program_state().impl->get_executables(aa)) {
auto tmp1 = read_agent_globals(aa, executable);
tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
make_move_iterator(tmp1.end()));
}
globals_from_process[aa].second = move(move(tmp0));
}, agent);
const auto it = globals_from_process.find(agent);
if (it == globals_from_process.cend()) return hipErrorNotInitialized;
std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(),
it->second.second.cend(), name);
return *dptr ? hipSuccess : hipErrorNotFound;
}
};
agent_globals::agent_globals() : impl(new agent_globals_impl()) {
if (!impl)
hip_throw(
std::runtime_error{"Error when constructing agent global data structures."});
}
agent_globals::~agent_globals() { delete impl; }
hipError_t agent_globals::read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
hipModule_t hmod, const char* name) {
return impl->read_agent_global_from_module(dptr, bytes, hmod, name);
}
hipError_t agent_globals::read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
const char* name) {
return impl->read_agent_global_from_process(dptr, bytes, name);
}
} // Namespace hip_impl.
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
hipModule_t hmod, const char* name) {
HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
if (!dptr || !bytes || !hmod) return hipErrorInvalidValue;
if (!name) return hipErrorNotInitialized;
return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name);
}
namespace {
inline void track(const Agent_global& x, hsa_agent_t agent) {
inline void track(const hip_impl::Agent_global& x, hsa_agent_t agent) {
tprintf(DB_MEM, " add variable '%s' with ptr=%p size=%u to tracker\n", x.name,
x.address, x.byte_cnt);
@@ -371,7 +510,7 @@ inline void track(const Agent_global& x, hsa_agent_t agent) {
}
template <typename Container = vector<Agent_global>>
template <typename Container = vector<hip_impl::Agent_global>>
inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t agent,
hsa_executable_symbol_t x, void* out) {
using namespace hip_impl;
@@ -382,7 +521,7 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag
hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &t);
if (t == HSA_SYMBOL_KIND_VARIABLE) {
Agent_global tmp(name(x).c_str(), address(x), size(x));
hip_impl::Agent_global tmp(name(x).c_str(), address(x), size(x));
static_cast<Container*>(out)->push_back(std::move(tmp));
track(static_cast<Container*>(out)->back(),agent);
projects/clr/hipamd/src/program_state.cpp
+74 -43
파일 보기
@@ -1,4 +1,6 @@
#include "../include/hip/hcc_detail/program_state.hpp"
// contains implementation of program_state_impl
#include "program_state.inl"
#include <hsa/hsa.h>
@@ -7,57 +9,86 @@
#include <unordered_map>
#include <vector>
// contains implementation of program_state_impl
#include "program_state.inl"
namespace hip_impl {
kernarg::kernarg() : impl(new kernarg_impl) {
}
std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{
return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].first;
}
kernarg::kernarg(kernarg&& k) : impl(k.impl) {
k.impl = nullptr;
}
std::size_t kernargs_size_align::alignment(std::size_t n) const{
return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].second;
}
kernarg::~kernarg() {
if (impl)
delete(impl);
}
program_state::program_state() :
impl(new program_state_impl) {
if (!impl) hip_throw(std::runtime_error {
"Unknown error when constructing program state."});
}
std::uint8_t* kernarg::data() {
return impl->v.data();
}
program_state::~program_state() {
delete(impl);
}
std::size_t kernarg::size() {
return impl->v.size();
}
void* program_state::global_addr_by_name(const char* name) {
const auto it = impl->get_globals().find(name);
if (it == impl->get_globals().end())
return nullptr;
else
return it->second;
}
void kernarg::reserve(std::size_t c) {
impl->v.reserve(c);
}
hsa_executable_t program_state::load_executable(const char* data,
const size_t data_size,
hsa_executable_t executable,
hsa_agent_t agent) {
return impl->load_executable(data, data_size, executable, agent);
}
void kernarg::resize(std::size_t c) {
impl->v.resize(c);
}
const std::vector<hsa_executable_t>& program_state::executables(hsa_agent_t agent) {
return impl->get_executables(agent);
}
std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{
return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].first;
}
hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address,
hsa_agent_t agent) {
auto& kd = impl->kernel_descriptor(function_address, agent);
return kd;
}
std::size_t kernargs_size_align::alignment(std::size_t n) const{
return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].second;
}
kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t kernel) {
kernargs_size_align t;
t.handle = reinterpret_cast<const void*>(&impl->kernargs_size_align(kernel));
return t;
}
program_state::program_state() : impl(new program_state_impl) {
if (!impl) hip_throw(std::runtime_error {
"Unknown error when constructing program state."});
}
program_state::~program_state() {
delete(impl);
}
void* program_state::global_addr_by_name(const char* name) {
const auto it = impl->get_globals().find(name);
if (it == impl->get_globals().end())
return nullptr;
else
return it->second;
}
hsa_executable_t program_state::load_executable(const char* data,
const size_t data_size,
hsa_executable_t executable,
hsa_agent_t agent) {
return impl->load_executable(data, data_size, executable, agent);
}
hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address,
hsa_agent_t agent) {
auto& kd = impl->kernel_descriptor(function_address, agent);
return kd;
}
kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t kernel) {
kernargs_size_align t;
t.handle = reinterpret_cast<const void*>(&impl->kernargs_size_align(kernel));
return t;
}
std::mutex executables_cache_mutex;
std::vector<hsa_executable_t>& executables_cache(
std::string elf, hsa_isa_t isa, hsa_agent_t agent) {
static std::unordered_map<std::string,
std::unordered_map<hsa_isa_t,
std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>>> cache;
return cache[elf][isa][agent];
}
};
projects/clr/hipamd/src/program_state.inl
+29
파일 보기
@@ -27,12 +27,36 @@
#include <cstdio>
#include <memory>
#include <mutex>
#include <stdexcept>
#include <string>
#include <sstream>
#include <unordered_map>
#include <utility>
#include <vector>
namespace std {
template<>
struct hash<hsa_agent_t> {
size_t operator()(hsa_agent_t x) const {
return hash<decltype(x.handle)>{}(x.handle);
}
};
template<>
struct hash<hsa_isa_t> {
size_t operator()(hsa_isa_t x) const {
return hash<decltype(x.handle)>{}(x.handle);
}
};
} // namespace std
inline constexpr bool operator==(hsa_agent_t x, hsa_agent_t y) {
return x.handle == y.handle;
}
inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) {
return x.handle == y.handle;
}
namespace hip_impl {
[[noreturn]]
@@ -801,4 +825,9 @@ public:
}
}; // class program_state_impl
struct kernarg_impl {
std::vector<std::uint8_t> v;
};
};