From cfdacab32f2b637868b90f8391e670cef15d176d Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 8 Aug 2016 11:55:57 -0500 Subject: [PATCH] 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 --- include/hcc_detail/hip_hcc.h | 170 ++++++---- include/hcc_detail/hip_runtime_api.h | 4 +- src/hip_device.cpp | 13 +- src/hip_hcc.cpp | 451 +++++++++++++++------------ src/hip_memory.cpp | 439 +++++++++++++------------- src/hip_peer.cpp | 111 +++++-- src/hip_stream.cpp | 30 +- 7 files changed, 687 insertions(+), 531 deletions(-) diff --git a/include/hcc_detail/hip_hcc.h b/include/hcc_detail/hip_hcc.h index 8202b279b6..fb8e11aa14 100644 --- a/include/hcc_detail/hip_hcc.h +++ b/include/hcc_detail/hip_hcc.h @@ -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 ihipDeviceCriticalBase_t : LockedBase +//---- +// 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 +class ihipCtxCriticalBase_t : LockedBase +{ +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; + // Streams: + void addStream(ihipStream_t *stream); std::list &streams() { return _streams; }; const std::list &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 > _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; +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 _peers; // list of enabled peer devices. + std::list _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 ihipDeviceCritical_t; +// Note Mutex type Real/Fake selected based on CtxMutex +typedef ihipCtxCriticalBase_t ihipCtxCritical_t; // This type is used by functions that need access to the critical device structures. -typedef LockedAccessor LockedAccessor_DeviceCrit_t; +typedef LockedAccessor 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 (_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 for examples; do not access _criticalData directly. - ihipDeviceCritical_t _criticalData; + // Search for LockedAccessor 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 diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 1ee7d70ef5..8209645fce 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/include/hcc_detail/hip_runtime_api.h @@ -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); diff --git a/src/hip_device.cpp b/src/hip_device.cpp index bc8879a120..dff736bb94 100644 --- a/src/hip_device.cpp +++ b/src/hip_device.cpp @@ -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); -} +}; diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 30d0e301e6..00b13143f7 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -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 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::recomputePeerAgents() -{ - _peerCnt = 0; - std::for_each (_peers.begin(), _peers.end(), [this](ihipCtx_t* ctx) { - _peerAgents[_peerCnt++] = ctx->_hsa_agent; - }); -} +//============================================================================= -template<> -bool ihipDeviceCriticalBase_t::isPeer(const ihipCtx_t *peer) -{ - auto match = std::find(_peers.begin(), _peers.end(), peer); - return (match != std::end(_peers)); -} - - -template<> -bool ihipDeviceCriticalBase_t::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::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::resetPeers(ihipCtx_t *thisDevice) -{ - _peers.clear(); - _peerCnt = 0; - addPeer(thisDevice); // peer-list always contains self agent. -} - - -template<> -void ihipDeviceCriticalBase_t::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::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::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::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::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::resetPeers(ihipCtx_t *thisDevice) +{ + _peers.clear(); + _peerCnt = 0; + addPeer(thisDevice); // peer-list always contains self agent. +} + + +template<> +void ihipCtxCriticalBase_t::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 (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 (-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(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(data))++; + if (device_type == HSA_DEVICE_TYPE_CPU) { + (*static_cast(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 (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(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_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); diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 91a27ce08b..7146bf76a3 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -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::vectorvecAcc; for(int i=0;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 diff --git a/src/hip_peer.cpp b/src/hip_peer.cpp index b7475d1d38..f5f9287baa 100644 --- a/src/hip_peer.cpp +++ b/src/hip_peer.cpp @@ -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 */ diff --git a/src/hip_stream.cpp b/src/hip_stream.cpp index 2d8efdffb9..a204e2f79a 100644 --- a/src/hip_stream.cpp +++ b/src/hip_stream.cpp @@ -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);