diff --git a/projects/clr/hipamd/include/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime.h index 7c5a2f2e36..aeed53348e 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_runtime.h @@ -481,7 +481,8 @@ __device__ inline float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); }; #ifdef __HCC_CPP__ -hc::accelerator_view *ihipLaunchKernel(hipStream_t stream); +hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av); +void ihipPostLaunchKernel(hipStream_t stream, hc::completion_future &cf); #if not defined(DISABLE_GRID_LAUNCH) #define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \ @@ -496,12 +497,13 @@ do {\ lp.groupMemBytes = _groupMemBytes;\ hc::completion_future cf;\ lp.cf = &cf; \ - lp.av = (ihipLaunchKernel(_stream)); \ + hipStream_t trueStream = (ihipPreLaunchKernel(_stream, &lp.av)); \ if (HIP_TRACE_API) {\ fprintf(stderr, "hiptrace1: launch '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n", \ #_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\ }\ _kernelName (lp, __VA_ARGS__);\ + ihipPostLaunchKernel(trueStream, cf);\ } while(0) #else @@ -519,12 +521,13 @@ do {\ lp.groupMemBytes = _groupMemBytes;\ hc::completion_future cf;\ lp.cf = &cf; \ - lp.av = (ihipLaunchKernel(_stream)); \ + hipStream_t trueStream = (ihipPreLaunchKernel(_stream, &lp.av)); \ if (HIP_TRACE_API) {\ fprintf(stderr, "hiptrace1: launch '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n", \ #_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\ }\ _kernelName (lp, __VA_ARGS__);\ + ihipPostLaunchKernel(trueStream, cf);\ } while(0) /*end hipLaunchKernel */ #endif diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index f814f99ddf..2d7650b6ed 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -31,6 +31,7 @@ THE SOFTWARE. #include #include #include +#include #include #include @@ -91,28 +92,35 @@ struct ihipDevice_t; enum ihipCommand_t { ihipCommandKernel, - ihipCommandData, + ihipCommandCopyH2D, + ihipCommandCopyD2H, +}; + +const char* ihipCommandName[] = { + "Kernel", "CopyH2D", "CopyD2H" }; // Small wrapper around signals. // Designed to be used from stream. struct ihipSignal_t { - hsa_signal_t _hsa_signal; - int _refCnt; + hsa_signal_t _hsa_signal; // hsa signal handle + int _ref_cnt; // reference count, 0 == signal is free. + uint64_t _seq_id; // unique sequentially increasig ID. - ihipSignal_t() : _refCnt(0) { + ihipSignal_t() : _ref_cnt(0), _seq_id(0) { if (hsa_signal_create(1, 0, NULL, &_hsa_signal) != HSA_STATUS_SUCCESS) { throw; } + tprintf (TRACE_SYNC, " allocated hsa_signal=%lu\n", (_hsa_signal.handle)); } ~ihipSignal_t() { if (hsa_signal_destroy(_hsa_signal) != HSA_STATUS_SUCCESS) { throw; } - // _refCnt should be 0, unless we are shutting down... - _refCnt = 0; + // _ref_cnt should be 0, unless we are shutting down... + _ref_cnt = 0; }; }; @@ -120,23 +128,35 @@ struct ihipSignal_t { // Internal stream structure. class ihipStream_t { public: - unsigned _device_index; - hc::accelerator_view _av; - unsigned _flags; - ihipCommand_t _last_command; - //ihipStream_t() : _av(){ }; ihipStream_t(unsigned device_index, hc::accelerator_view av, unsigned int flags); ~ihipStream_t(); + inline void wait(); + inline ihipDevice_t * getDevice() const; - ihipSignal_t * getSignal() ; + ihipSignal_t * getSignal() ; void releaseSignal(ihipSignal_t *signal) ; + inline bool preKernelCommand(); + inline void postKernelCommand(hc::completion_future &kernel_future); + inline int copyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType); + + //--- + unsigned _device_index; + hc::accelerator_view _av; + unsigned _flags; private: + void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal); + + uint64_t _seq_signal_id; // Monotonically increasing unique signal id. + ihipCommand_t _last_command_type; // type of the last command + ihipSignal_t *_last_copy_signal; // signal of last copy command sent to the stream. Copy can be either H2D or D2H. + hc::completion_future _last_kernel_future; // Completion future of last kernel command sent to GPU. + int _signalCursor; - std::vector _signalPool; + std::deque _signalPool; }; @@ -239,22 +259,20 @@ unsigned g_deviceCnt; //================================================================================================= //--- ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, unsigned int flags) : - _device_index(device_index), _av(av), _flags(flags), _last_command(ihipCommandKernel), + _device_index(device_index), + _av(av), + _flags(flags), + _seq_signal_id(0), + _last_command_type(ihipCommandCopyH2D), + _last_copy_signal (NULL), _signalCursor(0) { + tprintf(TRACE_SYNC, " streamCreate: stream=%p\n", this); _signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1); - -#if 0 - auto s = this; - std::for_each(_signalPool.begin(), _signalPool.end(), - [s](ihipSignal_t &iter) { - printf (" stream:%p allocated hsa_signal=%lu\n", s, (iter._hsa_signal.handle)); - }); -#endif - }; + //--- ihipStream_t::~ihipStream_t() { @@ -262,6 +280,17 @@ ihipStream_t::~ihipStream_t() } +void ihipStream_t::wait() { + tprintf (TRACE_SYNC, "stream %p wait for queue-empty and lastCopy:#%lu...\n", this, _last_copy_signal ? _last_copy_signal->_seq_id: 0x0 ); + _av.wait(); + if (_last_copy_signal) { + hsa_signal_wait_acquire(_last_copy_signal->_hsa_signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); + // TODO-stream : reset ? + } + +}; + + //--- inline ihipDevice_t * ihipStream_t::getDevice() const { @@ -280,8 +309,9 @@ ihipSignal_t *ihipStream_t::getSignal() _signalCursor = 0; } - if (_signalPool[thisCursor]._refCnt == 0) { - _signalPool[thisCursor]._refCnt ++; // allocate it + if (_signalPool[thisCursor]._ref_cnt == 0) { + _signalPool[thisCursor]._ref_cnt ++; // allocate it + _signalPool[thisCursor]._seq_id = ++_seq_signal_id; // allocate it return &_signalPool[thisCursor]; } @@ -291,19 +321,19 @@ ihipSignal_t *ihipStream_t::getSignal() assert(numToScan == 0); // Have to grow the pool: - printf ("Grow signal pool\n"); _signalCursor = _signalPool.size(); // set to the beginning of the new entries: _signalPool.resize(_signalPool.size() * 2); + tprintf (TRACE_SYNC, "grow signal pool to %zu entries, cursor=%d\n", _signalPool.size(), _signalCursor); return getSignal(); // try again, - // Shouldnever reach here. + // Should never reach here. assert(0); } void ihipStream_t::releaseSignal(ihipSignal_t *signal) { - if (--signal->_refCnt <= 0) { + if (--signal->_ref_cnt <= 0) { // restore signal to the initial value 1 hsa_signal_store_release(signal->_hsa_signal, 1); } @@ -714,7 +744,7 @@ static inline void ihipWaitAllStreams(ihipDevice_t *device) { tprintf(TRACE_SYNC, "waitAllStream\n"); for (auto streamI=device->_streams.begin(); streamI!=device->_streams.end(); streamI++) { - (*streamI)->_av.wait(); + (*streamI)->wait(); } } @@ -730,7 +760,7 @@ inline void ihipWaitNullStream(ihipDevice_t *device) if (!(stream->_flags & hipStreamNonBlocking)) { // TODO-hcc - use blocking or active wait here? // TODO-sync - cudaDeviceBlockingSync - stream->_av.wait(); + stream->wait(); } } } @@ -753,17 +783,9 @@ inline hipStream_t ihipSyncAndResolveStream(hipStream_t stream) } } -#if 0 -inline hsa_status_t -HSABarrier::enqueueBarrier(hsa_queue_t* queue) { - hsa_status_t status = HSA_STATUS_SUCCESS; - hc::completion_future marker = stream->_av.create_marker(); - - // Create a signal to wait for the barrier to finish. - std::pair ret = Kalmar::ctx.getSignal(); - signal = ret.first; - signalIndex = ret.second; +void +ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal) { // Obtain the write index for the command queue uint64_t index = hsa_queue_load_write_index_relaxed(queue); @@ -776,21 +798,20 @@ HSABarrier::enqueueBarrier(hsa_queue_t* queue) { // 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; + //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->completion_signal = signal; + barrier->dep_signal[0] = depSignal->_hsa_signal; + barrier->completion_signal.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); - - isDispatched = true; - - return status; } -#endif + //-- //When the commands in a stream change types (ie kernel command follows a data command, @@ -798,36 +819,93 @@ HSABarrier::enqueueBarrier(hsa_queue_t* queue) { //into the stream to mimic CUDA stream semantics. (some hardware uses separate //queues for data commands and kernel commands, and no implicit ordering is provided). // -inline bool ihipCheckCommandSwitchSync(hipStream_t stream, ihipCommand_t new_command, hc::completion_future *marker) +inline bool ihipStream_t::preKernelCommand() { bool addedSync = false; // If switching command types, we need to add a barrier packet to synchronize things. - if (stream->_last_command != new_command) { - addedSync = true; - *marker = stream->_av.create_marker(); + if (_last_command_type != ihipCommandKernel) { + if (_last_copy_signal) { + addedSync = true; - tprintf (TRACE_SYNC, "stream %p switch %s to %s (barrier pkt inserted)\n", - (void*)stream, - stream->_last_command == ihipCommandKernel ? "Kernel" : "Data", - new_command == ihipCommandKernel ? "Kernel" : "Data"); - stream->_last_command = new_command; + hsa_queue_t * q = (hsa_queue_t*)_av.get_hsa_queue(); + this->enqueueBarrier(q, _last_copy_signal); + + tprintf (TRACE_SYNC, "stream %p switch %s to %s (barrier pkt inserted)\n", + this, + ihipCommandName[_last_command_type], ihipCommandName[ihipCommandKernel]) + } + _last_command_type = ihipCommandKernel; } return addedSync; } +//--- +inline void ihipStream_t::postKernelCommand(hc::completion_future &kernelFuture) +{ + _last_kernel_future = kernelFuture; +}; + + + +//--- +// Called whenever a copy command is set to the stream. +// Examines the last command sent to this stream and returns a signal to wait on, if required. +inline int ihipStream_t::copyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType) +{ + int needSync = 0; + // If switching command types, we need to add a barrier packet to synchronize things. + if (_last_command_type != copyType) { + needSync = 1; + + + if (_last_command_type == ihipCommandKernel) { + tprintf (TRACE_SYNC, "stream %p switch %s to %s (async copy dep on prev kernel)\n", + this, ihipCommandName[_last_command_type], ihipCommandName[copyType]); + hsa_signal_t *hsaSignal = (static_cast (_last_kernel_future.get_native_handle())); + if (hsaSignal) { + *waitSignal = * hsaSignal; + } + } else if (_last_copy_signal) { + tprintf (TRACE_SYNC, "stream %p switch %s to %s (async copy dep on other copy)\n", + this, ihipCommandName[_last_command_type], ihipCommandName[copyType]); + assert (_last_copy_signal->_ref_cnt > 0); + *waitSignal = _last_copy_signal->_hsa_signal; + } + + _last_command_type = copyType; + } + + _last_copy_signal = lastCopy; + + return needSync; +} + + + + + +// TODO - data-up to data-down: // Called just before a kernel is launched from hipLaunchKernel. // Allows runtime to track some information about the stream. -hc::accelerator_view *ihipLaunchKernel(hipStream_t stream) +hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av) { - stream = ihipSyncAndResolveStream(stream); - hc::completion_future marker; - ihipCheckCommandSwitchSync(stream, ihipCommandKernel, &marker); + stream->preKernelCommand(); - return &(stream->_av); + *av = &stream->_av; + + return (stream); +} + + +//--- +//Called after kernel finishes execution. +void ihipPostLaunchKernel(hipStream_t stream, hc::completion_future &kernelFuture) +{ + stream->postKernelCommand(kernelFuture); } @@ -1202,7 +1280,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int { // Super-conservative version of this - TODO - remove me: - stream->_av.wait(); + stream->wait(); e = hipSuccess; } @@ -1220,7 +1298,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) ihipDevice_t *device = ihipGetTlsDefaultDevice(); ihipWaitNullStream(device); } else { - stream->_av.wait(); + stream->wait(); e = hipSuccess; } @@ -1389,7 +1467,7 @@ void ihipSetTs(hipEvent_t e) // already recorded, done: return; } else { - // Test this code: + // TODO - use completion-future functions to obtain ticks and timestamps: hsa_signal_t *sig = static_cast (eh->_marker.get_native_handle()); if (sig) { if (hsa_signal_load_acquire(*sig) == 0) { @@ -1694,8 +1772,9 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou } auto device = ihipGetTlsDefaultDevice(); - hc::completion_future marker; - ihipCheckCommandSwitchSync(device._null_stream, ihipCommandData, &marker); + //hsa_signal_t depSignal; + //int depSignalCnt = device._null_stream->copyCommand(NULL, &depSignal, ihipCommandCopyH2D); + assert(0); // Need to properly synchronize the copy - do something with depSignal if != NULL. device->_acc.memcpy_symbol(symbolName, (void*) src,count, offset); #endif @@ -1762,7 +1841,12 @@ void StagingBuffer::CopyHostToDevice(void* dst, const void* src, size_t sizeByte hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); #if USE_ROCR_V2 - hsa_status_t hsa_status = hsa_amd_memory_async_copy(dstp, _device->_hsa_agent, _pinnedStagingBuffer[bufferIndex], _device->_hsa_agent, theseBytes, 0, NULL, _completion_signal[bufferIndex]); + hsa_signal_t depSignal; + int depSignalCnt = 0; //stream->copyCommand(_completion_signal[bufferIndex], &depSignal, copyType); + + printf ("need sync\n"); + + hsa_status_t hsa_status = hsa_amd_memory_async_copy(dstp, _device->_hsa_agent, _pinnedStagingBuffer[bufferIndex], _device->_hsa_agent, theseBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, _completion_signal[bufferIndex]); #else hsa_status_t hsa_status = hsa_amd_memory_async_copy(dstp, _pinnedStagingBuffer[bufferIndex], theseBytes, _device->_hsa_agent, 0, NULL, _completion_signal[bufferIndex]); #endif @@ -1807,7 +1891,13 @@ void StagingBuffer::CopyDeviceToHost(void* dst, const void* src, size_t sizeByte tprintf (TRACE_COPY2, "D2H: async_copy %zu bytes src:%p to staging:%p\n", theseBytes, srcp0, _pinnedStagingBuffer[bufferIndex]); hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); #if USE_ROCR_V2 - hsa_status_t hsa_status = hsa_amd_memory_async_copy(_pinnedStagingBuffer[bufferIndex], _device->_hsa_agent, srcp0, _device->_hsa_agent, theseBytes, 0, NULL, _completion_signal[bufferIndex]); + hsa_signal_t depSignal; + // TODO + int depSignalCnt = 0; //stream->copyCommand(_completion_signal[bufferIndex], &depSignal, copyType); + + printf ("need sync\n"); + + hsa_status_t hsa_status = hsa_amd_memory_async_copy(_pinnedStagingBuffer[bufferIndex], _device->_hsa_agent, srcp0, _device->_hsa_agent, theseBytes, depSignalCnt, depSignalCnt ? &depSignal:0, _completion_signal[bufferIndex]); #else hsa_status_t hsa_status = hsa_amd_memory_async_copy(_pinnedStagingBuffer[bufferIndex], srcp0, theseBytes, _device->_hsa_agent, 0, NULL, _completion_signal[bufferIndex]); #endif @@ -1867,6 +1957,14 @@ void ihipSyncCopy(ihipDevice_t *device, void* dst, const void* src, size_t sizeB } } +#if 0 + + //TODO + hsa_signal_t depSignal; + int dep_signals = stream->commandCopy(&depSignal, ); + pass to CopyHostToDevice +#endif + if ((kind == hipMemcpyHostToDevice) && (srcNotTracked)) { if (useStagingBuffer) { std::lock_guard l (device->_copy_lock[0]); @@ -1918,7 +2016,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); hc::completion_future marker; - ihipCheckCommandSwitchSync(stream, ihipCommandData, &marker); + hipError_t e = hipSuccess; @@ -1955,44 +2053,52 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp hipError_t e = hipSuccess; - stream = ihipSyncAndResolveStream(stream); - - hc::completion_future marker; - ihipCheckCommandSwitchSync(stream, ihipCommandData, &marker); - - // Dispatch async memory copy to synchronize with items in the specified stream. - - // Async - need to set up dependency on the last command queued to the device? - + stream = ihipSyncAndResolveStream(stream); #if USE_AM_TRACKER - - hipStream_t s = ihipSyncAndResolveStream(stream); - - if (s) { - ihipDevice_t *device = s->getDevice(); + if (stream) { + ihipDevice_t *device = stream->getDevice(); if (kind == hipMemcpyDefault) { e = hipErrorInvalidMemcpyDirection; + + } else if (kind == hipMemcpyHostToHost) { + tprintf (TRACE_COPY2, "H2H copy with memcpy"); + + memcpy(dst, src, sizeBytes); + } else { - // Let HSA runtime handle it: - // TODO - need buffer pool for the signals rather than lock: ihipSignal_t *ihip_signal = stream->getSignal(); - //stream->saveLastSignal(ihipSignal); + ihipCommand_t copyType; + if ((kind == hipMemcpyHostToDevice) || (kind == hipMemcpyDeviceToDevice)) { + copyType = ihipCommandCopyH2D; + } else if (kind == hipMemcpyDeviceToHost) { + copyType = ihipCommandCopyD2H; + } else { + e = hipErrorInvalidMemcpyDirection; + copyType = ihipCommandCopyD2H; + } #if USE_ROCR_V2 - hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, device->_hsa_agent, src, device->_hsa_agent, sizeBytes, 0, NULL, ihip_signal->_hsa_signal); + hsa_signal_t depSignal; + int depSignalCnt = stream->copyCommand(ihip_signal, &depSignal, copyType); + + tprintf (TRACE_SYNC, " copy-async, waitFor=%d(%lu) completion=%lu\n", depSignalCnt, depSignal.handle, ihip_signal->_seq_id); + + hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, device->_hsa_agent, src, device->_hsa_agent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsa_signal); #else hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, src, sizeBytes, device->_hsa_agent, 0, NULL, ihip_signal->_hsa_signal); #endif - if (hsa_status == HSA_STATUS_SUCCESS) { + if (hsa_status == HSA_STATUS_SUCCESS) { if (HIP_LAUNCH_BLOCKING) { hsa_signal_wait_relaxed(ihip_signal->_hsa_signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); stream->releaseSignal(ihip_signal); - } + } else { + //stream->releaseSignal(ihip_signal); + } } else { // This path can be hit if src or dst point to unpinned host memory. // TODO - does async-copy fall back to sync if input pointers are not pinned? @@ -2014,6 +2120,8 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp // 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 ) { std::call_once(hip_initialized, ihipInit); @@ -2021,37 +2129,42 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s hipError_t e = hipSuccess; stream = ihipSyncAndResolveStream(stream); - hc::completion_future marker; - ihipCheckCommandSwitchSync(stream, ihipCommandData, &marker); + stream->preKernelCommand(); + if (stream) { - hc::completion_future cf ; + hc::completion_future cf ; - if ((sizeBytes & 0x3) == 0) { - // use a faster word-per-workitem copy: - try { - value = value & 0xff; - unsigned value32 = (value << 24) | (value << 16) | (value << 8) | (value) ; - cf = ihipMemsetKernel (stream, static_cast (dst), value32, sizeBytes/sizeof(unsigned)); - } - catch (std::exception &ex) { - e = hipErrorInvalidValue; + if ((sizeBytes & 0x3) == 0) { + // use a faster word-per-workitem copy: + try { + value = value & 0xff; + unsigned value32 = (value << 24) | (value << 16) | (value << 8) | (value) ; + cf = ihipMemsetKernel (stream, static_cast (dst), value32, sizeBytes/sizeof(unsigned)); + } + catch (std::exception &ex) { + e = hipErrorInvalidValue; + } + } else { + // use a slow byte-per-workitem copy: + try { + cf = ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); + } + catch (std::exception &ex) { + e = hipErrorInvalidValue; + } } + + stream->postKernelCommand(cf); + + + if (HIP_LAUNCH_BLOCKING) { + tprintf (TRACE_SYNC, "'%s' LAUNCH_BLOCKING wait for completion [stream:%p].\n", __func__, (void*)stream); + cf.wait(); + tprintf (TRACE_SYNC, "'%s' LAUNCH_BLOCKING completed [stream:%p].\n", __func__, (void*)stream); + } } else { - // use a slow byte-per-workitem copy: - try { - cf = ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); - } - catch (std::exception &ex) { - e = hipErrorInvalidValue; - } - } - - - if (HIP_LAUNCH_BLOCKING) { - tprintf (TRACE_SYNC, "'%s' LAUNCH_BLOCKING wait for completion [stream:%p].\n", __func__, (void*)stream); - cf.wait(); - tprintf (TRACE_SYNC, "'%s' LAUNCH_BLOCKING completed [stream:%p].\n", __func__, (void*)stream); + e = hipErrorInvalidValue; } diff --git a/projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp b/projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp index 19f1a94761..8669b986d8 100644 --- a/projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp +++ b/projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp @@ -36,24 +36,62 @@ void simpleNegTest() //Send many async copies to the same stream. //This requires runtime to keep track of many outstanding commands, and in the case of HCC requires growing/tracking the signal pool: template -void test_manyCopies(int nElements, size_t numCopies, int nStreams) +void test_manyCopies(int nElements, int numCopies) { size_t Nbytes = nElements*sizeof(T); - printf ("Nbytes=%zu (%6.1f MB)\n", Nbytes, (double)(Nbytes)/1024.0/1024.0); + size_t eachCopyElements = nElements / numCopies; + size_t eachCopyBytes = eachCopyElements * sizeof(T); - int *A_d, *B_d, *C_d; - int *A_h, *B_h, *C_h; + printf ("-----------------------------------------------------------------------------------------------\n"); + printf ("testing: %s Nbytes=%zu (%6.1f MB) numCopies=%d eachCopyElements=%zu eachCopyBytes=%zu\n", + __func__, Nbytes, (double)(Nbytes)/1024.0/1024.0, numCopies, eachCopyElements, eachCopyBytes); - HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, true); + T *A_d; + T *A_h1, *A_h2; - size_t eachCopyBytes = Nbytes / numCopies; - - for (size_t i=0; i (i); } - HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, true); + + hipStream_t stream; + HIPCHECK (hipStreamCreate(&stream)); + + //stream=0; // fixme TODO + + + for (int i=0; i(1024, 16); + test_manyCopies(1024, 4); + test_manyCopies(1024*4, 64); + } - test_chunkedAsyncExample(p_streams, true, true, true); // Easy sync version - test_chunkedAsyncExample(p_streams, false, true, true); // Easy sync version - test_chunkedAsyncExample(p_streams, false, false, true); // Some async - test_chunkedAsyncExample(p_streams, false, false, false); // All async + if (p_tests & 0x4) { + test_chunkedAsyncExample(p_streams, true, true, true); // Easy sync version + test_chunkedAsyncExample(p_streams, false, true, true); // Easy sync version + test_chunkedAsyncExample(p_streams, false, false, true); // Some async + test_chunkedAsyncExample(p_streams, false, false, false); // All async + } diff --git a/projects/clr/hipamd/tests/src/test_common.h b/projects/clr/hipamd/tests/src/test_common.h index 5b631d2c3a..f133696d78 100644 --- a/projects/clr/hipamd/tests/src/test_common.h +++ b/projects/clr/hipamd/tests/src/test_common.h @@ -25,7 +25,7 @@ printf (__VA_ARGS__);\ printf ("\n");\ printf ("error: TEST FAILED\n%s", KNRM );\ - exit(EXIT_FAILURE); + abort(); #define HIPCHECK(error) \