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

Split ihipCtx_t into ihipCtx_t and ihipDevice_t .

Browse Source 
Major change to existing code base.
    Ctx holds streams, enables peers, and flags.
    Device holds accelerator, hsa-agent, device props.

Add hipCtx_t.

Add peer APIs that accept hipCtx_t (in addition to deviceId)

Compiles and passes directed tests.

Change-Id: Iddab1eb9edbf90caad2ef5959c6b811d658197f1


[ROCm/hip commit: cfdacab32f]
This commit is contained in:
Ben Sander
2016-08-08 11:55:57 -05:00
parent 1f9d2201fe
commit 57df8a3967
7 changed files with 687 additions and 531 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/src/hip_device.cpp
+8 -5
View File
@@ -283,15 +283,18 @@ hipError_t hipSetDeviceFlags( unsigned int flags)
hipError_t e;
auto * hipDevice = ihipGetTlsDefaultCtx();
if(hipDevice){
hipDevice->_device_flags = hipDevice->_device_flags | flags;
auto * ctx = ihipGetTlsDefaultCtx();
// TODO : does this really OR in the flags or replaces previous flags:
// TODO : Review error handling behavior for this function, it often returns ErrorSetOnActiveProcess
if (ctx) {
ctx->_ctxFlags = ctx->_ctxFlags | flags;
e = hipSuccess;
}else{
} else {
e = hipErrorInvalidDevice;
}
return ihipLogStatus(e);
}
};
projects/hip/src/hip_hcc.cpp
+253 -198
View File
@@ -48,6 +48,11 @@ THE SOFTWARE.
extern const char *ihipErrorString(hipError_t hip_error);
#include "hcc_detail/trace_helper.h"
//=================================================================================================
//Global variables:
//=================================================================================================
const int release = 1;
#define MEMCPY_D2H_STAGING_VS_PININPLACE_COPY_THRESHOLD 4194304
@@ -81,19 +86,21 @@ thread_local hipError_t tls_lastHipError = hipSuccess;
//=================================================================================================
//Forward Declarations:
//=================================================================================================
std::once_flag hip_initialized;
// Array of primary contexts for each device:
ihipCtx_t *g_primaryCtxArray; ;
// Array of pointers to devices.
ihipDevice_t **g_deviceArray;
bool g_visible_device = false;
unsigned g_deviceCnt;
std::vector<int> g_hip_visible_devices;
hsa_agent_t g_cpu_agent;
// TODO, remove these if possible:
hsa_agent_t gpu_agent_;
hsa_amd_memory_pool_t gpu_pool_;
//=================================================================================================
// Implementation:
@@ -106,29 +113,29 @@ static inline bool ihipIsValidDevice(unsigned deviceIndex)
return (deviceIndex < g_deviceCnt);
}
//---
ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex)
ihipDevice_t * ihipGetDevice(int deviceIndex)
{
if (ihipIsValidDevice(deviceIndex)) {
return &g_primaryCtxArray[deviceIndex];
} else {
return NULL;
}
};
// FIXME- index the new g_deviceArray data structure
ihipCtx_t * ihipGetDevice(int deviceIndex)
{
if (ihipIsValidDevice(deviceIndex)) {
return &g_primaryCtxArray[deviceIndex];
return g_deviceArray[deviceIndex];
} else {
return NULL;
}
}
//---
//FIXME - is this function dead?
ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex)
{
ihipDevice_t *device = ihipGetDevice(deviceIndex);
return device ? device->getPrimaryCtx() : NULL;
};
//---
//FIXME - this needs to return the active context for this CPU thread - not primary for device.
ihipCtx_t *ihipGetTlsDefaultCtx()
{
// If this is invalid, the TLS state is corrupt.
@@ -136,10 +143,11 @@ ihipCtx_t *ihipGetTlsDefaultCtx()
// TODO - consider replacing assert with error code
assert (ihipIsValidDevice(tls_defaultDevice));
return &g_primaryCtxArray[tls_defaultDevice];
return ihipGetPrimaryCtx(tls_defaultDevice);
}
//=================================================================================================
// ihipSignal_t:
//=================================================================================================
@@ -249,84 +257,27 @@ void ihipStream_t::locked_wait(bool assertQueueEmpty)
};
// Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted.
// The packed _peerAgents can efficiently be used on each memory allocation.
template<>
void ihipDeviceCriticalBase_t<DeviceMutex>::recomputePeerAgents()
{
_peerCnt = 0;
std::for_each (_peers.begin(), _peers.end(), [this](ihipCtx_t* ctx) {
_peerAgents[_peerCnt++] = ctx->_hsa_agent;
});
}
//=============================================================================
template<>
bool ihipDeviceCriticalBase_t<DeviceMutex>::isPeer(const ihipCtx_t *peer)
{
auto match = std::find(_peers.begin(), _peers.end(), peer);
return (match != std::end(_peers));
}
template<>
bool ihipDeviceCriticalBase_t<DeviceMutex>::addPeer(ihipCtx_t *peer)
{
auto match = std::find(_peers.begin(), _peers.end(), peer);
if (match == std::end(_peers)) {
// Not already a peer, let's update the list:
_peers.push_back(peer);
recomputePeerAgents();
return true;
}
// If we get here - peer was already on list, silently ignore.
return false;
}
template<>
bool ihipDeviceCriticalBase_t<DeviceMutex>::removePeer(ihipCtx_t *peer)
{
auto match = std::find(_peers.begin(), _peers.end(), peer);
if (match != std::end(_peers)) {
// Found a valid peer, let's remove it.
_peers.remove(peer);
recomputePeerAgents();
return true;
} else {
return false;
}
}
template<>
void ihipDeviceCriticalBase_t<DeviceMutex>::resetPeers(ihipCtx_t *thisDevice)
{
_peers.clear();
_peerCnt = 0;
addPeer(thisDevice); // peer-list always contains self agent.
}
template<>
void ihipDeviceCriticalBase_t<DeviceMutex>::addStream(ihipStream_t *stream)
{
_streams.push_back(stream);
stream->_id = incStreamId();
}
//-------------------------------------------------------------------------------------------------
//---
ihipCtx_t * ihipStream_t::getDevice() const
const ihipDevice_t * ihipStream_t::getDevice() const
{
return _ctx->getDevice();
};
ihipCtx_t * ihipStream_t::getCtx() const
{
return _ctx;
};
#define HIP_NUM_SIGNALS_PER_STREAM 32
@@ -521,56 +472,80 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t
//=================================================================================================
//
//Reset the device - this is called from hipDeviceReset.
//Device may be reset multiple times, and may be reset after init.
void ihipCtx_t::locked_reset()
//=============================================================================
// Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted.
// The packed _peerAgents can efficiently be used on each memory allocation.
template<>
void ihipCtxCriticalBase_t<CtxMutex>::recomputePeerAgents()
{
// Obtain mutex access to the device critical data, release by destructor
LockedAccessor_DeviceCrit_t crit(_criticalData);
_peerCnt = 0;
std::for_each (_peers.begin(), _peers.end(), [this](ihipCtx_t* ctx) {
_peerAgents[_peerCnt++] = ctx->getDevice()->_hsa_agent;
});
}
//---
//Wait for pending activity to complete? TODO - check if this is required behavior:
tprintf(DB_SYNC, "locked_reset waiting for activity to complete.\n");
template<>
bool ihipCtxCriticalBase_t<CtxMutex>::isPeer(const ihipCtx_t *peer)
{
auto match = std::find(_peers.begin(), _peers.end(), peer);
return (match != std::end(_peers));
}
// Reset and remove streams:
// Delete all created streams including the default one.
for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) {
ihipStream_t *stream = *streamI;
(*streamI)->locked_wait();
tprintf(DB_SYNC, " delete stream=%p\n", stream);
delete stream;
template<>
bool ihipCtxCriticalBase_t<CtxMutex>::addPeer(ihipCtx_t *peer)
{
auto match = std::find(_peers.begin(), _peers.end(), peer);
if (match == std::end(_peers)) {
// Not already a peer, let's update the list:
_peers.push_back(peer);
recomputePeerAgents();
return true;
}
// Clear the list.
crit->streams().clear();
// If we get here - peer was already on list, silently ignore.
return false;
}
// Create a fresh default stream and add it:
_default_stream = new ihipStream_t(this, _acc.get_default_view(), hipStreamDefault);
crit->addStream(_default_stream);
// This resest peer list to just me:
crit->resetPeers(this);
// Reset and release all memory stored in the tracker:
// Reset will remove peer mapping so don't need to do this explicitly.
am_memtracker_reset(_acc);
};
//---
void ihipCtx_t::init(unsigned device_index, unsigned deviceCnt, hc::accelerator &acc, unsigned flags)
template<>
bool ihipCtxCriticalBase_t<CtxMutex>::removePeer(ihipCtx_t *peer)
{
_device_index = device_index;
_device_flags = flags;
_acc = acc;
auto match = std::find(_peers.begin(), _peers.end(), peer);
if (match != std::end(_peers)) {
// Found a valid peer, let's remove it.
_peers.remove(peer);
recomputePeerAgents();
return true;
} else {
return false;
}
}
template<>
void ihipCtxCriticalBase_t<CtxMutex>::resetPeers(ihipCtx_t *thisDevice)
{
_peers.clear();
_peerCnt = 0;
addPeer(thisDevice); // peer-list always contains self agent.
}
template<>
void ihipCtxCriticalBase_t<CtxMutex>::addStream(ihipStream_t *stream)
{
stream->_id = _streams.size();
_streams.push_back(stream);
}
//=============================================================================
//==============================================================================================
ihipDevice_t::ihipDevice_t(unsigned device_index, unsigned deviceCnt, hc::accelerator &acc) :
_device_index(device_index),
_acc(acc)
{
hsa_agent_t *agent = static_cast<hsa_agent_t*> (acc.get_hsa_agent());
if (agent) {
int err = hsa_agent_get_info(*agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, &_compute_units);
@@ -583,30 +558,17 @@ void ihipCtx_t::init(unsigned device_index, unsigned deviceCnt, hc::accelerator
_hsa_agent.handle = static_cast<uint64_t> (-1);
}
getProperties(&_props);
_criticalData.init(deviceCnt);
locked_reset();
tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream);
initProperties(&_props);
_staging_buffer[0] = new StagingBuffer(_hsa_agent,g_cpu_agent, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS);
_staging_buffer[1] = new StagingBuffer(_hsa_agent,g_cpu_agent, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS);
};
_primaryCtx = new ihipCtx_t(this, deviceCnt, hipDeviceMapHost);
}
ihipCtx_t::~ihipCtx_t()
ihipDevice_t::~ihipDevice_t()
{
if (_default_stream) {
delete _default_stream;
_default_stream = NULL;
}
for (int i=0; i<2; i++) {
if (_staging_buffer[i]) {
delete _staging_buffer[i];
@@ -615,19 +577,8 @@ ihipCtx_t::~ihipCtx_t()
}
}
//----
//=================================================================================================
// Utility functions, these are not part of the public HIP API
//=================================================================================================
//=================================================================================================
#define DeviceErrorCheck(x) if (x != HSA_STATUS_SUCCESS) { return hipErrorInvalidDevice; }
#define ErrorCheck(x) error_check(x, __LINE__, __FILE__)
void error_check(hsa_status_t hsa_error_code, int line_num, std::string str) {
@@ -636,8 +587,28 @@ void error_check(hsa_status_t hsa_error_code, int line_num, std::string str) {
}
}
hsa_agent_t gpu_agent_;
hsa_amd_memory_pool_t gpu_pool_;
//---
// Helper for initProperties
// Determines if the given agent is of type HSA_DEVICE_TYPE_GPU and counts it.
static hsa_status_t countGpuAgents(hsa_agent_t agent, void *data) {
if (data == NULL) {
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
hsa_device_type_t device_type;
hsa_status_t status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type);
if (status != HSA_STATUS_SUCCESS) {
return status;
}
if (device_type == HSA_DEVICE_TYPE_GPU) {
(*static_cast<int*>(data))++;
}
return HSA_STATUS_SUCCESS;
}
hsa_status_t FindGpuDevice(hsa_agent_t agent, void* data) {
if (data == NULL) {
@@ -717,24 +688,30 @@ hsa_status_t get_region_info(hsa_region_t region, void* data)
return HSA_STATUS_SUCCESS;
}
// Determines if the given agent is of type HSA_DEVICE_TYPE_GPU and counts it.
static hsa_status_t countGpuAgents(hsa_agent_t agent, void *data) {
if (data == NULL) {
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
static hsa_status_t findCpuAgent(hsa_agent_t agent, void *data)
{
hsa_device_type_t device_type;
hsa_status_t status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type);
if (status != HSA_STATUS_SUCCESS) {
return status;
}
if (device_type == HSA_DEVICE_TYPE_GPU) {
(*static_cast<int*>(data))++;
if (device_type == HSA_DEVICE_TYPE_CPU) {
(*static_cast<hsa_agent_t*>(data)) = agent;
return HSA_STATUS_INFO_BREAK;
}
return HSA_STATUS_SUCCESS;
}
// Internal version,
hipError_t ihipCtx_t::getProperties(hipDeviceProp_t* prop)
#define DeviceErrorCheck(x) if (x != HSA_STATUS_SUCCESS) { return hipErrorInvalidDevice; }
//---
// Initialize properties for the device.
// Call this once when the ihipDevice_t is created:
hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop)
{
hipError_t e = hipSuccess;
hsa_status_t err;
@@ -898,22 +875,111 @@ hipError_t ihipCtx_t::getProperties(hipDeviceProp_t* prop)
prop->arch.has3dGrid = 1;
prop->arch.hasDynamicParallelism = 0;
prop->concurrentKernels = 1; // All ROCR hardware supports executing multiple kernels concurrently
prop->concurrentKernels = 1; // All ROCm hardware supports executing multiple kernels concurrently
prop->canMapHostMemory = 1; // All ROCm devices can map host memory
#if 0
// TODO - code broken below since it always returns 1.
// Are the flags part of the context or part of the device?
if ( _device_flags | hipDeviceMapHost) {
prop->canMapHostMemory = 1;
} else {
prop->canMapHostMemory = 0;
}
#endif
return e;
}
//=================================================================================================
// ihipDevice_t
//=================================================================================================
//---
ihipCtx_t::ihipCtx_t(const ihipDevice_t *device, unsigned deviceCnt, unsigned flags) :
_ctxFlags(flags),
_device(device),
_criticalData(deviceCnt)
{
locked_reset();
tprintf(DB_SYNC, "created ctx with default_stream=%p\n", _default_stream);
};
ihipCtx_t::~ihipCtx_t()
{
if (_default_stream) {
delete _default_stream;
_default_stream = NULL;
}
}
//Reset the device - this is called from hipDeviceReset.
//Device may be reset multiple times, and may be reset after init.
void ihipCtx_t::locked_reset()
{
// Obtain mutex access to the device critical data, release by destructor
LockedAccessor_CtxCrit_t crit(_criticalData);
//---
//Wait for pending activity to complete? TODO - check if this is required behavior:
tprintf(DB_SYNC, "locked_reset waiting for activity to complete.\n");
// Reset and remove streams:
// Delete all created streams including the default one.
for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) {
ihipStream_t *stream = *streamI;
(*streamI)->locked_wait();
tprintf(DB_SYNC, " delete stream=%p\n", stream);
delete stream;
}
// Clear the list.
crit->streams().clear();
// Create a fresh default stream and add it:
_default_stream = new ihipStream_t(this, getDevice()->_acc.get_default_view(), hipStreamDefault);
crit->addStream(_default_stream);
// Reset peer list to just me:
crit->resetPeers(this);
// Reset and release all memory stored in the tracker:
// Reset will remove peer mapping so don't need to do this explicitly.
// FIXME - This is clearly a non-const action! Is this a context reset or a device reset - maybe should reference count?
ihipDevice_t *device = const_cast<ihipDevice_t*> (getDevice());
am_memtracker_reset(device->_acc);
};
//----
//=================================================================================================
// Utility functions, these are not part of the public HIP API
//=================================================================================================
//=================================================================================================
// Implement "default" stream syncronization
// This waits for all other streams to drain before continuing.
// If waitOnSelf is set, this additionally waits for the default stream to empty.
void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf)
{
LockedAccessor_DeviceCrit_t crit(_criticalData);
LockedAccessor_CtxCrit_t crit(_criticalData);
tprintf(DB_SYNC, "syncDefaultStream\n");
@@ -936,7 +1002,7 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf)
//---
void ihipCtx_t::locked_addStream(ihipStream_t *s)
{
LockedAccessor_DeviceCrit_t crit(_criticalData);
LockedAccessor_CtxCrit_t crit(_criticalData);
crit->addStream(s);
}
@@ -944,7 +1010,7 @@ void ihipCtx_t::locked_addStream(ihipStream_t *s)
//---
void ihipCtx_t::locked_removeStream(ihipStream_t *s)
{
LockedAccessor_DeviceCrit_t crit(_criticalData);
LockedAccessor_CtxCrit_t crit(_criticalData);
crit->streams().remove(s);
}
@@ -954,7 +1020,7 @@ void ihipCtx_t::locked_removeStream(ihipStream_t *s)
//Heavyweight synchronization that waits on all streams, ignoring hipStreamNonBlocking flag.
void ihipCtx_t::locked_waitAllStreams()
{
LockedAccessor_DeviceCrit_t crit(_criticalData);
LockedAccessor_CtxCrit_t crit(_criticalData);
tprintf(DB_SYNC, "waitAllStream\n");
for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) {
@@ -1029,21 +1095,6 @@ void ihipReadEnv_I(int *var_ptr, const char *var_name1, const char *var_name2, c
#endif
// Determines if the given agent is of type HSA_DEVICE_TYPE_GPU and counts it.
static hsa_status_t findCpuAgent(hsa_agent_t agent, void *data)
{
hsa_device_type_t device_type;
hsa_status_t status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type);
if (status != HSA_STATUS_SUCCESS) {
return status;
}
if (device_type == HSA_DEVICE_TYPE_CPU) {
(*static_cast<hsa_agent_t*>(data)) = agent;
return HSA_STATUS_INFO_BREAK;
}
return HSA_STATUS_SUCCESS;
}
//---
@@ -1145,7 +1196,7 @@ void ihipInit()
throw ihipException(hipErrorRuntimeOther);
}
g_primaryCtxArray = new ihipCtx_t[deviceCnt];
g_deviceArray = new ihipDevice_t* [deviceCnt];
g_deviceCnt = 0;
for (int i=0; i<accs.size(); i++) {
// check if the device id is included in the HIP_VISIBLE_DEVICES env variable
@@ -1155,7 +1206,7 @@ void ihipInit()
//If device is not in visible devices list, ignore
continue;
}
g_primaryCtxArray[g_deviceCnt].init(g_deviceCnt, deviceCnt, accs[i], hipDeviceMapHost);
g_deviceArray[g_deviceCnt] = new ihipDevice_t(g_deviceCnt, deviceCnt, accs[i]);
g_deviceCnt++;
}
}
@@ -1190,7 +1241,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
// Have to wait for legacy default stream to be empty:
if (!(stream->_flags & hipStreamNonBlocking)) {
tprintf(DB_SYNC, "stream %p wait default stream\n", stream);
stream->getDevice()->_default_stream->locked_wait();
stream->getCtx()->_default_stream->locked_wait();
}
return stream;
@@ -1428,7 +1479,7 @@ unsigned ihipStream_t::resolveMemcpyDirection(bool srcTracked, bool dstTracked,
void ihipStream_t::setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent)
{
// current* represents the device associated with the specified stream.
ihipCtx_t *streamDevice = this->getDevice();
const ihipDevice_t *streamDevice = this->getDevice();
hsa_agent_t streamAgent = streamDevice->_hsa_agent;
// ROCR runtime logic is :
@@ -1448,7 +1499,9 @@ void ihipStream_t::setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType,
void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind)
{
ihipCtx_t *device = this->getDevice();
ihipCtx_t *ctx = this->getCtx();
const ihipDevice_t *device = ctx->getDevice();
if (device == NULL) {
throw ihipException(hipErrorInvalidDevice);
}
@@ -1472,9 +1525,11 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
bool copyEngineCanSeeSrcAndDest = false;
if (kind == hipMemcpyDeviceToDevice) {
#if USE_PEER_TO_PEER>=2
// TODO - consider refactor. Do we need to support simul access of enable/disable peers with access?
LockedAccessor_DeviceCrit_t dcrit(device->criticalData());
if (dcrit->isPeer(ihipGetDevice(dstPtrInfo._appId)) && (dcrit->isPeer(ihipGetDevice(srcPtrInfo._appId)))) {
// Lock to prevent another thread from modifying peer list while we are trying to look at it.
LockedAccessor_CtxCrit_t dcrit(ctx->criticalData());
// FIXME - this assumes peer access only from primary context.
// Would need to change the tracker to store a void * parameter that we could map to the ctx where the pointer is allocated.
if (dcrit->isPeer(ihipGetPrimaryCtx(dstPtrInfo._appId)) && (dcrit->isPeer(ihipGetPrimaryCtx(srcPtrInfo._appId)))) {
copyEngineCanSeeSrcAndDest = true;
}
#endif
@@ -1658,9 +1713,9 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig
{
LockedAccessor_StreamCrit_t crit(_criticalData);
ihipCtx_t *device = this->getDevice();
const ihipCtx_t *ctx = this->getCtx();
if (device == NULL) {
if ((ctx == nullptr) || (ctx->getDevice() == nullptr)) {
throw ihipException(hipErrorInvalidDevice);
}
@@ -1744,12 +1799,12 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc)
{
HIP_INIT_API(deviceId, acc);
ihipCtx_t *d = ihipGetDevice(deviceId);
const ihipDevice_t *device = ihipGetDevice(deviceId);
hipError_t err;
if (d == NULL) {
if (device == NULL) {
err = hipErrorInvalidDevice;
} else {
*acc = d->_acc;
*acc = device->_acc;
err = hipSuccess;
}
return ihipLogStatus(err);
projects/hip/src/hip_memory.cpp
+225 -214
View File
@@ -1,21 +1,21 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in 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:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in 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:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
@@ -120,18 +120,19 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
auto ctx = ihipGetTlsDefaultCtx();
if (ctx) {
auto device = ctx->getWriteableDevice();
const unsigned am_flags = 0;
*ptr = hc::am_alloc(sizeBytes, ctx->_acc, am_flags);
*ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);
if (sizeBytes && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
} else {
hc::am_memtracker_update(*ptr, ctx->_device_index, 0);
hc::am_memtracker_update(*ptr, device->_device_index, 0);
{
LockedAccessor_DeviceCrit_t crit(ctx->criticalData());
LockedAccessor_CtxCrit_t crit(ctx->criticalData());
if (crit->peerCnt()) {
hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
}
@@ -155,22 +156,24 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
auto ctx = ihipGetTlsDefaultCtx();
if(ctx){
// am_alloc requires writeable __acc, perhaps could be refactored?
auto device = ctx->getWriteableDevice();
if(flags == hipHostMallocDefault){
*ptr = hc::am_alloc(sizeBytes, ctx->_acc, amHostPinned);
*ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
if(sizeBytes < 1 && (*ptr == NULL)){
hip_status = hipErrorMemoryAllocation;
}else{
hc::am_memtracker_update(*ptr, ctx->_device_index, amHostPinned);
} else {
hc::am_memtracker_update(*ptr, device->_device_index, amHostPinned);
}
tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
} else if(flags & hipHostMallocMapped){
*ptr = hc::am_alloc(sizeBytes, ctx->_acc, amHostPinned);
*ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
if(sizeBytes && (*ptr == NULL)){
hip_status = hipErrorMemoryAllocation;
}else{
hc::am_memtracker_update(*ptr, ctx->_device_index, flags);
hc::am_memtracker_update(*ptr, device->_device_index, flags);
{
LockedAccessor_DeviceCrit_t crit(ctx->criticalData());
LockedAccessor_CtxCrit_t crit(ctx->criticalData());
if (crit->peerCnt()) {
hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
}
@@ -199,58 +202,62 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
}
// width in bytes
hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) {
hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height)
{
HIP_INIT_API(ptr, pitch, width, height);
HIP_INIT_API(ptr, pitch, width, height);
hipError_t hip_status = hipSuccess;
hipError_t hip_status = hipSuccess;
if(width == 0 || height == 0)
return ihipLogStatus(hipErrorUnknown);
if(width == 0 || height == 0)
return ihipLogStatus(hipErrorUnknown);
// hardcoded 128 bytes
*pitch = ((((int)width-1)/128) + 1)*128;
const size_t sizeBytes = (*pitch)*height;
// hardcoded 128 bytes
*pitch = ((((int)width-1)/128) + 1)*128;
const size_t sizeBytes = (*pitch)*height;
auto ctx = ihipGetTlsDefaultCtx();
auto ctx = ihipGetTlsDefaultCtx();
//err = hipMalloc(ptr, (*pitch)*height);
if (ctx) {
auto device = ctx->getWriteableDevice();
//err = hipMalloc(ptr, (*pitch)*height);
if (ctx) {
const unsigned am_flags = 0;
*ptr = hc::am_alloc(sizeBytes, ctx->_acc, am_flags);
const unsigned am_flags = 0;
*ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);
if (sizeBytes && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
} else {
hc::am_memtracker_update(*ptr, ctx->_device_index, 0);
{
LockedAccessor_DeviceCrit_t crit(ctx->criticalData());
if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved:
hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
if (hsa_status != HSA_STATUS_SUCCESS) {
if (sizeBytes && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
}
} else {
hc::am_memtracker_update(*ptr, device->_device_index, 0);
{
LockedAccessor_CtxCrit_t crit(ctx->criticalData());
if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved:
hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
if (hsa_status != HSA_STATUS_SUCCESS) {
hip_status = hipErrorMemoryAllocation;
}
}
}
}
}
} else {
hip_status = hipErrorMemoryAllocation;
}
} else {
hip_status = hipErrorMemoryAllocation;
}
return ihipLogStatus(hip_status);
return ihipLogStatus(hip_status);
}
hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) {
hipChannelFormatDesc cd;
cd.x = x; cd.y = y; cd.z = z; cd.w = w;
cd.f = f;
return cd;
hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f)
{
hipChannelFormatDesc cd;
cd.x = x; cd.y = y; cd.z = z; cd.w = w;
cd.f = f;
return cd;
}
hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
size_t width, size_t height, unsigned int flags) {
size_t width, size_t height, unsigned int flags)
{
HIP_INIT_API(array, desc, width, height, flags);
hipError_t hip_status = hipSuccess;
@@ -266,40 +273,41 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
void ** ptr = &array[0]->data;
if (ctx) {
const unsigned am_flags = 0;
const size_t size = width*height;
auto device = ctx->getWriteableDevice();
const unsigned am_flags = 0;
const size_t size = width*height;
switch(desc->f) {
case hipChannelFormatKindSigned:
*ptr = hc::am_alloc(size*sizeof(int), ctx->_acc, am_flags);
break;
case hipChannelFormatKindUnsigned:
*ptr = hc::am_alloc(size*sizeof(unsigned int), ctx->_acc, am_flags);
break;
case hipChannelFormatKindFloat:
*ptr = hc::am_alloc(size*sizeof(float), ctx->_acc, am_flags);
break;
case hipChannelFormatKindNone:
*ptr = hc::am_alloc(size*sizeof(size_t), ctx->_acc, am_flags);
break;
default:
hip_status = hipErrorUnknown;
break;
}
if (size && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
} else {
hc::am_memtracker_update(*ptr, ctx->_device_index, 0);
{
LockedAccessor_DeviceCrit_t crit(ctx->criticalData());
if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved:
hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
if (hsa_status != HSA_STATUS_SUCCESS) {
hip_status = hipErrorMemoryAllocation;
}
}
}
}
switch(desc->f) {
case hipChannelFormatKindSigned:
*ptr = hc::am_alloc(size*sizeof(int), device->_acc, am_flags);
break;
case hipChannelFormatKindUnsigned:
*ptr = hc::am_alloc(size*sizeof(unsigned int), device->_acc, am_flags);
break;
case hipChannelFormatKindFloat:
*ptr = hc::am_alloc(size*sizeof(float), device->_acc, am_flags);
break;
case hipChannelFormatKindNone:
*ptr = hc::am_alloc(size*sizeof(size_t), device->_acc, am_flags);
break;
default:
hip_status = hipErrorUnknown;
break;
}
if (size && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
} else {
hc::am_memtracker_update(*ptr, device->_device_index, 0);
{
LockedAccessor_CtxCrit_t crit(ctx->criticalData());
if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved:
hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
if (hsa_status != HSA_STATUS_SUCCESS) {
hip_status = hipErrorMemoryAllocation;
}
}
}
}
} else {
hip_status = hipErrorMemoryAllocation;
@@ -314,24 +322,24 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr)
{
HIP_INIT_API(flagsPtr, hostPtr);
hipError_t hip_status = hipSuccess;
hipError_t hip_status = hipSuccess;
hc::accelerator acc;
hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
if(status == AM_SUCCESS){
*flagsPtr = amPointerInfo._appAllocationFlags;
if(*flagsPtr == 0){
hip_status = hipErrorInvalidValue;
}
else{
hip_status = hipSuccess;
}
tprintf(DB_MEM, " %s: host ptr=%p\n", __func__, hostPtr);
}else{
hip_status = hipErrorInvalidValue;
}
return ihipLogStatus(hip_status);
hc::accelerator acc;
hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
if(status == AM_SUCCESS){
*flagsPtr = amPointerInfo._appAllocationFlags;
if(*flagsPtr == 0){
hip_status = hipErrorInvalidValue;
}
else{
hip_status = hipSuccess;
}
tprintf(DB_MEM, " %s: host ptr=%p\n", __func__, hostPtr);
}else{
hip_status = hipErrorInvalidValue;
}
return ihipLogStatus(hip_status);
}
@@ -353,24 +361,25 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags)
if(am_status == AM_SUCCESS){
hip_status = hipErrorHostMemoryAlreadyRegistered;
}else{
} else {
auto ctx = ihipGetTlsDefaultCtx();
if(hostPtr == NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
if(ctx){
if (ctx) {
auto device = ctx->getWriteableDevice();
if(flags == hipHostRegisterDefault || flags == hipHostRegisterPortable || flags == hipHostRegisterMapped){
std::vector<hc::accelerator>vecAcc;
for(int i=0;i<g_deviceCnt;i++){
vecAcc.push_back(ihipGetDevice(i)->_acc);
}
am_status = hc::am_memory_host_lock(ctx->_acc, hostPtr, sizeBytes, &vecAcc[0], vecAcc.size());
am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0], vecAcc.size());
if(am_status == AM_SUCCESS){
hip_status = hipSuccess;
}else{
} else {
hip_status = hipErrorMemoryAllocation;
}
}else{
} else {
hip_status = hipErrorInvalidValue;
}
}
@@ -387,7 +396,8 @@ hipError_t hipHostUnregister(void *hostPtr)
if(hostPtr == NULL){
hip_status = hipErrorInvalidValue;
}else{
am_status_t am_status = hc::am_memory_host_unlock(ctx->_acc, hostPtr);
auto device = ctx->getWriteableDevice();
am_status_t am_status = hc::am_memory_host_unlock(device->_acc, hostPtr);
if(am_status != AM_SUCCESS){
hip_status = hipErrorHostMemoryNotRegistered;
}
@@ -402,17 +412,17 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou
HIP_INIT_API(symbolName, src, count, offset, kind);
#ifdef USE_MEMCPYTOSYMBOL
if(kind != hipMemcpyHostToDevice)
{
return ihipLogStatus(hipErrorInvalidValue);
}
auto ctx = ihipGetTlsDefaultCtx();
if(kind != hipMemcpyHostToDevice)
{
return ihipLogStatus(hipErrorInvalidValue);
}
auto ctx = ihipGetTlsDefaultCtx();
//hsa_signal_t depSignal;
//int depSignalCnt = ctx._default_stream->preCopyCommand(NULL, &depSignal, ihipCommandCopyH2D);
assert(0); // Need to properly synchronize the copy - do something with depSignal if != NULL.
ctx->_acc.memcpy_symbol(symbolName, (void*) src,count, offset);
ctx->_acc.memcpy_symbol(symbolName, (void*) src,count, offset);
#endif
return ihipLogStatus(hipSuccess);
}
@@ -476,110 +486,110 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
// dpitch, spitch, and width in bytes
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
size_t width, size_t height, hipMemcpyKind kind) {
size_t width, size_t height, hipMemcpyKind kind) {
HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind);
HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind);
if(width > dpitch || width > spitch)
return ihipLogStatus(hipErrorUnknown);
if(width > dpitch || width > spitch)
return ihipLogStatus(hipErrorUnknown);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hc::completion_future marker;
hipError_t e = hipSuccess;
hipError_t e = hipSuccess;
try {
for(int i = 0; i < height; ++i) {
stream->locked_copySync((unsigned char*)dst + i*dpitch, (unsigned char*)src + i*spitch, width, kind);
try {
for(int i = 0; i < height; ++i) {
stream->locked_copySync((unsigned char*)dst + i*dpitch, (unsigned char*)src + i*spitch, width, kind);
}
}
catch (ihipException ex) {
e = ex._code;
}
}
catch (ihipException ex) {
e = ex._code;
}
return ihipLogStatus(e);
return ihipLogStatus(e);
}
// wOffset, width, and spitch in bytes
hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
size_t spitch, size_t width, size_t height, hipMemcpyKind kind) {
size_t spitch, size_t width, size_t height, hipMemcpyKind kind) {
HIP_INIT_API(dst, wOffset, hOffset, src, spitch, width, height, kind);
HIP_INIT_API(dst, wOffset, hOffset, src, spitch, width, height, kind);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hc::completion_future marker;
hipError_t e = hipSuccess;
hipError_t e = hipSuccess;
size_t byteSize;
if(dst) {
switch(dst[0].f) {
case hipChannelFormatKindSigned:
byteSize = sizeof(int);
break;
case hipChannelFormatKindUnsigned:
byteSize = sizeof(unsigned int);
break;
case hipChannelFormatKindFloat:
byteSize = sizeof(float);
break;
case hipChannelFormatKindNone:
byteSize = sizeof(size_t);
break;
default:
byteSize = 0;
break;
size_t byteSize;
if(dst) {
switch(dst[0].f) {
case hipChannelFormatKindSigned:
byteSize = sizeof(int);
break;
case hipChannelFormatKindUnsigned:
byteSize = sizeof(unsigned int);
break;
case hipChannelFormatKindFloat:
byteSize = sizeof(float);
break;
case hipChannelFormatKindNone:
byteSize = sizeof(size_t);
break;
default:
byteSize = 0;
break;
}
} else {
return ihipLogStatus(hipErrorUnknown);
}
} else {
return ihipLogStatus(hipErrorUnknown);
}
if((wOffset + width > (dst->width * byteSize)) || width > spitch) {
return ihipLogStatus(hipErrorUnknown);
}
size_t src_w = spitch;
size_t dst_w = (dst->width)*byteSize;
try {
for(int i = 0; i < height; ++i) {
stream->locked_copySync((unsigned char*)dst->data + i*dst_w, (unsigned char*)src + i*src_w, width, kind);
if((wOffset + width > (dst->width * byteSize)) || width > spitch) {
return ihipLogStatus(hipErrorUnknown);
}
}
catch (ihipException ex) {
e = ex._code;
}
return ihipLogStatus(e);
size_t src_w = spitch;
size_t dst_w = (dst->width)*byteSize;
try {
for(int i = 0; i < height; ++i) {
stream->locked_copySync((unsigned char*)dst->data + i*dst_w, (unsigned char*)src + i*src_w, width, kind);
}
}
catch (ihipException ex) {
e = ex._code;
}
return ihipLogStatus(e);
}
hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset,
const void* src, size_t count, hipMemcpyKind kind) {
const void* src, size_t count, hipMemcpyKind kind) {
HIP_INIT_API(dst, wOffset, hOffset, src, count, kind);
HIP_INIT_API(dst, wOffset, hOffset, src, count, kind);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hc::completion_future marker;
hipError_t e = hipSuccess;
hipError_t e = hipSuccess;
try {
stream->locked_copySync((char *)dst->data + wOffset, src, count, kind);
}
catch (ihipException ex) {
e = ex._code;
}
try {
stream->locked_copySync((char *)dst->data + wOffset, src, count, kind);
}
catch (ihipException ex) {
e = ex._code;
}
return ihipLogStatus(e);
return ihipLogStatus(e);
}
// TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
/** @return #hipErrorInvalidValue
*/
*/
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream )
{
HIP_INIT_API(dst, value, sizeBytes, stream);
@@ -694,17 +704,18 @@ hipError_t hipMemGetInfo (size_t *free, size_t *total)
ihipCtx_t * ctx = ihipGetTlsDefaultCtx();
if (ctx) {
auto device = ctx->getWriteableDevice();
if (total) {
*total = ctx->_props.totalGlobalMem;
*total = device->_props.totalGlobalMem;
}
if (free) {
// TODO - replace with kernel-level for reporting free memory:
size_t deviceMemSize, hostMemSize, userMemSize;
hc::am_memtracker_sizeinfo(ctx->_acc, &deviceMemSize, &hostMemSize, &userMemSize);
hc::am_memtracker_sizeinfo(device->_acc, &deviceMemSize, &hostMemSize, &userMemSize);
printf ("deviceMemSize=%zu\n", deviceMemSize);
*free = ctx->_props.totalGlobalMem - deviceMemSize;
*free = device->_props.totalGlobalMem - deviceMemSize;
}
} else {
@@ -722,7 +733,7 @@ hipError_t hipFree(void* ptr)
hipError_t hipStatus = hipErrorInvalidDevicePointer;
// Synchronize to ensure all work has finished.
// Synchronize to ensure all work has finished.
ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
if (ptr) {
@@ -748,7 +759,7 @@ hipError_t hipHostFree(void* ptr)
{
HIP_INIT_API(ptr);
// Synchronize to ensure all work has finished.
// Synchronize to ensure all work has finished.
ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
@@ -780,26 +791,26 @@ hipError_t hipFreeHost(void* ptr)
hipError_t hipFreeArray(hipArray* array)
{
HIP_INIT_API(array);
HIP_INIT_API(array);
hipError_t hipStatus = hipErrorInvalidDevicePointer;
hipError_t hipStatus = hipErrorInvalidDevicePointer;
// Synchronize to ensure all work has finished.
ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
// Synchronize to ensure all work has finished.
ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
if(array->data) {
hc::accelerator acc;
hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, array->data);
if(status == AM_SUCCESS){
if(amPointerInfo._hostPointer == NULL){
hc::am_free(array->data);
hipStatus = hipSuccess;
}
if(array->data) {
hc::accelerator acc;
hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, array->data);
if(status == AM_SUCCESS){
if(amPointerInfo._hostPointer == NULL){
hc::am_free(array->data);
hipStatus = hipSuccess;
}
}
}
}
return ihipLogStatus(hipStatus);
return ihipLogStatus(hipStatus);
}
// Stubs of threadfence operations
projects/hip/src/hip_peer.cpp
+78 -33
View File
@@ -23,25 +23,32 @@ THE SOFTWARE.
#include "hcc_detail/hip_hcc.h"
#include "hcc_detail/trace_helper.h"
// Peer access functions.
// There are two flavors:
// - one where contexts are specified with hipCtx_t type.
// - one where contexts are specified with integer deviceIds, that are mapped to the primary context for that device.
// The implementation contains a set of internal ihip* functions which operate on contexts. Then the
// public APIs are thin wrappers which call into this internal implementations.
// TODO - actually not yet - currently the integer deviceId flavors just call the context APIs. need to fix.
/**
* HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P
*/
//---
hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId)
hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t *thisCtx, hipCtx_t *peerCtx)
{
HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId);
HIP_INIT_API(canAccessPeer, thisCtx, peerCtx);
hipError_t err = hipSuccess;
auto thisDevice = ihipGetDevice(deviceId);
auto peerDevice = ihipGetDevice(peerDeviceId);
if ((thisDevice != NULL) && (peerDevice != NULL)) {
if (deviceId == peerDeviceId) {
if ((thisCtx != NULL) && (peerCtx != NULL)) {
if (thisCtx == peerCtx) {
*canAccessPeer = 0;
} else {
#if USE_PEER_TO_PEER>=2
*canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc);
*canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc);
#else
*canAccessPeer = 0;
#endif
@@ -60,33 +67,32 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDe
//---
// Disable visibility of this device into memory allocated on peer device.
// Remove this device from peer device peerlist.
hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
hipError_t hipDeviceDisablePeerAccess (hipCtx_t *peerCtx)
{
HIP_INIT_API(peerDeviceId);
HIP_INIT_API(peerCtx);
hipError_t err = hipSuccess;
auto thisDevice = ihipGetTlsDefaultCtx();
auto peerDevice = ihipGetDevice(peerDeviceId);
if ((thisDevice != NULL) && (peerDevice != NULL)) {
auto thisCtx = ihipGetTlsDefaultCtx();
if ((thisCtx != NULL) && (peerCtx != NULL)) {
#if USE_PEER_TO_PEER>=2
// Return true if thisDevice can access peerDevice's memory:
bool canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc);
// Return true if thisCtx can access peerCtx's memory:
bool canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc);
#else
bool canAccessPeer = 0;
#endif
if (! canAccessPeer) {
err = hipErrorInvalidDevice; // P2P not allowed between these devices.
} else if (thisDevice == peerDevice) {
} else if (thisCtx == peerCtx) {
err = hipErrorInvalidDevice; // Can't disable peer access to self.
} else {
LockedAccessor_DeviceCrit_t peerCrit(peerDevice->criticalData());
bool changed = peerCrit->removePeer(thisDevice);
LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData());
bool changed = peerCrit->removePeer(thisCtx);
if (changed) {
#if USE_PEER_TO_PEER>=3
// Update the peers for all memory already saved in the tracker:
am_memtracker_update_peers(peerDevice->_acc, peerCrit->peerCnt(), peerCrit->peerAgents());
am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents());
#endif
} else {
err = hipErrorPeerAccessNotEnabled; // never enabled P2P access.
@@ -103,24 +109,23 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
//---
// Allow the current device to see all memory allocated on peerDevice.
// This should add this device to the peer-device peer list.
hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
hipError_t hipDeviceEnablePeerAccess (hipCtx_t *peerCtx, unsigned int flags)
{
HIP_INIT_API(peerDeviceId, flags);
HIP_INIT_API(peerCtx, flags);
hipError_t err = hipSuccess;
if (flags != 0) {
err = hipErrorInvalidValue;
} else {
auto thisDevice = ihipGetTlsDefaultCtx();
auto peerDevice = ihipGetDevice(peerDeviceId);
if (thisDevice == peerDevice) {
auto thisCtx = ihipGetTlsDefaultCtx();
if (thisCtx == peerCtx) {
err = hipErrorInvalidDevice; // Can't enable peer access to self.
} else if ((thisDevice != NULL) && (peerDevice != NULL)) {
LockedAccessor_DeviceCrit_t peerCrit(peerDevice->criticalData());
bool isNewPeer = peerCrit->addPeer(thisDevice);
} else if ((thisCtx != NULL) && (peerCtx != NULL)) {
LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData());
bool isNewPeer = peerCrit->addPeer(thisCtx);
if (isNewPeer) {
#if USE_PEER_TO_PEER>=3
am_memtracker_update_peers(peerDevice->_acc, peerCrit->peerCnt(), peerCrit->peerAgents());
am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents());
#endif
} else {
err = hipErrorPeerAccessAlreadyEnabled;
@@ -135,20 +140,17 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
//---
hipError_t hipMemcpyPeer (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes)
hipError_t hipMemcpyPeer (void* dst, hipCtx_t *dstCtx, const void* src, hipCtx_t *srcCtx, size_t sizeBytes)
{
HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes);
HIP_INIT_API(dst, dstCtx, src, srcCtx, sizeBytes);
// HCC has a unified memory architecture so device specifiers are not required.
return hipMemcpy(dst, src, sizeBytes, hipMemcpyDefault);
};
/**
* This function uses a synchronous copy
*/
//---
hipError_t hipMemcpyPeerAsync (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream)
hipError_t hipMemcpyPeerAsync (void* dst, hipCtx_t *dstDevice, const void* src, hipCtx_t *srcDevice, size_t sizeBytes, hipStream_t stream)
{
HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream);
// HCC has a unified memory architecture so device specifiers are not required.
@@ -156,6 +158,49 @@ hipError_t hipMemcpyPeerAsync (void* dst, int dstDevice, const void* src, int
};
//=============================================================================
// These are the flavors that accept integer deviceIDs.
// Implementations map these to primary contexts and call the internal functions above.
//=============================================================================
hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId)
{
HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId);
return hipDeviceCanAccessPeer(canAccessPeer, ihipGetPrimaryCtx(deviceId), ihipGetPrimaryCtx(peerDeviceId));
}
hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
{
HIP_INIT_API(peerDeviceId);
return hipDeviceDisablePeerAccess(ihipGetPrimaryCtx(peerDeviceId));
}
hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
{
HIP_INIT_API(peerDeviceId, flags);
return hipDeviceEnablePeerAccess(ihipGetPrimaryCtx(peerDeviceId), flags);
}
hipError_t hipMemcpyPeer (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes)
{
HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes);
return hipMemcpyPeer(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes);
}
hipError_t hipMemcpyPeerAsync (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream)
{
HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream);
return hipMemcpyPeerAsync(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes, stream);
}
/**
* @return #hipSuccess
*/
projects/hip/src/hip_stream.cpp
+18 -12
View File
@@ -31,22 +31,28 @@ THE SOFTWARE.
hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags)
{
ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
hc::accelerator acc = ctx->_acc;
// TODO - se try-catch loop to detect memory exception?
//
//
//Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete:
//This matches CUDA stream behavior:
hipError_t e = hipSuccess;
auto istream = new ihipStream_t(ctx, acc.create_view(), flags);
if (ctx) {
hc::accelerator acc = ctx->getWriteableDevice()->_acc;
ctx->locked_addStream(istream);
// TODO - se try-catch loop to detect memory exception?
//
//Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete:
//This matches CUDA stream behavior:
*stream = istream;
tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream);
auto istream = new ihipStream_t(ctx, acc.create_view(), flags);
return hipSuccess;
ctx->locked_addStream(istream);
*stream = istream;
tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream);
} else {
e = hipErrorInvalidDevice;
}
return ihipLogStatus(e);
}
@@ -129,7 +135,7 @@ hipError_t hipStreamDestroy(hipStream_t stream)
e = hipSuccess;
}
ihipCtx_t *ctx = stream->getDevice();
ihipCtx_t *ctx = stream->getCtx();
if (ctx) {
ctx->locked_removeStream(stream);