From b9a0abe32eb08963b14afed7e4fd3cf7cda0f18d Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 19 Sep 2016 17:49:25 -0500 Subject: [PATCH] Cleanup. Remove cfs, ihipSignal_t, staging buffer calls. Change-Id: I8bb67c484e3a65be06a03665f059217930da2bed [ROCm/hip commit: e0ce1d395412bea66ee23962de11e0f8e994fa50] --- projects/hip/include/hcc_detail/hip_hcc.h | 48 +----- projects/hip/src/hip_hcc.cpp | 173 +--------------------- 2 files changed, 7 insertions(+), 214 deletions(-) diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h index 3d1b43994e..39c04a1afc 100644 --- a/projects/hip/include/hcc_detail/hip_hcc.h +++ b/projects/hip/include/hcc_detail/hip_hcc.h @@ -227,40 +227,6 @@ extern "C" { const hipStream_t hipStreamNull = 0x0; -enum ihipCommand_t { - ihipCommandCopyH2H, - ihipCommandCopyH2D, - ihipCommandCopyD2H, - ihipCommandCopyD2D, - ihipCommandCopyP2P, - ihipCommandKernel, -}; - -static const char* ihipCommandName[] = { - "CopyH2H", "CopyH2D", "CopyD2H", "CopyD2D", "CopyP2P", "Kernel" -}; - - - -typedef uint64_t SIGSEQNUM; - -//--- -// Small wrapper around signals. -// Designed to be used from stream. -// TODO-someday refactor this class so it can be stored in a vector<> -// we already store the index here so we can use for garbage collection. -struct ihipSignal_t { - hsa_signal_t _hsaSignal; // hsa signal handle - int _index; // Index in pool, used for garbage collection. - SIGSEQNUM _sigId; // unique sequentially increasing ID. - - ihipSignal_t(); - ~ihipSignal_t(); - - void release(); -}; - - // Used to remove lock, for performance or stimulating bugs. class FakeMutex { @@ -384,13 +350,8 @@ public: public: // TODO - remove _kernelCnt mechanism: - uint32_t _kernelCnt; // Count of inflight kernels in this stream. Reset at ::wait(). - hc::accelerator_view _av; - - std::vector _cfs; - }; @@ -426,8 +387,6 @@ typedef uint64_t SeqNum_t ; void locked_waitEvent(hipEvent_t event); void locked_recordEvent(hipEvent_t event); - void addCFtoStream(LockedAccessor_StreamCrit_t &crit, hc::completion_future* cf); - void waitOnAllCFs(LockedAccessor_StreamCrit_t &crit); //--- @@ -456,17 +415,14 @@ public: private: - 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); private: // Data - // Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t + // Critical Data - MUST be accessed through LockedAccessor_StreamCrit_t ihipStreamCritical_t _criticalData; ihipCtx_t *_ctx; // parent context that owns this stream. @@ -525,7 +481,7 @@ public: unsigned _computeUnits; hipDeviceProp_t _props; // saved device properties. - UnpinnedCopyEngine *_stagingBuffer[2]; // one buffer for each direction. + // TODO - report this through device properties, base on HCC API call. int _isLargeBar; ihipCtx_t *_primaryCtx; diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 36146868db..c3624c5fe3 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -67,21 +67,8 @@ int HIP_TRACE_API= 0; std::string HIP_TRACE_API_COLOR("green"); int HIP_ATP_MARKER= 0; int HIP_DB= 0; -int HIP_STAGING_SIZE = 64; /* size of staging buffers, in KB */ -static const int HIP_STAGING_BUFFERS = 2; -int HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING = 0; -int HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE = 0; -int HIP_D2H_MEM_TRANSFER_THRESHOLD = 0; -int HIP_STREAM_SIGNALS = 32; /* number of signals to allocate at stream creation */ int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU identifiers */ - - -//--- -// Chicken bits for disabling functionality to work around potential issues: -int HIP_DISABLE_HW_KERNEL_DEP = 0; - - - +int HIP_NUM_KERNELS_INFLIGHT = 128; std::once_flag hip_initialized; @@ -96,9 +83,6 @@ 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_; //================================================================================================= // Thread-local storage: @@ -166,29 +150,6 @@ hipError_t ihipSynchronize(void) return (hipSuccess); } -//================================================================================================= -// ihipSignal_t: -//================================================================================================= -// -//--- -ihipSignal_t::ihipSignal_t() : _sigId(0) -{ - if (hsa_signal_create(0/*value*/, 0, NULL, &_hsaSignal) != HSA_STATUS_SUCCESS) { - throw ihipException(hipErrorRuntimeMemory); - } - //tprintf (DB_SIGNAL, " allocated hsa_signal=%lu\n", (_hsaSignal.handle)); -} - -//--- -ihipSignal_t::~ihipSignal_t() -{ - tprintf (DB_SIGNAL, " destroy hsa_signal #%lu (#%lu)\n", (_hsaSignal.handle), _sigId); - if (hsa_signal_destroy(_hsaSignal) != HSA_STATUS_SUCCESS) { - throw ihipException(hipErrorRuntimeOther); - } -}; - - //================================================================================================= // ihipStream_t: @@ -210,47 +171,18 @@ ihipStream_t::~ihipStream_t() } -//--- -void ihipStream_t::waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal) -{ - SIGSEQNUM sigNum = signal->_sigId; - tprintf(DB_SYNC, "waitCopy signal:#%lu\n", sigNum); - - hsa_signal_wait_acquire(signal->_hsaSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); - - - tprintf(DB_SIGNAL, "waitCopy reclaim signal #%lu\n", sigNum); -} - //Wait for all kernel and data copy commands in this stream to complete. //This signature should be used in routines that already have locked the stream mutex void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty) { if (! assertQueueEmpty) { tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this); -// crit->_av.wait(); - waitOnAllCFs(crit); + crit->_av.wait(); } crit->_kernelCnt = 0; } -void ihipStream_t::addCFtoStream(LockedAccessor_StreamCrit_t &crit, hc::completion_future *cf) -{ - crit->_cfs.push_back(cf); -} - -void ihipStream_t::waitOnAllCFs(LockedAccessor_StreamCrit_t &crit) -{ - for(uint32_t i=0;i_cfs.size();i++){ - if(crit->_cfs[i] != NULL){ - crit->_cfs[i]->wait(); - delete crit->_cfs[i]; - } - } - crit->_cfs.clear(); -} - //--- //Wait for all kernel and data copy commands in this stream to complete. void ihipStream_t::locked_wait(bool assertQueueEmpty) @@ -301,41 +233,8 @@ ihipCtx_t * ihipStream_t::getCtx() const }; -#define HIP_NUM_SIGNALS_PER_STREAM 32 -//--- -void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal) -{ - - // Obtain the write index for the command queue - uint64_t index = hsa_queue_load_write_index_relaxed(queue); - const uint32_t queueMask = queue->size - 1; - - // Define the barrier packet to be at the calculated queue index address - hsa_barrier_and_packet_t* barrier = &(((hsa_barrier_and_packet_t*)(queue->base_address))[index&queueMask]); - memset(barrier, 0, sizeof(hsa_barrier_and_packet_t)); - - // setup header - uint16_t header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; - header |= 1 << HSA_PACKET_HEADER_BARRIER; - //header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; - //header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; - barrier->header = header; - - barrier->dep_signal[0].handle = depSignal ? depSignal->_hsaSignal.handle: 0; - - barrier->completion_signal.handle = completionSignal ? completionSignal->_hsaSignal.handle : 0; - - // TODO - check queue overflow, return error: - // Increment write index and ring doorbell to dispatch the kernel - hsa_queue_store_write_index_relaxed(queue, index+1); - hsa_signal_store_relaxed(queue->doorbell_signal, index); -} - - -int HIP_NUM_KERNELS_INFLIGHT = 128; - //-- // Lock the stream to prevent other threads from intervening. LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand() @@ -511,16 +410,6 @@ ihipDevice_t::ihipDevice_t(unsigned deviceId, unsigned deviceCnt, hc::accelerato initProperties(&_props); - _stagingBuffer[0] = new UnpinnedCopyEngine(_hsaAgent,g_cpu_agent, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS, - _isLargeBar, - HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING, - HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE, - HIP_D2H_MEM_TRANSFER_THRESHOLD); - _stagingBuffer[1] = new UnpinnedCopyEngine(_hsaAgent,g_cpu_agent, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS, - _isLargeBar, - HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING, - HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE, - HIP_D2H_MEM_TRANSFER_THRESHOLD); _primaryCtx = new ihipCtx_t(this, deviceCnt, hipDeviceMapHost); } @@ -528,12 +417,8 @@ ihipDevice_t::ihipDevice_t(unsigned deviceId, unsigned deviceCnt, hc::accelerato ihipDevice_t::~ihipDevice_t() { - for (int i=0; i<2; i++) { - if (_stagingBuffer[i]) { - delete _stagingBuffer[i]; - _stagingBuffer[i] = NULL; - } - } + delete _primaryCtx; + _primaryCtx = NULL; } @@ -607,14 +492,6 @@ hsa_status_t GetDevicePool(hsa_amd_memory_pool_t pool, void* data) { return HSA_STATUS_SUCCESS; } -void FindDevicePool() -{ - hsa_status_t err = hsa_iterate_agents(FindGpuDevice, &gpu_agent_); - ErrorCheck(err); - - err = hsa_amd_agent_iterate_memory_pools(gpu_agent_, GetDevicePool, &gpu_pool_); - ErrorCheck(err); -} int checkAccess(hsa_agent_t agent, hsa_amd_memory_pool_t pool) { @@ -767,14 +644,7 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) /* Computemode for HSA Devices is always : cudaComputeModeDefault */ prop->computeMode = 0; - FindDevicePool(); - int access=checkAccess(g_cpu_agent, gpu_pool_); - if (0 != access){ - _isLargeBar= 1; - } else { - _isLargeBar=0; - } - + _isLargeBar = _acc.has_cpu_accessible_am(); // Get Max Threads Per Multiprocessor @@ -1115,14 +985,8 @@ void ihipInit() READ_ENV_I(release, HIP_TRACE_API, 0, "Trace each HIP API call. Print function name and return code to stderr as program executes."); READ_ENV_S(release, HIP_TRACE_API_COLOR, 0, "Color to use for HIP_API. None/Red/Green/Yellow/Blue/Magenta/Cyan/White"); READ_ENV_I(release, HIP_ATP_MARKER, 0, "Add HIP function begin/end to ATP file generated with CodeXL"); - READ_ENV_I(release, HIP_STAGING_SIZE, 0, "Size of each staging buffer (in KB)" ); - READ_ENV_I(release, HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING, 0, "Threshold value for H2D unpinned memory transfer decision between direct copy or staging buffer usage,Under testing."); - READ_ENV_I(release, HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE, 0, "Threshold value for H2D unpinned memory transfer decision between staging buffer usage or pininplace usage .Under testing."); - READ_ENV_I(release, HIP_D2H_MEM_TRANSFER_THRESHOLD, 0, "Threshold value for D2H unpinned memory transfer decision between staging buffer usage or pininplace usage .Under testing."); - READ_ENV_I(release, HIP_STREAM_SIGNALS, 0, "Number of signals to allocate when new stream is created (signal pool will grow on demand)"); READ_ENV_I(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the secquence are visible to HIP applications and they are enumerated in the order of secquence" ); - READ_ENV_I(release, HIP_DISABLE_HW_KERNEL_DEP, 0, "Disable HW dependencies before kernel commands - instead wait for dependency on host. -1 means ignore these dependencies. (debug mode)"); READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Number of kernels per stream "); @@ -1281,7 +1145,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_ auto crit = stream->lockopen_preKernelCommand(); lp->av = &(crit->_av); lp->cf = new hc::completion_future; - stream->addCFtoStream(crit, lp->cf); ihipPrintKernelLaunch(kernelNameStr, lp, stream); return (stream); @@ -1304,7 +1167,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri auto crit = stream->lockopen_preKernelCommand(); lp->av = &(crit->_av); lp->cf = new hc::completion_future; - stream->addCFtoStream(crit, lp->cf); ihipPrintKernelLaunch(kernelNameStr, lp, stream); return (stream); } @@ -1326,7 +1188,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri auto crit = stream->lockopen_preKernelCommand(); lp->av = &(crit->_av); lp->cf = new hc::completion_future; - stream->addCFtoStream(crit, lp->cf); ihipPrintKernelLaunch(kernelNameStr, lp, stream); return (stream); } @@ -1349,7 +1210,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g lp->av = &(crit->_av); lp->cf = new hc::completion_future; // TODO, is this necessary? - stream->addCFtoStream(crit, lp->cf); ihipPrintKernelLaunch(kernelNameStr, lp, stream); return (stream); @@ -1498,29 +1358,6 @@ unsigned ihipStream_t::resolveMemcpyDirection(bool srcTracked, bool dstTracked, } -// Setup the copyCommandType and the copy agents (for hsa_amd_memory_async_copy) -// srcPhysAcc is the physical location of the src data. For many copies this is the -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. - const ihipDevice_t *streamDevice = this->getDevice(); - hsa_agent_t streamAgent = streamDevice->_hsaAgent; - - // ROCR runtime logic is : - // - If both src and dst are cpu agent, launch thread and memcpy. We want to avoid this. - // - If either/both src or dst is a gpu agent, use the first gpu agent’s DMA engine to perform the copy. - - switch (kind) { - //case hipMemcpyHostToHost : *commandType = ihipCommandCopyH2H; *srcAgent=streamAgent; *dstAgent=streamAgent; break; // TODO - enable me, for async copy use SDMA. - case hipMemcpyHostToHost : *commandType = ihipCommandCopyH2H; *srcAgent=g_cpu_agent; *dstAgent=g_cpu_agent; break; - case hipMemcpyHostToDevice : *commandType = ihipCommandCopyH2D; *srcAgent=g_cpu_agent; *dstAgent=streamAgent; break; - case hipMemcpyDeviceToHost : *commandType = ihipCommandCopyD2H; *srcAgent=streamAgent; *dstAgent=g_cpu_agent; break; - case hipMemcpyDeviceToDevice : *commandType = ihipCommandCopyD2D; *srcAgent=streamAgent; *dstAgent=streamAgent; break; - default: throw ihipException(hipErrorInvalidMemcpyDirection); - }; -} - - void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn) { ihipCtx_t *ctx = this->getCtx();