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

Update code object/isa/loader to hsa v1.1

Browse Source 
- Includes Sean's latest changes
- Cleanups/improvements
- Fixes for few bugs that crept over from previous releases

Change-Id: I839dc4895bf13ebd0afc8843424387a9fef667b0


[ROCm/ROCR-Runtime commit: c2c993e0d8]
This commit is contained in:
Konstantin Zhuravlyov
2016-08-22 01:40:26 -04:00
parent a9dff11965
commit e10cd184ef
32 changed files with 5831 additions and 2010 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp
+9 -9
View File
@@ -603,7 +603,7 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand(void* dst, const void* src,
}
// Wait for the packet to finish.
if (HSA::hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_LT,
if (HSA::hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_LT,
1, uint64_t(-1),
HSA_WAIT_STATE_ACTIVE) != 0) {
// Signal wait returned unexpected value.
@@ -627,8 +627,8 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand(
const uint16_t kBarrierPacketHeader =
(HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE) |
(1 << HSA_PACKET_HEADER_BARRIER) |
(HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
(HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) |
(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
hsa_barrier_and_packet_t barrier_packet = {0};
barrier_packet.header = HSA_PACKET_TYPE_INVALID;
@@ -766,7 +766,7 @@ hsa_status_t BlitKernel::SubmitLinearFillCommand(void* ptr, uint32_t value,
ReleaseWriteIndex(write_index, 1);
// Wait for the packet to finish.
if (HSA::hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_LT,
if (HSA::hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_LT,
1, uint64_t(-1),
HSA_WAIT_STATE_ACTIVE) != 0) {
// Signal wait returned unexpected value.
@@ -813,8 +813,8 @@ hsa_status_t BlitKernel::FenceRelease(uint64_t write_index,
const uint16_t kBarrierPacketHeader =
(HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE) |
(1 << HSA_PACKET_HEADER_BARRIER) |
(HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
(fence << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
(HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) |
(fence << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
hsa_barrier_and_packet_t packet = {0};
packet.header = kInvalidPacketHeader;
@@ -842,7 +842,7 @@ hsa_status_t BlitKernel::FenceRelease(uint64_t write_index,
ReleaseWriteIndex(write_index, num_copy_packet + 1);
// Wait for the packet to finish.
if (HSA::hsa_signal_wait_acquire(packet.completion_signal,
if (HSA::hsa_signal_wait_scacquire(packet.completion_signal,
HSA_SIGNAL_CONDITION_LT, 1, uint64_t(-1),
HSA_WAIT_STATE_ACTIVE) != 0) {
// Signal wait returned unexpected value.
@@ -862,8 +862,8 @@ void BlitKernel::PopulateQueue(uint64_t index, uint64_t code_handle, void* args,
static const uint16_t kDispatchPacketHeader =
(HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
(((completion_signal.handle != 0) ? 1 : 0) << HSA_PACKET_HEADER_BARRIER) |
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) |
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
packet.header = kInvalidPacketHeader;
packet.kernel_object = code_handle;
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_cpu_agent.cpp
+21 -1
View File
@@ -125,6 +125,15 @@ void CpuAgent::InitCacheList() {
}
}
}
//Update cache objects
caches_.clear();
caches_.resize(cache_props_.size());
char name[64];
GetInfo(HSA_AGENT_INFO_NAME, name);
std::string deviceName=name;
for(size_t i=0; i<caches_.size(); i++)
caches_[i].reset(new core::Cache(deviceName+" L"+std::to_string(cache_props_[i].CacheLevel), cache_props_[i].CacheLevel, cache_props_[i].CacheSize));
}
hsa_status_t CpuAgent::VisitRegion(bool include_peer,
@@ -167,6 +176,17 @@ hsa_status_t CpuAgent::IterateRegion(
return VisitRegion(true, callback, data);
}
hsa_status_t CpuAgent::IterateCache(hsa_status_t (*callback )(hsa_cache_t cache, void *data), void* data) const
{
for(size_t i=0; i<caches_.size(); i++)
{
hsa_status_t stat = callback(core::Cache::Convert(caches_[i].get()), data);
if(stat!=HSA_STATUS_SUCCESS)
return stat;
}
return HSA_STATUS_SUCCESS;
}
hsa_status_t CpuAgent::GetInfo(hsa_agent_info_t attribute, void* value) const {
const size_t kNameSize = 64; // agent, and vendor name size limit
@@ -263,7 +283,7 @@ hsa_status_t CpuAgent::GetInfo(hsa_agent_info_t attribute, void* value) const {
*((uint16_t*)value) = 1;
break;
case HSA_AGENT_INFO_VERSION_MINOR:
*((uint16_t*)value) = 0;
*((uint16_t*)value) = 1;
break;
case HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS:
case HSA_EXT_AGENT_INFO_IMAGE_1DA_MAX_ELEMENTS:
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
+36 -12
View File
@@ -377,6 +377,15 @@ void GpuAgent::InitCacheList() {
}
}
}
//Update cache objects
caches_.clear();
caches_.resize(cache_props_.size());
char name[64];
GetInfo(HSA_AGENT_INFO_NAME, name);
std::string deviceName=name;
for(size_t i=0; i<caches_.size(); i++)
caches_[i].reset(new core::Cache(deviceName+" L"+std::to_string(cache_props_[i].CacheLevel), cache_props_[i].CacheLevel, cache_props_[i].CacheSize));
}
bool GpuAgent::InitEndTsPool() {
@@ -436,6 +445,17 @@ hsa_status_t GpuAgent::IterateRegion(
return VisitRegion(true, callback, data);
}
hsa_status_t GpuAgent::IterateCache(hsa_status_t (*callback )(hsa_cache_t cache, void *data), void* data) const
{
for(size_t i=0; i<caches_.size(); i++)
{
hsa_status_t stat = callback(core::Cache::Convert(caches_[i].get()), data);
if(stat!=HSA_STATUS_SUCCESS)
return stat;
}
return HSA_STATUS_SUCCESS;
}
hsa_status_t GpuAgent::VisitRegion(bool include_peer,
hsa_status_t (*callback)(hsa_region_t region,
void* data),
@@ -738,24 +758,28 @@ hsa_status_t GpuAgent::GetInfo(hsa_agent_info_t attribute, void* value) const {
*((hsa_isa_t*)value) = core::Isa::Handle(isa_);
break;
case HSA_AGENT_INFO_EXTENSIONS:
memset(value, 0, sizeof(uint8_t) * 128);
{
memset(value, 0, sizeof(uint8_t) * 128);
if (core::hsa_internal_api_table_.finalizer_api.hsa_ext_program_finalize_fn != NULL) {
*((uint8_t*)value) = 1 << HSA_EXTENSION_FINALIZER;
auto setFlag = [&](uint32_t bit) { assert(bit<128*8 && "Extension value exceeds extension bitmask"); uint index = bit/8; uint subBit = bit%8; ((uint8_t*)value)[index] |= 1<<subBit; };
if (core::hsa_internal_api_table_.finalizer_api.hsa_ext_program_finalize_fn != NULL) {
setFlag(HSA_EXTENSION_FINALIZER);
}
if (core::hsa_internal_api_table_.image_api.hsa_ext_image_create_fn != NULL) {
setFlag(HSA_EXTENSION_IMAGES);
}
setFlag(HSA_EXTENSION_AMD_PROFILER);
break;
}
if (core::hsa_internal_api_table_.image_api.hsa_ext_image_create_fn != NULL) {
*((uint8_t*)value) |= 1 << HSA_EXTENSION_IMAGES;
}
*((uint8_t*)value) |= 1 << HSA_EXTENSION_AMD_PROFILER;
break;
case HSA_AGENT_INFO_VERSION_MAJOR:
*((uint16_t*)value) = 1;
break;
case HSA_AGENT_INFO_VERSION_MINOR:
*((uint16_t*)value) = 0;
*((uint16_t*)value) = 1;
break;
case HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS:
case HSA_EXT_AGENT_INFO_IMAGE_1DA_MAX_ELEMENTS:
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_loader_context.cpp
+46 -33
View File
@@ -370,6 +370,26 @@ bool RegionMemory::Freeze() {
return true;
}
hsa_status_t IsIsaEquivalent(hsa_isa_t isa, void *data) {
assert(data);
std::pair<hsa_isa_t, bool> *data_pair = (std::pair<hsa_isa_t, bool>*)data;
assert(data_pair);
assert(data_pair->first.handle != 0);
assert(data_pair->second != true);
const core::Isa *isa1 = core::Isa::Object(isa);
assert(isa1);
const core::Isa *isa2 = core::Isa::Object(data_pair->first);
assert(isa2);
if (isa1->version() == isa2->version()) {
data_pair->second = true;
return HSA_STATUS_INFO_BREAK;
}
return HSA_STATUS_SUCCESS;
}
} // namespace anonymous
namespace amd {
@@ -392,25 +412,14 @@ hsa_isa_t LoaderContext::IsaFromName(const char *name) {
bool LoaderContext::IsaSupportedByAgent(hsa_agent_t agent,
hsa_isa_t code_object_isa) {
assert(agent.handle);
assert(agent.handle != 0);
hsa_status_t hsa_status = HSA_STATUS_SUCCESS;
hsa_isa_t agent_isa;
agent_isa.handle = 0;
hsa_status = HSA::hsa_agent_get_info(agent, HSA_AGENT_INFO_ISA, &agent_isa);
if (HSA_STATUS_SUCCESS != hsa_status) {
std::pair<hsa_isa_t, bool> data(code_object_isa, false);
hsa_status_t status = HSA::hsa_agent_iterate_isas(agent, IsIsaEquivalent, &data);
if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) {
return false;
}
bool result = false;
hsa_status = HSA::hsa_isa_compatible(code_object_isa, agent_isa, &result);
if (HSA_STATUS_SUCCESS != hsa_status) {
return false;
}
return result;
return data.second;
}
void* LoaderContext::SegmentAlloc(amdgpu_hsa_elf_segment_t segment,
@@ -421,14 +430,16 @@ void* LoaderContext::SegmentAlloc(amdgpu_hsa_elf_segment_t segment,
{
assert(0 < size);
assert(0 < align && 0 == (align & (align - 1)));
hsa_profile_t agent_profile;
if (HSA_STATUS_SUCCESS != HSA::hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_profile)) {
return nullptr;
}
SegmentMemory *mem = nullptr;
switch (segment) {
case AMDGPU_HSA_SEGMENT_GLOBAL_AGENT:
case AMDGPU_HSA_SEGMENT_READONLY_AGENT:
case AMDGPU_HSA_SEGMENT_READONLY_AGENT: {
hsa_profile_t agent_profile;
if (HSA_STATUS_SUCCESS != HSA::hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_profile)) {
return nullptr;
}
switch (agent_profile) {
case HSA_PROFILE_BASE:
mem = new (std::nothrow) RegionMemory(RegionMemory::AgentLocal(agent));
@@ -440,19 +451,17 @@ void* LoaderContext::SegmentAlloc(amdgpu_hsa_elf_segment_t segment,
assert(false);
}
break;
case AMDGPU_HSA_SEGMENT_GLOBAL_PROGRAM:
switch (agent_profile) {
case HSA_PROFILE_BASE:
mem = new (std::nothrow) RegionMemory(RegionMemory::System());
break;
case HSA_PROFILE_FULL:
mem = new (std::nothrow) MallocedMemory();
break;
default:
assert(false);
}
}
case AMDGPU_HSA_SEGMENT_GLOBAL_PROGRAM: {
mem = new (std::nothrow) RegionMemory(RegionMemory::System());
break;
case AMDGPU_HSA_SEGMENT_CODE_AGENT:
}
case AMDGPU_HSA_SEGMENT_CODE_AGENT: {
hsa_profile_t agent_profile;
if (HSA_STATUS_SUCCESS != HSA::hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_profile)) {
return nullptr;
}
switch (agent_profile) {
case HSA_PROFILE_BASE:
mem = new (std::nothrow) RegionMemory(IsDebuggerRegistered() ?
@@ -470,16 +479,20 @@ void* LoaderContext::SegmentAlloc(amdgpu_hsa_elf_segment_t segment,
((GpuAgentInt*)core::Agent::Convert(agent))->InvalidateCodeCaches();
break;
}
default:
assert(false);
}
if (nullptr == mem) {
return nullptr;
}
if (!mem->Allocate(size, align, zero)) {
delete mem;
return nullptr;
}
return mem;
}
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/cache.cpp
+71
View File
@@ -0,0 +1,71 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with the Software without restriction, including without limitation
// the rights to use, copy, modify, merge, publish, distribute, sublicense,
// and/or sell copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following conditions:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
// DEALINGS WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#include "core/inc/cache.h"
#include "assert.h"
namespace core
{
hsa_status_t Cache::GetInfo(hsa_cache_info_t attribute, void* value)
{
switch(attribute)
{
case HSA_CACHE_INFO_NAME_LENGTH:
*(uint32_t*)value=name_.size();
break;
case HSA_CACHE_INFO_NAME:
*(const char**)value=name_.c_str();
break;
case HSA_CACHE_INFO_LEVEL:
*(uint8_t*)value=level_;
break;
case HSA_CACHE_INFO_SIZE:
*(uint32_t*)value=size_;
break;
default:
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
return HSA_STATUS_SUCCESS;
}
}
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp
+1033 -324
View File
File diff suppressed because it is too large Load Diff
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp
+125 -46
View File
@@ -146,38 +146,38 @@ void HsaApiTable::UpdateCore() {
core_api.hsa_soft_queue_create_fn = HSA::hsa_soft_queue_create;
core_api.hsa_queue_destroy_fn = HSA::hsa_queue_destroy;
core_api.hsa_queue_inactivate_fn = HSA::hsa_queue_inactivate;
core_api.hsa_queue_load_read_index_acquire_fn =
HSA::hsa_queue_load_read_index_acquire;
core_api.hsa_queue_load_read_index_scacquire_fn =
HSA::hsa_queue_load_read_index_scacquire;
core_api.hsa_queue_load_read_index_relaxed_fn =
HSA::hsa_queue_load_read_index_relaxed;
core_api.hsa_queue_load_write_index_acquire_fn =
HSA::hsa_queue_load_write_index_acquire;
core_api.hsa_queue_load_write_index_scacquire_fn =
HSA::hsa_queue_load_write_index_scacquire;
core_api.hsa_queue_load_write_index_relaxed_fn =
HSA::hsa_queue_load_write_index_relaxed;
core_api.hsa_queue_store_write_index_relaxed_fn =
HSA::hsa_queue_store_write_index_relaxed;
core_api.hsa_queue_store_write_index_release_fn =
HSA::hsa_queue_store_write_index_release;
core_api.hsa_queue_cas_write_index_acq_rel_fn =
HSA::hsa_queue_cas_write_index_acq_rel;
core_api.hsa_queue_cas_write_index_acquire_fn =
HSA::hsa_queue_cas_write_index_acquire;
core_api.hsa_queue_store_write_index_screlease_fn =
HSA::hsa_queue_store_write_index_screlease;
core_api.hsa_queue_cas_write_index_scacq_screl_fn =
HSA::hsa_queue_cas_write_index_scacq_screl;
core_api.hsa_queue_cas_write_index_scacquire_fn =
HSA::hsa_queue_cas_write_index_scacquire;
core_api.hsa_queue_cas_write_index_relaxed_fn =
HSA::hsa_queue_cas_write_index_relaxed;
core_api.hsa_queue_cas_write_index_release_fn =
HSA::hsa_queue_cas_write_index_release;
core_api.hsa_queue_add_write_index_acq_rel_fn =
HSA::hsa_queue_add_write_index_acq_rel;
core_api.hsa_queue_add_write_index_acquire_fn =
HSA::hsa_queue_add_write_index_acquire;
core_api.hsa_queue_cas_write_index_screlease_fn =
HSA::hsa_queue_cas_write_index_screlease;
core_api.hsa_queue_add_write_index_scacq_screl_fn =
HSA::hsa_queue_add_write_index_scacq_screl;
core_api.hsa_queue_add_write_index_scacquire_fn =
HSA::hsa_queue_add_write_index_scacquire;
core_api.hsa_queue_add_write_index_relaxed_fn =
HSA::hsa_queue_add_write_index_relaxed;
core_api.hsa_queue_add_write_index_release_fn =
HSA::hsa_queue_add_write_index_release;
core_api.hsa_queue_add_write_index_screlease_fn =
HSA::hsa_queue_add_write_index_screlease;
core_api.hsa_queue_store_read_index_relaxed_fn =
HSA::hsa_queue_store_read_index_relaxed;
core_api.hsa_queue_store_read_index_release_fn =
HSA::hsa_queue_store_read_index_release;
core_api.hsa_queue_store_read_index_screlease_fn =
HSA::hsa_queue_store_read_index_screlease;
core_api.hsa_agent_iterate_regions_fn = HSA::hsa_agent_iterate_regions;
core_api.hsa_region_get_info_fn = HSA::hsa_region_get_info;
core_api.hsa_memory_register_fn = HSA::hsa_memory_register;
@@ -189,52 +189,72 @@ void HsaApiTable::UpdateCore() {
core_api.hsa_signal_create_fn = HSA::hsa_signal_create;
core_api.hsa_signal_destroy_fn = HSA::hsa_signal_destroy;
core_api.hsa_signal_load_relaxed_fn = HSA::hsa_signal_load_relaxed;
core_api.hsa_signal_load_acquire_fn = HSA::hsa_signal_load_acquire;
core_api.hsa_signal_load_scacquire_fn = HSA::hsa_signal_load_scacquire;
core_api.hsa_signal_store_relaxed_fn = HSA::hsa_signal_store_relaxed;
core_api.hsa_signal_store_release_fn = HSA::hsa_signal_store_release;
core_api.hsa_signal_store_screlease_fn = HSA::hsa_signal_store_screlease;
core_api.hsa_signal_wait_relaxed_fn = HSA::hsa_signal_wait_relaxed;
core_api.hsa_signal_wait_acquire_fn = HSA::hsa_signal_wait_acquire;
core_api.hsa_signal_wait_scacquire_fn = HSA::hsa_signal_wait_scacquire;
core_api.hsa_signal_and_relaxed_fn = HSA::hsa_signal_and_relaxed;
core_api.hsa_signal_and_acquire_fn = HSA::hsa_signal_and_acquire;
core_api.hsa_signal_and_release_fn = HSA::hsa_signal_and_release;
core_api.hsa_signal_and_acq_rel_fn = HSA::hsa_signal_and_acq_rel;
core_api.hsa_signal_and_scacquire_fn = HSA::hsa_signal_and_scacquire;
core_api.hsa_signal_and_screlease_fn = HSA::hsa_signal_and_screlease;
core_api.hsa_signal_and_scacq_screl_fn = HSA::hsa_signal_and_scacq_screl;
core_api.hsa_signal_or_relaxed_fn = HSA::hsa_signal_or_relaxed;
core_api.hsa_signal_or_acquire_fn = HSA::hsa_signal_or_acquire;
core_api.hsa_signal_or_release_fn = HSA::hsa_signal_or_release;
core_api.hsa_signal_or_acq_rel_fn = HSA::hsa_signal_or_acq_rel;
core_api.hsa_signal_or_scacquire_fn = HSA::hsa_signal_or_scacquire;
core_api.hsa_signal_or_screlease_fn = HSA::hsa_signal_or_screlease;
core_api.hsa_signal_or_scacq_screl_fn = HSA::hsa_signal_or_scacq_screl;
core_api.hsa_signal_xor_relaxed_fn = HSA::hsa_signal_xor_relaxed;
core_api.hsa_signal_xor_acquire_fn = HSA::hsa_signal_xor_acquire;
core_api.hsa_signal_xor_release_fn = HSA::hsa_signal_xor_release;
core_api.hsa_signal_xor_acq_rel_fn = HSA::hsa_signal_xor_acq_rel;
core_api.hsa_signal_xor_scacquire_fn = HSA::hsa_signal_xor_scacquire;
core_api.hsa_signal_xor_screlease_fn = HSA::hsa_signal_xor_screlease;
core_api.hsa_signal_xor_scacq_screl_fn = HSA::hsa_signal_xor_scacq_screl;
core_api.hsa_signal_exchange_relaxed_fn = HSA::hsa_signal_exchange_relaxed;
core_api.hsa_signal_exchange_acquire_fn = HSA::hsa_signal_exchange_acquire;
core_api.hsa_signal_exchange_release_fn = HSA::hsa_signal_exchange_release;
core_api.hsa_signal_exchange_acq_rel_fn = HSA::hsa_signal_exchange_acq_rel;
core_api.hsa_signal_exchange_scacquire_fn = HSA::hsa_signal_exchange_scacquire;
core_api.hsa_signal_exchange_screlease_fn = HSA::hsa_signal_exchange_screlease;
core_api.hsa_signal_exchange_scacq_screl_fn = HSA::hsa_signal_exchange_scacq_screl;
core_api.hsa_signal_add_relaxed_fn = HSA::hsa_signal_add_relaxed;
core_api.hsa_signal_add_acquire_fn = HSA::hsa_signal_add_acquire;
core_api.hsa_signal_add_release_fn = HSA::hsa_signal_add_release;
core_api.hsa_signal_add_acq_rel_fn = HSA::hsa_signal_add_acq_rel;
core_api.hsa_signal_add_scacquire_fn = HSA::hsa_signal_add_scacquire;
core_api.hsa_signal_add_screlease_fn = HSA::hsa_signal_add_screlease;
core_api.hsa_signal_add_scacq_screl_fn = HSA::hsa_signal_add_scacq_screl;
core_api.hsa_signal_subtract_relaxed_fn = HSA::hsa_signal_subtract_relaxed;
core_api.hsa_signal_subtract_acquire_fn = HSA::hsa_signal_subtract_acquire;
core_api.hsa_signal_subtract_release_fn = HSA::hsa_signal_subtract_release;
core_api.hsa_signal_subtract_acq_rel_fn = HSA::hsa_signal_subtract_acq_rel;
core_api.hsa_signal_subtract_scacquire_fn = HSA::hsa_signal_subtract_scacquire;
core_api.hsa_signal_subtract_screlease_fn = HSA::hsa_signal_subtract_screlease;
core_api.hsa_signal_subtract_scacq_screl_fn = HSA::hsa_signal_subtract_scacq_screl;
core_api.hsa_signal_cas_relaxed_fn = HSA::hsa_signal_cas_relaxed;
core_api.hsa_signal_cas_acquire_fn = HSA::hsa_signal_cas_acquire;
core_api.hsa_signal_cas_release_fn = HSA::hsa_signal_cas_release;
core_api.hsa_signal_cas_acq_rel_fn = HSA::hsa_signal_cas_acq_rel;
core_api.hsa_signal_cas_scacquire_fn = HSA::hsa_signal_cas_scacquire;
core_api.hsa_signal_cas_screlease_fn = HSA::hsa_signal_cas_screlease;
core_api.hsa_signal_cas_scacq_screl_fn = HSA::hsa_signal_cas_scacq_screl;
//===--- Instruction Set Architecture -----------------------------------===//
core_api.hsa_isa_from_name_fn = HSA::hsa_isa_from_name;
// Deprecated since v1.1.
core_api.hsa_isa_get_info_fn = HSA::hsa_isa_get_info;
// Deprecated since v1.1.
core_api.hsa_isa_compatible_fn = HSA::hsa_isa_compatible;
//===--- Code Objects (deprecated) --------------------------------------===//
// Deprecated since v1.1.
core_api.hsa_code_object_serialize_fn = HSA::hsa_code_object_serialize;
// Deprecated since v1.1.
core_api.hsa_code_object_deserialize_fn = HSA::hsa_code_object_deserialize;
// Deprecated since v1.1.
core_api.hsa_code_object_destroy_fn = HSA::hsa_code_object_destroy;
// Deprecated since v1.1.
core_api.hsa_code_object_get_info_fn = HSA::hsa_code_object_get_info;
// Deprecated since v1.1.
core_api.hsa_code_object_get_symbol_fn = HSA::hsa_code_object_get_symbol;
// Deprecated since v1.1.
core_api.hsa_code_symbol_get_info_fn = HSA::hsa_code_symbol_get_info;
// Deprecated since v1.1.
core_api.hsa_code_object_iterate_symbols_fn =
HSA::hsa_code_object_iterate_symbols;
//===--- Executable -----------------------------------------------------===//
// Deprecated since v1.1.
core_api.hsa_executable_create_fn = HSA::hsa_executable_create;
core_api.hsa_executable_destroy_fn = HSA::hsa_executable_destroy;
// Deprecated since v1.1.
core_api.hsa_executable_load_code_object_fn =
HSA::hsa_executable_load_code_object;
core_api.hsa_executable_freeze_fn = HSA::hsa_executable_freeze;
@@ -246,10 +266,69 @@ void HsaApiTable::UpdateCore() {
core_api.hsa_executable_readonly_variable_define_fn =
HSA::hsa_executable_readonly_variable_define;
core_api.hsa_executable_validate_fn = HSA::hsa_executable_validate;
// Deprecated since v1.1.
core_api.hsa_executable_get_symbol_fn = HSA::hsa_executable_get_symbol;
core_api.hsa_executable_symbol_get_info_fn = HSA::hsa_executable_symbol_get_info;
core_api.hsa_executable_iterate_symbols_fn = HSA::hsa_executable_iterate_symbols;
core_api.hsa_executable_symbol_get_info_fn =
HSA::hsa_executable_symbol_get_info;
// Deprecated since v1.1.
core_api.hsa_executable_iterate_symbols_fn =
HSA::hsa_executable_iterate_symbols;
//===--- Runtime Notifications ------------------------------------------===//
core_api.hsa_status_string_fn = HSA::hsa_status_string;
//Start HSA v1.1 additions
core_api.hsa_extension_get_name_fn = HSA::hsa_extension_get_name;
core_api.hsa_system_major_extension_supported_fn = HSA::hsa_system_major_extension_supported;
core_api.hsa_system_get_major_extension_table_fn = HSA::hsa_system_get_major_extension_table;
core_api.hsa_agent_major_extension_supported_fn = HSA::hsa_agent_major_extension_supported;
core_api.hsa_cache_get_info_fn = HSA::hsa_cache_get_info;
core_api.hsa_agent_iterate_caches_fn = HSA::hsa_agent_iterate_caches;
//Silent store optimization is present in all signal ops when no agents are sleeping.
core_api.hsa_signal_silent_store_relaxed_fn = HSA::hsa_signal_store_relaxed;
core_api.hsa_signal_silent_store_screlease_fn = HSA::hsa_signal_store_screlease;
core_api.hsa_signal_group_create_fn = HSA::hsa_signal_group_create;
core_api.hsa_signal_group_destroy_fn = HSA::hsa_signal_group_destroy;
core_api.hsa_signal_group_wait_any_scacquire_fn = HSA::hsa_signal_group_wait_any_scacquire;
core_api.hsa_signal_group_wait_any_relaxed_fn = HSA::hsa_signal_group_wait_any_relaxed;
//===--- Instruction Set Architecture - HSA v1.1 additions --------------===//
core_api.hsa_agent_iterate_isas_fn = HSA::hsa_agent_iterate_isas;
core_api.hsa_isa_get_info_alt_fn = HSA::hsa_isa_get_info_alt;
core_api.hsa_isa_get_exception_policies_fn =
HSA::hsa_isa_get_exception_policies;
core_api.hsa_isa_get_round_method_fn = HSA::hsa_isa_get_round_method;
core_api.hsa_wavefront_get_info_fn = HSA::hsa_wavefront_get_info;
core_api.hsa_isa_iterate_wavefronts_fn = HSA::hsa_isa_iterate_wavefronts;
//===--- Code Objects (deprecated) - HSA v1.1 additions -----------------===//
// Deprecated since v1.1.
core_api.hsa_code_object_get_symbol_from_name_fn =
HSA::hsa_code_object_get_symbol_from_name;
//===--- Executable - HSA v1.1 additions --------------------------------===//
core_api.hsa_code_object_reader_create_from_file_fn =
HSA::hsa_code_object_reader_create_from_file;
core_api.hsa_code_object_reader_create_from_memory_fn =
HSA::hsa_code_object_reader_create_from_memory;
core_api.hsa_code_object_reader_destroy_fn =
HSA::hsa_code_object_reader_destroy;
core_api.hsa_executable_create_alt_fn = HSA::hsa_executable_create_alt;
core_api.hsa_executable_load_program_code_object_fn =
HSA::hsa_executable_load_program_code_object;
core_api.hsa_executable_load_agent_code_object_fn =
HSA::hsa_executable_load_agent_code_object;
core_api.hsa_executable_validate_alt_fn = HSA::hsa_executable_validate_alt;
core_api.hsa_executable_get_symbol_by_name_fn =
HSA::hsa_executable_get_symbol_by_name;
core_api.hsa_executable_iterate_agent_symbols_fn =
HSA::hsa_executable_iterate_agent_symbols;
core_api.hsa_executable_iterate_program_symbols_fn =
HSA::hsa_executable_iterate_program_symbols;
}
// Update Api table for Amd Extensions.
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/isa.cpp
+86 -9
View File
@@ -47,6 +47,24 @@
namespace core {
bool Wavefront::GetInfo(
const hsa_wavefront_info_t &attribute,
void *value) const {
if (!value) {
return false;
}
switch (attribute) {
case HSA_WAVEFRONT_INFO_SIZE: {
*((uint32_t*)value) = 64;
return true;
}
default: {
return false;
}
}
}
std::string Isa::GetFullName() const {
std::stringstream full_name;
full_name << GetVendor() << ":" << GetArchitecture() << ":"
@@ -63,27 +81,82 @@ bool Isa::GetInfo(const hsa_isa_info_t &attribute, void *value) const {
switch (attribute) {
case HSA_ISA_INFO_NAME_LENGTH: {
std::string full_name = GetFullName();
*((uint32_t *)value) = static_cast<uint32_t>(full_name.size());
*((uint32_t*)value) = static_cast<uint32_t>(full_name.size() + 1);
return true;
}
case HSA_ISA_INFO_NAME: {
std::string full_name = GetFullName();
memset(value, 0x0, full_name.size() + 1);
memcpy(value, full_name.c_str(), full_name.size());
return true;
}
// @todo: following case needs to be removed.
// deprecated.
case HSA_ISA_INFO_CALL_CONVENTION_COUNT: {
*((uint32_t *)value) = 1;
*((uint32_t*)value) = 1;
return true;
}
// @todo: following case needs to be removed.
// deprecated.
case HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONT_SIZE: {
*((uint32_t *)value) = 64;
*((uint32_t*)value) = 64;
return true;
}
// @todo: following needs to be removed.
// deprecated.
case HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONTS_PER_COMPUTE_UNIT: {
*((uint32_t *)value) = 40;
*((uint32_t*)value) = 40;
return true;
}
case HSA_ISA_INFO_MACHINE_MODELS: {
const bool machine_models[2] = {false, true};
memcpy(value, machine_models, sizeof(machine_models));
return true;
}
case HSA_ISA_INFO_PROFILES: {
bool profiles[2] = {true, false};
if (this->version() == Version(7, 0, 0) ||
this->version() == Version(8, 0, 1)) {
profiles[1] = true;
}
memcpy(value, profiles, sizeof(profiles));
return true;
}
case HSA_ISA_INFO_DEFAULT_FLOAT_ROUNDING_MODES: {
const bool rounding_modes[3] = {false, false, true};
memcpy(value, rounding_modes, sizeof(rounding_modes));
return true;
}
case HSA_ISA_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES: {
const bool rounding_modes[3] = {false, false, true};
memcpy(value, rounding_modes, sizeof(rounding_modes));
return true;
}
case HSA_ISA_INFO_FAST_F16_OPERATION: {
if (this->GetMajorVersion() >= 8) {
*((bool*)value) = true;
} else {
*((bool*)value) = false;
}
return true;
}
case HSA_ISA_INFO_WORKGROUP_MAX_DIM: {
const uint16_t workgroup_max_dim[3] = {1024, 1024, 1024};
memcpy(value, workgroup_max_dim, sizeof(workgroup_max_dim));
return true;
}
case HSA_ISA_INFO_WORKGROUP_MAX_SIZE: {
*((uint32_t*)value) = 1024;
return true;
}
case HSA_ISA_INFO_GRID_MAX_DIM: {
const hsa_dim3_t grid_max_dim = {UINT32_MAX, UINT32_MAX, UINT32_MAX};
memcpy(value, &grid_max_dim, sizeof(grid_max_dim));
return true;
}
case HSA_ISA_INFO_GRID_MAX_SIZE: {
*((uint64_t*)value) = UINT64_MAX;
return true;
}
case HSA_ISA_INFO_FBARRIER_MAX_SIZE: {
*((uint32_t*)value) = 32;
return true;
}
default: {
@@ -92,6 +165,12 @@ bool Isa::GetInfo(const hsa_isa_info_t &attribute, void *value) const {
}
}
hsa_round_method_t Isa::GetRoundMethod(
hsa_fp_type_t fp_type,
hsa_flush_mode_t flush_mode) const {
return HSA_ROUND_METHOD_SINGLE;
}
const Isa *IsaRegistry::GetIsa(const std::string &full_name) {
auto isareg_iter = supported_isas_.find(full_name);
return isareg_iter == supported_isas_.end() ? nullptr : &isareg_iter->second;
@@ -117,11 +196,9 @@ const IsaRegistry::IsaMap IsaRegistry::GetSupportedIsas() {
ISAREG_ENTRY_GEN(7, 0, 0)
ISAREG_ENTRY_GEN(7, 0, 1)
ISAREG_ENTRY_GEN(8, 0, 0)
ISAREG_ENTRY_GEN(8, 0, 1)
ISAREG_ENTRY_GEN(8, 0, 2)
ISAREG_ENTRY_GEN(8, 0, 3)
ISAREG_ENTRY_GEN(8, 1, 0)
ISAREG_ENTRY_GEN(9, 0, 0)
return supported_isas;
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp
+16 -12
View File
@@ -63,7 +63,7 @@
#include "core/inc/hsa_api_trace_int.h"
#define HSA_VERSION_MAJOR 1
#define HSA_VERSION_MINOR 0
#define HSA_VERSION_MINOR 1
const char rocrbuildid[] = "ROCR BUILD ID: " STRING(ROCR_BUILD_ID);
@@ -526,19 +526,23 @@ hsa_status_t Runtime::GetSystemInfo(hsa_system_info_t attribute, void* value) {
#endif
break;
case HSA_SYSTEM_INFO_EXTENSIONS:
memset(value, 0, sizeof(uint8_t) * 128);
{
memset(value, 0, sizeof(uint8_t) * 128);
if (hsa_internal_api_table_.finalizer_api.hsa_ext_program_finalize_fn != NULL) {
*((uint8_t*)value) = 1 << HSA_EXTENSION_FINALIZER;
auto setFlag = [&](uint32_t bit) { assert(bit<128*8 && "Extension value exceeds extension bitmask"); uint index = bit/8; uint subBit = bit%8; ((uint8_t*)value)[index] |= 1<<subBit; };
if (hsa_internal_api_table_.finalizer_api.hsa_ext_program_finalize_fn != NULL) {
setFlag(HSA_EXTENSION_FINALIZER);
}
if (hsa_internal_api_table_.image_api.hsa_ext_image_create_fn != NULL) {
setFlag(HSA_EXTENSION_IMAGES);
}
setFlag(HSA_EXTENSION_AMD_PROFILER);
break;
}
if (hsa_internal_api_table_.image_api.hsa_ext_image_create_fn != NULL) {
*((uint8_t*)value) |= 1 << HSA_EXTENSION_IMAGES;
}
*((uint8_t*)value) |= 1 << HSA_EXTENSION_AMD_PROFILER;
break;
default:
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/signal.cpp
+16 -4
View File
@@ -49,13 +49,13 @@
namespace core {
uint32_t Signal::WaitAny(uint32_t signal_count, hsa_signal_t* hsa_signals,
hsa_signal_condition_t* conds,
hsa_signal_value_t* values, uint64_t timeout,
uint32_t Signal::WaitAny(uint32_t signal_count, const hsa_signal_t* hsa_signals,
const hsa_signal_condition_t* conds,
const hsa_signal_value_t* values, uint64_t timeout,
hsa_wait_state_t wait_hint,
hsa_signal_value_t* satisfying_value) {
hsa_signal_handle* signals =
reinterpret_cast<hsa_signal_handle*>(hsa_signals);
reinterpret_cast<hsa_signal_handle*>(const_cast<hsa_signal_t*>(hsa_signals));
uint32_t prior = 0;
for (uint32_t i = 0; i < signal_count; i++)
prior = Max(prior, atomic::Increment(&signals[i]->waiting_));
@@ -182,6 +182,18 @@ uint32_t Signal::WaitAny(uint32_t signal_count, hsa_signal_t* hsa_signals,
}
}
SignalGroup::SignalGroup(uint32_t num_signals, const hsa_signal_t* hsa_signals) : count(num_signals)
{
if(count!=0)
signals=new hsa_signal_t[count];
else
signals=NULL;
if(signals==NULL)
return;
for(int i=0; i<count; i++)
signals[i]=hsa_signals[i];
}
} // namespace core
#endif // header guard