AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: I7fd7b61a4a78eeb26600bfe18593c0fd67a008b8


[ROCm/clr commit: 4f32124b90]
This commit is contained in:
Jenkins
2019-03-16 16:37:02 -04:00
parent 0ea08121b1 8b550ca263
commit d65aae2a3e
9 changed files with 751 additions and 887 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp
+42 -1
View File
@@ -33,7 +33,48 @@ THE SOFTWARE.
#include <vector>
namespace hip_impl {
hsa_isa_t triple_to_hsa_isa(const std::string& triple);
inline
std::string transmogrify_triple(const std::string& triple)
{
static constexpr const char old_prefix[]{"hcc-amdgcn--amdhsa-gfx"};
static constexpr const char new_prefix[]{"hcc-amdgcn-amd-amdhsa--gfx"};
if (triple.find(old_prefix) == 0) {
return new_prefix + triple.substr(sizeof(old_prefix) - 1);
}
return (triple.find(new_prefix) == 0) ? triple : "";
}
inline
std::string isa_name(std::string triple)
{
static constexpr const char offload_prefix[]{"hcc-"};
triple = transmogrify_triple(triple);
if (triple.empty()) return {};
triple.erase(0, sizeof(offload_prefix) - 1);
return triple;
}
inline
hsa_isa_t triple_to_hsa_isa(const std::string& triple) {
const std::string isa{isa_name(std::move(triple))};
if (isa.empty()) return hsa_isa_t({});
hsa_isa_t r{};
if(HSA_STATUS_SUCCESS != hsa_isa_from_name(isa.c_str(), &r)) {
r.handle = 0;
}
return r;
}
struct Bundled_code {
union Header {
projects/clr/hipamd/include/hip/hcc_detail/hip_prof_str.h
+654 -814
View File
File diff suppressed because it is too large Load Diff
projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
+39 -5
View File
@@ -2506,7 +2506,41 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con
hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func);
struct Agent_global {
std::string name;
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;
};
@@ -2518,13 +2552,13 @@ struct Agent_global {
namespace hip_impl {
hsa_executable_t executable_for(hipModule_t);
const std::string& hash_for(hipModule_t);
const char* hash_for(hipModule_t);
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 x.name == name;
return strcmp(x.name, name) == 0;
});
return it == l ?
@@ -2543,7 +2577,7 @@ hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
// hipModule_t instance
static std::unordered_map<
std::string, std::vector<Agent_global>> agent_globals;
auto key = hash_for(hmod);
std::string key(hash_for(hmod));
if (agent_globals.count(key) == 0) {
static std::mutex mtx;
@@ -2890,7 +2924,7 @@ hipError_t hipRemoveApiCallback(uint32_t id);
hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg);
hipError_t hipRemoveActivityCallback(uint32_t id);
static inline const char* hipApiName(const uint32_t& id) { return hip_api_name(id); }
const char* hipKernelNameRef(hipFunction_t f);
const char* hipKernelNameRef(const hipFunction_t f);
#ifdef __cplusplus
} /* extern "C" */
#endif
projects/clr/hipamd/samples/0_Intro/module_api_global/runKernel.cpp
+4 -2
View File
@@ -31,7 +31,6 @@ THE SOFTWARE.
#define SIZE LEN * sizeof(float)
#define fileName "vcpy_kernel.code"
float myDeviceGlobal;
float myDeviceGlobalArray[16];
#define HIP_CHECK(cmd) \
{ \
@@ -69,7 +68,10 @@ int main() {
HIP_CHECK(hipModuleLoad(&Module, fileName));
float myDeviceGlobal_h = 42.0;
myDeviceGlobal = 42.0;
float* deviceGlobal;
size_t deviceGlobalSize;
HIP_CHECK(hipModuleGetGlobal((void**)&deviceGlobal, &deviceGlobalSize, Module, "myDeviceGlobal"));
*deviceGlobal = 42.0;
#define ARRAY_SIZE 16
projects/clr/hipamd/samples/0_Intro/module_api_global/vcpy_kernel.cpp
+1 -1
View File
@@ -24,7 +24,7 @@ THE SOFTWARE.
#define ARRAY_SIZE (16)
extern float myDeviceGlobal;
__device__ float myDeviceGlobal;
extern float myDeviceGlobalArray[16];
;
projects/clr/hipamd/src/code_object_bundle.cpp
-54
View File
@@ -10,60 +10,6 @@
using namespace std;
inline
std::string transmogrify_triple(const std::string& triple)
{
static constexpr const char old_prefix[]{"hcc-amdgcn--amdhsa-gfx"};
static constexpr const char new_prefix[]{"hcc-amdgcn-amd-amdhsa--gfx"};
if (triple.find(old_prefix) == 0) {
return new_prefix + triple.substr(sizeof(old_prefix) - 1);
}
return (triple.find(new_prefix) == 0) ? triple : "";
}
inline
std::string isa_name(std::string triple)
{
static constexpr const char offload_prefix[]{"hcc-"};
triple = transmogrify_triple(triple);
if (triple.empty()) return {};
triple.erase(0, sizeof(offload_prefix) - 1);
static hsa_isa_t tmp{};
static const bool is_old_rocr{
hsa_isa_from_name(triple.c_str(), &tmp) != HSA_STATUS_SUCCESS};
if (is_old_rocr) {
std::string tmp{triple.substr(triple.rfind('x') + 1)};
triple.replace(0, std::string::npos, "AMD:AMDGPU");
for (auto&& x : tmp) {
triple.push_back(':');
triple.push_back(x);
}
}
return triple;
}
hsa_isa_t hip_impl::triple_to_hsa_isa(const std::string& triple) {
const std::string isa{isa_name(std::move(triple))};
if (isa.empty()) return hsa_isa_t({});
hsa_isa_t r{};
if(HSA_STATUS_SUCCESS != hsa_isa_from_name(isa.c_str(), &r)) {
r.handle = 0;
}
return r;
}
// DATA - STATICS
constexpr const char hip_impl::Bundled_code_header::magic_string_[];
projects/clr/hipamd/src/hip_intercept.cpp
+1 -1
View File
@@ -27,7 +27,7 @@ THE SOFTWARE.
api_callbacks_table_t callbacks_table;
extern std::string& FunctionSymbol(hipFunction_t f);
extern std::string& FunctionSymbol(const hipFunction_t f);
const char* hipKernelNameRef(const hipFunction_t f) { return FunctionSymbol(f).c_str(); }
hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg) {
projects/clr/hipamd/src/hip_module.cpp
+7 -6
View File
@@ -99,7 +99,7 @@ string ToString(hipFunction_t v) {
return ss.str();
};
std::string& FunctionSymbol(hipFunction_t f) { return f->_name; };
const std::string& FunctionSymbol(const hipFunction_t f) { return f->_name; };
#define CHECK_HSA(hsaStatus, hipStatus) \
if (hsaStatus != HSA_STATUS_SUCCESS) { \
@@ -262,7 +262,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) {
HIP_INIT_API(hipHccModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX,
HIP_INIT_API(hipExtModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX,
localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra);
return ihipLogStatus(ihipModuleLaunchKernel(
f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
@@ -287,8 +287,8 @@ namespace hip_impl {
return hmod->executable;
}
const std::string& hash_for(hipModule_t hmod) {
return hmod->hash;
const char* hash_for(hipModule_t hmod) {
return hmod->hash.c_str();
}
hsa_agent_t this_agent() {
@@ -310,7 +310,7 @@ namespace hip_impl {
namespace {
inline void track(const Agent_global& x, hsa_agent_t agent) {
tprintf(DB_MEM, " add variable '%s' with ptr=%p size=%u to tracker\n", x.name.c_str(),
tprintf(DB_MEM, " add variable '%s' with ptr=%p size=%u to tracker\n", x.name,
x.address, x.byte_cnt);
int deviceIndex =0;
@@ -341,7 +341,8 @@ 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) {
static_cast<Container*>(out)->push_back(Agent_global{name(x), address(x), size(x)});
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/hip_peer.cpp
+3 -3
View File
@@ -73,7 +73,7 @@ hipError_t ihipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_
*/
//---
hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) {
HIP_INIT_API(hipDeviceCanAccessPeer2, canAccessPeer, thisCtx, peerCtx);
HIP_INIT_API(NONE, canAccessPeer, thisCtx, peerCtx);
return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, thisCtx, peerCtx));
}
@@ -150,7 +150,7 @@ hipError_t ihipEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) {
//---
hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t srcCtx,
size_t sizeBytes) {
HIP_INIT_API(hipMemcpyPeer2, dst, dstCtx, src, srcCtx, sizeBytes);
HIP_INIT_API(NONE, dst, dstCtx, src, srcCtx, sizeBytes);
// TODO - move to ihip memory copy implementaion.
// HCC has a unified memory architecture so device specifiers are not required.
@@ -161,7 +161,7 @@ hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t s
//---
hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hipCtx_t srcDevice,
size_t sizeBytes, hipStream_t stream) {
HIP_INIT_API(hipMemcpyPeerAsync2, dst, dstDevice, src, srcDevice, sizeBytes, stream);
HIP_INIT_API(NONE, dst, dstDevice, src, srcDevice, sizeBytes, stream);
// TODO - move to ihip memory copy implementaion.
// HCC has a unified memory architecture so device specifiers are not required.