AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

Remove HIP_MAX_QUEUES (replaced with HCC_MAX_QUEUES)

Sfoglia il codice sorgente
This commit is contained in:
Ben Sander
2017-05-23 23:47:56 -05:00
parent 236ce70e94
commit d0ef9d8462
4 ha cambiato i file con 4 aggiunte e 108 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
src/hip_hcc.cpp
+2 -89
Vedi File
@@ -64,7 +64,6 @@ std::string HIP_LAUNCH_BLOCKING_KERNELS;
std::vector<std::string> g_hipLaunchBlockingKernels;
int HIP_API_BLOCKING = 0;
int HIP_MAX_QUEUES = 0;
int HIP_PRINT_ENV = 0;
int HIP_TRACE_API= 0;
@@ -267,31 +266,6 @@ ihipStream_t::~ihipStream_t()
}
inline void ihipStream_t::ensureHaveQueue(LockedAccessor_StreamCrit_t &streamCrit)
{
if (HIP_MAX_QUEUES && !streamCrit->_hasQueue) {
// To avoid deadlock, we have to release the stream lock before acquiring context lock.
// Else we can get hung if another thread has the context lock is trying to get lock for this stream.
// We lock it again below.
streamCrit->munlock();
// Obtain mutex access to the device critical data, release by destructor
LockedAccessor_CtxCrit_t ctxCrit(this->_ctx->criticalData());
// TODO
auto needyCritPtr = this->_criticalData.mlock();
// Second test to ensure we still need to steal the queue - another thread may have
// snuck in here and already solved the issue.
if (!needyCritPtr->_hasQueue) {
needyCritPtr->_av = this->_ctx->stealActiveQueue(ctxCrit, this);
}
streamCrit->_hasQueue = true;
}
assert(streamCrit->_hasQueue);
}
hc::hcWaitMode ihipStream_t::waitMode() const
{
hc::hcWaitMode waitMode = hc::hcWaitModeActive;
@@ -323,13 +297,9 @@ hc::hcWaitMode ihipStream_t::waitMode() const
//This signature should be used in routines that already have locked the stream mutex
void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit)
{
if (crit->_hasQueue) {
tprintf (DB_SYNC, "%s wait for queue-empty..\n", ToString(this).c_str());
tprintf (DB_SYNC, "%s wait for queue-empty..\n", ToString(this).c_str());
crit->_av.wait(waitMode());
} else {
tprintf (DB_SYNC, "%s wait for queue empty (done since stream has no physical queue).\n", ToString(this).c_str());
}
crit->_av.wait(waitMode());
crit->_kernelCnt = 0;
}
@@ -350,7 +320,6 @@ void ihipStream_t::locked_waitEvent(hipEvent_t event)
{
LockedAccessor_StreamCrit_t crit(_criticalData);
this->ensureHaveQueue(crit);
crit->_av.create_blocking_marker(event->_marker, hc::accelerator_scope);
}
@@ -362,7 +331,6 @@ void ihipStream_t::locked_recordEvent(hipEvent_t event)
// Lock the stream to prevent simultaneous access
LockedAccessor_StreamCrit_t crit(_criticalData);
this->ensureHaveQueue(crit);
#if USE_NO_SCOPE
printf ("create_marker, flags = %x\n", event->_flags);
event->_marker = crit->_av.create_marker((event->_flags & hipEventDisableSystemRelease) ? hc::no_scope : hc::system_scope);
@@ -406,7 +374,6 @@ LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand()
crit->_kernelCnt = 0;
}
this->ensureHaveQueue(crit);
@@ -1001,55 +968,6 @@ std::string ihipCtx_t::toString() const
};
hc::accelerator_view
ihipCtx_t::stealActiveQueue(LockedAccessor_CtxCrit_t &ctxCrit, ihipStream_t *needyStream)
{
// TODO - review handling if queue can't be found.
while (1) {
for (auto iter=ctxCrit->streams().begin(); iter != ctxCrit->streams().end(); iter++) {
if (*iter != needyStream) {
auto victimCritPtr = (*iter)->_criticalData.mtry_lock();
if (victimCritPtr) {
// try-lock succeeded:
if (victimCritPtr->_hasQueue && (victimCritPtr->_kernelCnt == 0)) {
victimCritPtr->_hasQueue = false;
tprintf(DB_SYNC, " stealActiveQueue from victim:%s to needy:%s\n",
ToString(*iter).c_str(), ToString(needyStream).c_str());
hc::accelerator_view av = victimCritPtr->_av;
// TODO - cleanup to remove forced setting to N
uint64_t *p = (uint64_t*)(&victimCritPtr->_av);
*p = 0; // damage the victim av so attempt to use it will fault.
(*iter)->_criticalData.munlock();
return av;
}
(*iter)->_criticalData.munlock();
}
}
}
}
}
hc::accelerator_view
ihipCtx_t::createOrStealQueue(LockedAccessor_CtxCrit_t &ctxCrit)
{
if (HIP_MAX_QUEUES && (ctxCrit->streams().size() >= HIP_MAX_QUEUES)) {
// Steal a queue from an existing stream:
hc::accelerator_view av = this->stealActiveQueue (ctxCrit, nullptr);
return av;
} else {
// Create a new view
return getWriteableDevice()->_acc.create_view();
}
}
//----
@@ -1279,7 +1197,6 @@ void HipReadEnv()
READ_ENV_I(release, HIP_API_BLOCKING, 0, "Make HIP APIs 'host-synchronous', so they block until completed. Impacts hipMemcpyAsync, hipMemsetAsync." );
READ_ENV_I(release, HIP_MAX_QUEUES, 0, "Maximum number of queues that this app will use per-device. Additional streams will share the specified number of queues. 0=no limit.");
READ_ENV_C(release, HIP_DB, 0, "Print debug info. Bitmask (HIP_DB=0xff) or flags separated by '+' (HIP_DB=api+sync+mem+copy)", HIP_DB_callback);
if ((HIP_DB & (1<<DB_API)) && (HIP_TRACE_API == 0)) {
@@ -1971,7 +1888,6 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,
printPointerInfo(DB_COPY, " dst", dst, dstPtrInfo);
printPointerInfo(DB_COPY, " src", src, srcPtrInfo);
this->ensureHaveQueue(crit);
crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? &copyDevice->getDevice()->_acc : nullptr, forceUnpinnedCopy);
}
@@ -2078,7 +1994,6 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
// Perform fast asynchronous copy - we know copyDevice != NULL based on check above
try {
this->ensureHaveQueue(crit);
if (HIP_FORCE_SYNC_COPY) {
crit->_av.copy_ext (src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, &copyDevice->getDevice()->_acc, forceUnpinnedCopy);
@@ -2115,7 +2030,6 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
// Perform slow synchronous copy:
LockedAccessor_StreamCrit_t crit(_criticalData);
this->ensureHaveQueue(crit);
crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? &copyDevice->getDevice()->_acc : nullptr, forceUnpinnedCopy);
}
@@ -2170,7 +2084,6 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc)
//---
// Warning - with HIP_MAX_QUEUES!=0 there is no mechanism to prevent accelerator_view from being re-assigned...
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av)
{
HIP_INIT_API(stream, av);
src/hip_hcc_internal.h
-12
Vedi File
@@ -447,7 +447,6 @@ public:
ihipStreamCriticalBase_t(ihipStream_t *parentStream, hc::accelerator_view av) :
_kernelCnt(0),
_av(av),
_hasQueue(true),
_parent(parentStream)
{
};
@@ -473,11 +472,6 @@ public:
uint32_t _kernelCnt; // Count of inflight kernels in this stream. Reset at ::wait().
hc::accelerator_view _av;
// True if the stream has an allocated queue (accelerato_view) for its use:
// Always true at ihipStream creation but queue may later be stolen.
// This acts as a valid bit for the _av.
bool _hasQueue;
private:
};
@@ -544,8 +538,6 @@ public:
const ihipDevice_t * getDevice() const;
ihipCtx_t * getCtx() const;
void ensureHaveQueue(LockedAccessor_StreamCrit_t &streamCrit);
public:
//---
//Public member vars - these are set at initialization and never change:
@@ -792,10 +784,6 @@ public: // Functions:
void locked_waitAllStreams();
void locked_syncDefaultStream(bool waitOnSelf, bool syncHost);
// Will allocate a queue and assign it to the needyStream:
hc::accelerator_view stealActiveQueue(LockedAccessor_CtxCrit_t &ctxCrit, ihipStream_t *needyStream);
hc::accelerator_view createOrStealQueue(LockedAccessor_CtxCrit_t &ctxCrit);
ihipCtxCritical_t &criticalData() { return _criticalData; };
const ihipDevice_t *getDevice() const { return _device; };
src/hip_memory.cpp
-3
Vedi File
@@ -948,7 +948,6 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s
if (stream) {
auto crit = stream->lockopen_preKernelCommand();
stream->ensureHaveQueue(crit);
hc::completion_future cf ;
@@ -1000,7 +999,6 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes )
if (stream) {
auto crit = stream->lockopen_preKernelCommand();
stream->ensureHaveQueue(crit);
hc::completion_future cf ;
if ((sizeBytes & 0x3) == 0) {
@@ -1053,7 +1051,6 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeByte
if (stream) {
auto crit = stream->lockopen_preKernelCommand();
stream->ensureHaveQueue(crit);
hc::completion_future cf ;
if ((sizeBytes & 0x3) == 0) {
src/hip_stream.cpp
+2 -4
Vedi File
@@ -49,7 +49,7 @@ hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags)
// Obtain mutex access to the device critical data, release by destructor
LockedAccessor_CtxCrit_t ctxCrit(ctx->criticalData());
auto istream = new ihipStream_t(ctx, ctx->createOrStealQueue(ctxCrit), flags);
auto istream = new ihipStream_t(ctx, acc.create_view(), flags);
ctxCrit->addStream(istream);
*stream = istream;
@@ -129,9 +129,7 @@ hipError_t hipStreamQuery(hipStream_t stream)
{
LockedAccessor_StreamCrit_t crit(stream->_criticalData);
if (crit->_hasQueue) {
pendingOps = crit->_av.get_pending_async_ops();
}
pendingOps = crit->_av.get_pending_async_ops();
}