AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

Merge 'master' into 'amd-master'

Bladeren bron 
Change-Id: I599b0d1176a11c635c8b16c395d15677bd431fb5


[ROCm/hip commit: 1964bb1115]
This commit is contained in:
Jenkins
2018-04-17 04:09:48 -05:00
bovenliggende 40c2180e83 e119cc4f85
commit aad9b2ffee
2 gewijzigde bestanden met toevoegingen van 78 en 44 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
projects/hip/src/code_object_bundle.cpp
+45 -13
Bestand weergeven
@@ -10,23 +10,55 @@
using namespace std;
hsa_isa_t hip_impl::triple_to_hsa_isa(const std::string& triple) {
static constexpr const char prefix[] = "hcc-amdgcn--amdhsa-gfx";
static constexpr size_t prefix_sz = sizeof(prefix) - 1;
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"};
hsa_isa_t r = {};
if (triple.find(old_prefix) == 0) {
return new_prefix + triple.substr(sizeof(old_prefix) - 1);
}
auto idx = triple.find(prefix);
return (triple.find(new_prefix) == 0) ? triple : "";
}
if (idx != string::npos) {
idx += prefix_sz;
string tmp = "AMD:AMDGPU";
while (idx != triple.size()) {
tmp.push_back(':');
tmp.push_back(triple[idx++]);
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) {
auto 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);
}
}
hsa_isa_from_name(tmp.c_str(), &r);
return triple;
}
hsa_isa_t hip_impl::triple_to_hsa_isa(const std::string& triple) {
const auto 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;
@@ -59,4 +91,4 @@ hip_impl::Bundled_code_header::Bundled_code_header(
}
read(static_cast<const char*>(p), static_cast<const char*>(p) + sz, *this);
}
}
projects/hip/src/hip_module.cpp
+33 -31
Bestand weergeven
@@ -429,30 +429,27 @@ string code_object_blob_for_agent(const void* maybe_bundled_code, hsa_agent_t ag
} // namespace
hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const char* name) {
HIP_INIT_API(func, hmod, name);
if (!func || !name) return ihipLogStatus(hipErrorInvalidValue);
if (!func || !name) return hipErrorInvalidValue;
auto ctx = ihipGetTlsDefaultCtx();
if (!ctx) return ihipLogStatus(hipErrorInvalidContext);
hipError_t ret = hipSuccess;
if (!ctx) return hipErrorInvalidContext;
*func = new ihipModuleSymbol_t;
if (!*func) return ihipLogStatus(hipErrorInvalidValue);
if (!*func) return hipErrorInvalidValue;
auto kernel = find_kernel_by_name(hmod->executable, name);
if (kernel.handle == 0u) return ihipLogStatus(hipErrorNotFound);
if (kernel.handle == 0u) return hipErrorNotFound;
(*func)->_object = kernel_object(kernel);
(*func)->_groupSegmentSize = group_size(kernel);
(*func)->_privateSegmentSize = private_size(kernel);
(*func)->_name = name;
return ihipLogStatus(hipSuccess);
return hipSuccess;
}
hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const char* name) {
@@ -474,6 +471,31 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h
return ihipLogStatus(r);
}
hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
if (!module) return hipErrorInvalidValue;
*module = new ihipModule_t;
auto ctx = ihipGetTlsDefaultCtx();
if (!ctx) return hipErrorInvalidContext;
hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr,
&(*module)->executable);
auto tmp = code_object_blob_for_agent(image, this_agent());
(*module)->executable = hip_impl::load_executable(
tmp.empty() ? read_elf_file_as_string(image) : tmp, (*module)->executable, this_agent());
return (*module)->executable.handle ? hipSuccess : hipErrorUnknown;
}
hipError_t hipModuleLoadData(hipModule_t* module, const void* image) {
HIP_INIT_API(module, image);
return ihipLogStatus(ihipModuleLoadData(module,image));
}
hipError_t hipModuleLoad(hipModule_t* module, const char* fname) {
HIP_INIT_API(module, fname);
@@ -485,33 +507,13 @@ hipError_t hipModuleLoad(hipModule_t* module, const char* fname) {
vector<char> tmp{istreambuf_iterator<char>{file}, istreambuf_iterator<char>{}};
return hipModuleLoadData(module, tmp.data());
}
hipError_t hipModuleLoadData(hipModule_t* module, const void* image) {
HIP_INIT_API(module, image);
if (!module) return ihipLogStatus(hipErrorInvalidValue);
*module = new ihipModule_t;
auto ctx = ihipGetTlsDefaultCtx();
if (!ctx) return ihipLogStatus(hipErrorInvalidContext);
hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr,
&(*module)->executable);
auto tmp = code_object_blob_for_agent(image, this_agent());
(*module)->executable = hip_impl::load_executable(
tmp.empty() ? read_elf_file_as_string(image) : tmp, (*module)->executable, this_agent());
return ihipLogStatus((*module)->executable.handle ? hipSuccess : hipErrorUnknown);
return ihipLogStatus(ihipModuleLoadData(module, tmp.data()));
}
hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned int numOptions,
hipJitOption* options, void** optionValues) {
return hipModuleLoadData(module, image);
HIP_INIT_API(module, image, numOptions, options, optionValues);
return ihipLogStatus(ihipModuleLoadData(module, image));
}
hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name) {