AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Split ihipCtx_t into ihipCtx_t and ihipDevice_t .

Просмотр исходного кода 
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/clr commit: d09b19bb6c]
Этот коммит содержится в:
Ben Sander
2016-08-08 11:55:57 -05:00
родитель 4c7160a134
Коммит 3ee6e7f8ab
7 изменённых файлов: 687 добавлений и 531 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/clr/hipamd/include/hcc_detail/hip_hcc.h
+102 -68
Просмотреть файл
@@ -69,6 +69,7 @@ extern int HIP_DISABLE_HW_COPY_DEP;
extern thread_local int tls_defaultDevice;
extern thread_local hipError_t tls_lastHipError;
class ihipStream_t;
class ihipDevice_t;
class ihipCtx_t;
@@ -90,7 +91,7 @@ class ihipCtx_t;
#define STREAM_THREAD_SAFE 1
#define DEVICE_THREAD_SAFE 1
#define CTX_THREAD_SAFE 1
// If FORCE_COPY_DEP=1 , HIP runtime will add
// synchronization for copy commands in the same stream, regardless of command type.
@@ -227,7 +228,6 @@ public:
extern "C" {
#endif
typedef class ihipStream_t* hipStream_t;
#ifdef __cplusplus
}
@@ -287,10 +287,11 @@ typedef std::mutex StreamMutex;
typedef FakeMutex StreamMutex;
#endif
#if DEVICE_THREAD_SAFE
typedef std::mutex DeviceMutex;
// Pair Device and Ctx together, these could also be toggled separately if desired.
#if CTX_THREAD_SAFE
typedef std::mutex CtxMutex;
#else
typedef FakeMutex DeviceMutex;
typedef FakeMutex CtxMutex;
#warning "Device thread-safe disabled"
#endif
@@ -429,7 +430,8 @@ typedef uint64_t SeqNum_t ;
//-- Non-racy accessors:
// These functions access fields set at initialization time and are non-racy (so do not acquire mutex)
ihipCtx_t * getDevice() const;
const ihipDevice_t * getDevice() const;
ihipCtx_t * getCtx() const;
public:
@@ -440,21 +442,22 @@ public:
unsigned _flags;
private:
// Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t
ihipStreamCritical_t _criticalData;
private:
void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal);
void waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal);
void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal);
void waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal);
// The unsigned return is hipMemcpyKind
unsigned resolveMemcpyDirection(bool srcTracked, bool dstTracked, bool srcInDeviceMem, bool dstInDeviceMem);
void setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent);
void setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent);
private: // Data
// Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t
ihipStreamCritical_t _criticalData;
ihipCtx_t *_ctx; // parent context that owns this stream.
// Friends:
friend std::ostream& operator<<(std::ostream& os, const ihipStream_t& s);
};
@@ -507,117 +510,145 @@ struct ihipEvent_t {
//---
// Data that must be protected with thread-safe access
// All members are private - this class must be accessed through friend LockedAccessor which
// will lock the mutex on construction and unlock on destruction.
//
// MUTEX_TYPE is template argument so can easily convert to FakeMutex for performance or stress testing.
template <class MUTEX_TYPE>
class ihipDeviceCriticalBase_t : LockedBase<MUTEX_TYPE>
//----
// Properties of the HIP device.
// Multiple contexts can point to same device.
class ihipDevice_t
{
public:
ihipDeviceCriticalBase_t() : _stream_id(0), _peerAgents(nullptr) {};
void init(unsigned deviceCnt) {
assert(_peerAgents == nullptr);
ihipDevice_t(unsigned deviceIndex, unsigned deviceCnt, hc::accelerator &acc);
~ihipDevice_t();
// Accessors:
ihipCtx_t *getPrimaryCtx() const { return _primaryCtx; };
public:
unsigned _device_index; // device ID
hc::accelerator _acc;
hsa_agent_t _hsa_agent; // hsa agent handle
//! Number of compute units supported by the device:
unsigned _compute_units;
hipDeviceProp_t _props; // saved device properties.
StagingBuffer *_staging_buffer[2]; // one buffer for each direction.
int isLargeBar;
ihipCtx_t *_primaryCtx;
private:
hipError_t initProperties(hipDeviceProp_t* prop);
};
//=============================================================================
//=============================================================================
//class ihipCtxCriticalBase_t
template <typename MUTEX_TYPE>
class ihipCtxCriticalBase_t : LockedBase<MUTEX_TYPE>
{
public:
ihipCtxCriticalBase_t(unsigned deviceCnt) :
_peerCnt(0)
{
_peerAgents = new hsa_agent_t[deviceCnt];
};
~ihipDeviceCriticalBase_t() {
~ihipCtxCriticalBase_t() {
if (_peerAgents != nullptr) {
delete _peerAgents;
_peerAgents = nullptr;
}
_peerCnt = 0;
}
friend class LockedAccessor<ihipDeviceCriticalBase_t>;
// Streams:
void addStream(ihipStream_t *stream);
std::list<ihipStream_t*> &streams() { return _streams; };
const std::list<ihipStream_t*> &const_streams() const { return _streams; };
// "Allocate" a stream ID:
ihipStream_t::SeqNum_t incStreamId() { return _stream_id++; };
// Peer Accessor classes:
bool isPeer(const ihipCtx_t *peer); // returns Trus if peer has access to memory physically located on this device.
bool addPeer(ihipCtx_t *peer);
bool removePeer(ihipCtx_t *peer);
void resetPeers(ihipCtx_t *thisDevice);
void addStream(ihipStream_t *stream);
uint32_t peerCnt() const { return _peerCnt; };
hsa_agent_t *peerAgents() const { return _peerAgents; };
private:
//std::list< std::shared_ptr<ihipStream_t> > _streams; // streams associated with this device. TODO - convert to shared_ptr.
std::list< ihipStream_t* > _streams; // streams associated with this device.
ihipStream_t::SeqNum_t _stream_id;
friend class LockedAccessor<ihipCtxCriticalBase_t>;
private:
//--- Stream Tracker:
std::list< ihipStream_t* > _streams; // streams associated with this device.
//--- Peer Tracker:
// These reflect the currently Enabled set of peers for this GPU:
// Enabled peers have permissions to access the memory physically allocated on this device.
std::list<ihipCtx_t*> _peers; // list of enabled peer devices.
std::list<ihipCtx_t*> _peers; // list of enabled peer devices.
uint32_t _peerCnt; // number of enabled peers
hsa_agent_t *_peerAgents; // efficient packed array of enabled agents (to use for allocations.)
private:
void recomputePeerAgents();
};
// Note Mutex selected based on DeviceMutex
typedef ihipDeviceCriticalBase_t<DeviceMutex> ihipDeviceCritical_t;
// Note Mutex type Real/Fake selected based on CtxMutex
typedef ihipCtxCriticalBase_t<CtxMutex> ihipCtxCritical_t;
// This type is used by functions that need access to the critical device structures.
typedef LockedAccessor<ihipDeviceCritical_t> LockedAccessor_DeviceCrit_t;
typedef LockedAccessor<ihipCtxCritical_t> LockedAccessor_CtxCrit_t;
//=============================================================================
//-------------------------------------------------------------------------------------------------
// Functions which read or write the critical data are named locked_.
// ihipCtx_t does not use recursive locks so the ihip implementation must avoid calling a locked_ function from within a locked_ function.
// External functions which call several locked_ functions will acquire and release the lock for each function. if this occurs in
// performance-sensitive code we may want to refactor by adding non-locked functions and creating a new locked_ member function to call them all.
//=============================================================================
//class ihipCtx_t:
// A HIP CTX (context) points at one of the existing devices and contains the streams,
// peer-to-peer mappings, creation flags. Multiple contexts can point to the same
// device.
//
class ihipCtx_t
{
public: // Functions:
ihipCtx_t() {}; // note: calls constructor for _criticalData
void init(unsigned device_index, unsigned deviceCnt, hc::accelerator &acc, unsigned flags);
ihipCtx_t(const ihipDevice_t *device, unsigned deviceCnt, unsigned flags); // note: calls constructor for _criticalData
~ihipCtx_t();
// Functions which read or write the critical data are named locked_.
// ihipCtx_t does not use recursive locks so the ihip implementation must avoid calling a locked_ function from within a locked_ function.
// External functions which call several locked_ functions will acquire and release the lock for each function. if this occurs in
// performance-sensitive code we may want to refactor by adding non-locked functions and creating a new locked_ member function to call them all.
void locked_addStream(ihipStream_t *s);
void locked_removeStream(ihipStream_t *s);
void locked_reset();
void locked_waitAllStreams();
void locked_syncDefaultStream(bool waitOnSelf);
ihipDeviceCritical_t &criticalData() { return _criticalData; }; // TODO, move private. Fix P2P.
ihipCtxCritical_t &criticalData() { return _criticalData; }; // TODO, move private. Fix P2P.
public: // Data, set at initialization:
unsigned _device_index; // device ID
const ihipDevice_t *getDevice() const { return _device; };
hipDeviceProp_t _props; // saved device properties.
hc::accelerator _acc;
hsa_agent_t _hsa_agent; // hsa agent handle
// TODO - review uses of getWriteableDevice(), can these be converted to getDevice()
ihipDevice_t *getWriteableDevice() const { return const_cast<ihipDevice_t*> (_device); };
public: // Data
// The NULL stream is used if no other stream is specified.
// Default stream has special synchronization properties with other streams.
ihipStream_t *_default_stream;
unsigned _compute_units;
StagingBuffer *_staging_buffer[2]; // one buffer for each direction.
int isLargeBar;
unsigned _device_flags;
// Flags specified when the context is created:
unsigned _ctxFlags;
private:
hipError_t getProperties(hipDeviceProp_t* prop);
const ihipDevice_t *_device;
private: // Critical data, protected with locked access:
// Members of _protected data MUST be accessed through the LockedAccessor.
// Search for LockedAccessor<ihipDeviceCritical_t> for examples; do not access _criticalData directly.
ihipDeviceCritical_t _criticalData;
// Search for LockedAccessor<ihipCtxCritical_t> for examples; do not access _criticalData directly.
ihipCtxCritical_t _criticalData;
};
@@ -633,8 +664,11 @@ extern hsa_agent_t g_cpu_agent ; // the CPU agent.
// Extern functions:
extern void ihipInit();
extern const char *ihipErrorString(hipError_t);
extern ihipCtx_t *ihipGetTlsDefaultCtx();
extern ihipCtx_t *ihipGetDevice(int);
extern ihipCtx_t *ihipGetTlsDefaultCtx();
extern ihipDevice_t *ihipGetDevice(int);
ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex);
extern void ihipSetTs(hipEvent_t e);
template<typename T>
projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h
+3 -1
Просмотреть файл
@@ -43,6 +43,7 @@ THE SOFTWARE.
extern "C" {
#endif
typedef struct ihipCtx_t hipCtx_t;
typedef struct ihipStream_t *hipStream_t;
typedef struct hipEvent_t {
struct ihipEvent_t *_handle;
@@ -417,7 +418,6 @@ const char *hipGetErrorString(hipError_t hip_error);
*
* even if the handle goes out-of-scope. To release the memory used by the stream, applicaiton must call hipStreamDestroy.
* Flags controls behavior of the stream. See #hipStreamDefault, #hipStreamNonBlocking.
* @error hipStream_t are under development - with current HIP use the NULL stream.
*/
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags);
@@ -437,6 +437,8 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags);
*
* @see hipStreamDestroy
*
* @return
*
*/
hipError_t hipStreamCreate(hipStream_t *stream);
projects/clr/hipamd/src/hip_device.cpp
+8 -5
Просмотреть файл
@@ -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/clr/hipamd/src/hip_hcc.cpp
+253 -198
Просмотреть файл
@@ -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/clr/hipamd/src/hip_memory.cpp
+225 -214
Просмотреть файл
@@ -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/clr/hipamd/src/hip_peer.cpp
+78 -33
Просмотреть файл
@@ -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/clr/hipamd/src/hip_stream.cpp
+18 -12
Просмотреть файл
@@ -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);