From e9764276c0f8304b6428a8a62cf783292bf3cdd3 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 24 Mar 2016 14:33:11 -0500 Subject: [PATCH 01/16] add note on using HIP_PLATFORM to force hcc path [ROCm/clr commit: 01ac10e345f943aa5d0d5e5aedeab00141eda019] --- projects/clr/hipamd/docs/markdown/hip_faq.md | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/projects/clr/hipamd/docs/markdown/hip_faq.md b/projects/clr/hipamd/docs/markdown/hip_faq.md index 734ced94ed..2d00e7f60c 100644 --- a/projects/clr/hipamd/docs/markdown/hip_faq.md +++ b/projects/clr/hipamd/docs/markdown/hip_faq.md @@ -106,3 +106,9 @@ HIP is a source-portable language that can be compiled to run on either the HCC HIP is a portable C++ language that supports a strong subset of the CUDA run-time APIs and device-kernel language. It's designed to simplify CUDA conversion to portable C++. HIP provides a C-compatible run-time API, C-compatible kernel-launch mechanism, C++ kernel language and pointer-based memory management. A C++ dialect, hc is supported by the AMD HCC compiler. It provides C++ run time, C++ kernel-launch APIs (parallel_for_each), C++ kernel language, and several memory-management options, including pointers, arrays and array_view (with implicit data synchronization). It's intended to be a leading indicator of the ISO C++ standard. + +### What if I have a CUDA SDK installed but want to use HCC? +If HIP sees the CUDA SDK installed at /usr/local/cuda, it assumes the platform is nvcc. Sometimes this isn't what you want - you can force HIP to recognize the platform by setting HIP_PLATFORM to hcc (or nvcc) +``` +export HIP_PLATFORM=hcc +``` From ecd13f750d6442dae03b45c81b04a2a74d755ec1 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 24 Mar 2016 22:12:41 -0500 Subject: [PATCH 02/16] report linux distro if possible [ROCm/clr commit: 120d45b0dcd52a7cd1e26556e93e62a178bb4f47] --- projects/clr/hipamd/bin/hipconfig | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/projects/clr/hipamd/bin/hipconfig b/projects/clr/hipamd/bin/hipconfig index d5f068c097..af85e335a6 100755 --- a/projects/clr/hipamd/bin/hipconfig +++ b/projects/clr/hipamd/bin/hipconfig @@ -116,6 +116,10 @@ if (!$printed or $p_full) { print "== Linux Kernel\n"; system ("uname -a"); + if (-e "/usr/bin/lsb_release") { + system ("/usr/bin/lsb_release -a"); + } + print "\n" ; $printed = 1; } From 8cf3f24d5c2e32a7863a05927a6b5a45804c651c Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 25 Mar 2016 17:08:34 -0500 Subject: [PATCH 03/16] Improve automated platform detection. If AMD GPU is installed and detected by driver, default HIP_PLATFORM to hcc. [ROCm/clr commit: 6d9eafa18ae96e60d94ec91b9042607eb3d7d368] --- projects/clr/hipamd/bin/hipconfig | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/bin/hipconfig b/projects/clr/hipamd/bin/hipconfig index af85e335a6..2d880f157d 100755 --- a/projects/clr/hipamd/bin/hipconfig +++ b/projects/clr/hipamd/bin/hipconfig @@ -45,10 +45,15 @@ $HSA_PATH='/opt/hsa' unless defined $HSA_PATH; #--- #HIP_PLATFORM controls whether to use NVCC or HCC for compilation: $HIP_PLATFORM=$ENV{'HIP_PLATFORM'}; -if (not defined $HIP_PLATFORM and (-e "$CUDA_PATH/bin/nvcc")) { - $HIP_PLATFORM="nvcc"; +if (not defined $HIP_PLATFORM) { + $NAMDGPUNODES=`cat /sys/class/kfd/kfd/topology/nodes/*/properties 2>/dev/null | grep -c 'simd_count [1-9]'`; + + if ($NAMDGPUNODES > 0) { + $HIP_PLATFORM = "hcc" + } else { + $HIP_PLATFORM = "nvcc"; + } } -$HIP_PLATFORM="hcc" unless defined $HIP_PLATFORM; $HIP_PATH=$ENV{'HIP_PATH'}; $HIP_PATH=Cwd::realpath (dirname (dirname $0)) unless defined $HIP_PATH; # use parent directory of this tool From bdc842f7f2eacc16083badaac67ed668517bd9c8 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 25 Mar 2016 17:11:49 -0500 Subject: [PATCH 04/16] describe HIP_PLATFORM [ROCm/clr commit: f8f839d9dd837d0301f86fd46a4018cb5f9f5b5e] --- projects/clr/hipamd/docs/markdown/hip_faq.md | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/hip_faq.md b/projects/clr/hipamd/docs/markdown/hip_faq.md index 2d00e7f60c..1a62784100 100644 --- a/projects/clr/hipamd/docs/markdown/hip_faq.md +++ b/projects/clr/hipamd/docs/markdown/hip_faq.md @@ -107,8 +107,9 @@ HIP is a portable C++ language that supports a strong subset of the CUDA run-tim A C++ dialect, hc is supported by the AMD HCC compiler. It provides C++ run time, C++ kernel-launch APIs (parallel_for_each), C++ kernel language, and several memory-management options, including pointers, arrays and array_view (with implicit data synchronization). It's intended to be a leading indicator of the ISO C++ standard. -### What if I have a CUDA SDK installed but want to use HCC? -If HIP sees the CUDA SDK installed at /usr/local/cuda, it assumes the platform is nvcc. Sometimes this isn't what you want - you can force HIP to recognize the platform by setting HIP_PLATFORM to hcc (or nvcc) +### HIP detected my platform (hcc vs nvcc) incorrectly - what should I do? +HIP will set the platform to HCC if it sees that the AMD graphics driver is installed and has detected an AMD GPU. +Sometimes this isn't what you want - you can force HIP to recognize the platform by setting HIP_PLATFORM to hcc (or nvcc) ``` export HIP_PLATFORM=hcc ``` From e16e848d5562f18b50114a37d82d8a995f6b5b87 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 26 Mar 2016 10:46:20 -0500 Subject: [PATCH 05/16] Make ihipDevice_t thread-safe. Move critical data into separate class and protect with LockAccessor wrapper class. For device, the streams list is the critical data since it is modified when streams are created or destroyed. The streams list is accessed in several places including when synchronizing across all streams on the device (ie from the default stream). Other device data is set once by the device cosntructor and is not critical so All functions which acquire the LockAccessor now named with "locked_" prefix. [ROCm/clr commit: 4dd77c66124d4543c05e169483f0b360560a6618] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 88 +++++++++-- projects/clr/hipamd/src/hip_device.cpp | 14 +- projects/clr/hipamd/src/hip_event.cpp | 4 +- projects/clr/hipamd/src/hip_hcc.cpp | 54 +++++-- projects/clr/hipamd/src/hip_memory.cpp | 2 +- projects/clr/hipamd/src/hip_stream.cpp | 10 +- projects/clr/hipamd/tests/src/CMakeLists.txt | 9 +- .../hipamd/tests/src/hipThreadSafeDevice.cpp | 137 ++++++++++++++++++ 8 files changed, 279 insertions(+), 39 deletions(-) create mode 100644 projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 9000602ff1..409573d2dc 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -63,8 +63,8 @@ extern int HIP_DISABLE_HW_COPY_DEP; extern thread_local int tls_defaultDevice; extern thread_local hipError_t tls_lastHipError; -struct ihipStream_t; -struct ihipDevice_t; +class ihipStream_t; +class ihipDevice_t; // Color defs for debug messages: @@ -86,6 +86,9 @@ struct ihipDevice_t; // Stream functions will acquire a mutex before entering critical sections. #define STREAM_THREAD_SAFE 1 + +#define DEVICE_THREAD_SAFE 1 + // If FORCE_COPY_DEP=1 , HIP runtime will add // synchronization for copy commands in the same stream, regardless of command type. // If FORCE_COPY_DEP=0 data copies of the same kind (H2H, H2D, D2H, D2D) are assumed to be implicitly ordered. @@ -376,12 +379,76 @@ struct ihipEvent_t { } ; +//--- +// Protects access to the member _data with a lock acquired on contruction/destruction. +// T must contain a _mutex field which meets the BasicLockable requirements (lock/unlock) +template +class LockedAccessor +{ +public: + LockedAccessor(T &data) : _data(&data) + { + _data->_mutex.lock(); + }; + + ~LockedAccessor() + { + _data->_mutex.unlock(); + } + + // Syntactic sugar so -> can be used to get the underlying type. + T *operator->() { return _data; }; + +private: + T *_data; +}; + + + +//--- +// 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 +{ +public: + friend class LockedAccessor; + + std::list &streams() { return _streams; }; + const std::list &const_streams() const { return _streams; }; + +private: + std::list _streams; // streams associated with this device. + MUTEX_TYPE _mutex; +}; + +// Typedefs for common definitions. +#if DEVICE_THREAD_SAFE +typedef ihipDeviceCriticalBase_t ihipDeviceCritical_t; // Use real mutex +#else +#warning "Device thread-safe disabled" +typedef ihipDeviceCriticalBase_t ihipDeviceCritical_t; // Fake mutex, for testing +#endif + +typedef LockedAccessor Locked_ihipDeviceCritical_t; //------------------------------------------------------------------------------------------------- -struct ihipDevice_t +// Functions which read or write the critical data are named locked_. +// ihipDevice_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 ihipDevice_t { +public: // Functions: + void locked_addStream(ihipStream_t *s); + void locked_removeStream(ihipStream_t *s); + +public: // Data, set at initialization: unsigned _device_index; // index into g_devices. hipDeviceProp_t _props; // saved device properties. @@ -392,7 +459,6 @@ struct ihipDevice_t // NULL has special synchronization properties with other streams. ihipStream_t *_default_stream; - std::list _streams; // streams associated with this device. unsigned _compute_units; @@ -403,20 +469,24 @@ struct ihipDevice_t unsigned _device_flags; public: - void init(unsigned device_index, hc::accelerator acc, unsigned flags); + void locked_init(unsigned device_index, hc::accelerator acc, unsigned flags); ~ihipDevice_t(); - void reset(); hipError_t getProperties(hipDeviceProp_t* prop); - void waitAllStreams(); - void syncDefaultStream(bool waitOnSelf); + void locked_reset(); + void locked_waitAllStreams(); + void locked_syncDefaultStream(bool waitOnSelf); private: + // Members of _protected data MUST be accessed through the LockedAccessor. + // Search for LockedAccessor for examples; do not access _criticalData directly. + ihipDeviceCritical_t _criticalData; }; -// Global initialization. + +// Global variable definition: extern std::once_flag hip_initialized; extern ihipDevice_t *g_devices; // Array of all non-emulated (ie GPU) accelerators in the system. extern bool g_visible_device; // Set the flag when HIP_VISIBLE_DEVICES is set diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index b819673967..c935f20d34 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/projects/clr/hipamd/src/hip_device.cpp @@ -150,8 +150,7 @@ hipError_t hipDeviceSynchronize(void) { HIP_INIT_API(); - ihipGetTlsDefaultDevice()->waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. - + ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. return ihipLogStatus(hipSuccess); } @@ -174,15 +173,12 @@ hipError_t hipDeviceReset(void) if (device) { //--- - //Wait for pending activity to complete? - //TODO - check if this is required behavior: - for (auto streamI=device->_streams.begin(); streamI!=device->_streams.end(); streamI++) { - ihipStream_t *stream = *streamI; - stream->wait(); - } + //Wait for pending activity to complete? TODO - check if this is required behavior: + + device->locked_waitAllStreams(); // Release device resources (streams and memory): - device->reset(); + device->locked_reset(); } return ihipLogStatus(hipSuccess); diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 27232875b9..66a7c7c7fc 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -67,7 +67,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) // TODO-HCC fix this - is CUDA this conservative or still uses device timestamps? // TODO-HCC can we use barrier or event marker to implement better solution? ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->syncDefaultStream(true); + device->locked_syncDefaultStream(true); eh->_timestamp = hc::get_system_ticks(); eh->_state = hipEventStatusRecorded; @@ -117,7 +117,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) return ihipLogStatus(hipSuccess); } else if (eh->_stream == NULL) { ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->syncDefaultStream(true); + device->locked_syncDefaultStream(true); return ihipLogStatus(hipSuccess); } else { #if __hcc_workweek__ >= 16033 diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 1c93ae48dd..cb216919e2 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -145,6 +145,8 @@ ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, SeqNu ihipStream_t::~ihipStream_t() { _signalPool.clear(); + // Hack to catch memory issues, in particular accesses to _acc after it has been destroyed. + //memset (&_av, 0x0, sizeof(hc::accelerator_view&)); } @@ -348,6 +350,8 @@ int ihipStream_t::preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSigna hsa_signal_t *hsaSignal = (static_cast (_last_kernel_future.get_native_handle())); if (hsaSignal) { *waitSignal = * hsaSignal; + } else { + assert(0); // if NULL signal, and we return 1, hsa_amd_memory_copy_async will fail. Confirm this never happens. } } else if (_last_copy_signal) { needSync = 1; @@ -383,10 +387,13 @@ int ihipStream_t::preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSigna // //Reset the device - this is called from hipDeviceReset. //Device may be reset multiple times, and may be reset after init. -void ihipDevice_t::reset() +void ihipDevice_t::locked_reset() { + // Obtain mutex access to the device critical data, release by destructor + Locked_ihipDeviceCritical_t l(_criticalData); + // Reset and remove streams: - _streams.clear(); + l->streams().clear(); // Reset and release all memory stored in the tracker: am_memtracker_reset(_acc); @@ -395,7 +402,7 @@ void ihipDevice_t::reset() //--- -void ihipDevice_t::init(unsigned device_index, hc::accelerator acc, unsigned flags) +void ihipDevice_t::locked_init(unsigned device_index, hc::accelerator acc, unsigned flags) { _stream_id = 0; @@ -416,8 +423,12 @@ void ihipDevice_t::init(unsigned device_index, hc::accelerator acc, unsigned fla getProperties(&_props); - _default_stream = new ihipStream_t(device_index, acc.get_default_view(), _stream_id++, hipStreamDefault); - this->_streams.push_back(_default_stream); + { + Locked_ihipDeviceCritical_t l(_criticalData); + _default_stream = new ihipStream_t(device_index, acc.get_default_view(), _stream_id++, hipStreamDefault); + l->streams().push_back(_default_stream); + } + tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream); @@ -670,11 +681,13 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) // 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 ihipDevice_t::syncDefaultStream(bool waitOnSelf) +void ihipDevice_t::locked_syncDefaultStream(bool waitOnSelf) { + Locked_ihipDeviceCritical_t l(_criticalData); + tprintf(DB_SYNC, "syncDefaultStream\n"); - for (auto streamI=_streams.begin(); streamI!=_streams.end(); streamI++) { + for (auto streamI=l->const_streams().begin(); streamI!=l->const_streams().end(); streamI++) { ihipStream_t *stream = *streamI; // Don't wait for streams that have "opted-out" of syncing with NULL stream. @@ -690,13 +703,31 @@ void ihipDevice_t::syncDefaultStream(bool waitOnSelf) } } +//--- +void ihipDevice_t::locked_addStream(ihipStream_t *s) +{ + Locked_ihipDeviceCritical_t l(_criticalData); + + l->streams().push_back(s); +} + +//--- +void ihipDevice_t::locked_removeStream(ihipStream_t *s) +{ + Locked_ihipDeviceCritical_t l(_criticalData); + + l->streams().remove(s); +} + //--- //Heavyweight synchronization that waits on all streams, ignoring hipStreamNonBlocking flag. -void ihipDevice_t::waitAllStreams() +void ihipDevice_t::locked_waitAllStreams() { + Locked_ihipDeviceCritical_t l(_criticalData); + tprintf(DB_SYNC, "waitAllStream\n"); - for (auto streamI=_streams.begin(); streamI!=_streams.end(); streamI++) { + for (auto streamI=l->const_streams().begin(); streamI!=l->const_streams().end(); streamI++) { (*streamI)->wait(); } } @@ -864,7 +895,7 @@ void ihipInit() //If device is not in visible devices list, ignore continue; } - g_devices[g_deviceCnt].init(g_deviceCnt, accs[i], hipDeviceMapHost); + g_devices[g_deviceCnt].locked_init(g_deviceCnt, accs[i], hipDeviceMapHost); g_deviceCnt++; } } @@ -933,7 +964,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream) ihipDevice_t *device = ihipGetTlsDefaultDevice(); #ifndef HIP_API_PER_THREAD_DEFAULT_STREAM - device->syncDefaultStream(false); + device->locked_syncDefaultStream(false); #endif return device->_default_stream; } else { @@ -1302,6 +1333,7 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a // TODO - describe naming convention. ihip _. No accessors. No early returns from functions. Set status to success at top, only set error codes in implementation. No tabs. // Caps convention _ or camelCase // if { } +// Should use ihip* data structures inside code rather than app-facing hip. For example, use ihipDevice_t (rather than hipDevice_t), ihipStream_t (rather than hipStream_t). // TODO - describe MT strategy // //// TODO - add identifier numbers for streams and devices to help with debugging. diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 46fbd38fd5..c633d23611 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -459,7 +459,7 @@ hipError_t hipFree(void* ptr) hipError_t hipStatus = hipErrorInvalidDevicePointer; // Synchronize to ensure all work has finished. - ihipGetTlsDefaultDevice()->waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. + ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. if (ptr) { hc::accelerator acc; diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 6dbf7458ce..784fa9224c 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -42,7 +42,9 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) //This matches CUDA stream behavior: auto istream = new ihipStream_t(device->_device_index, acc.create_view(), device->_stream_id, flags); - device->_streams.push_back(istream); + + device->locked_addStream(istream); + *stream = istream; tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream); @@ -82,7 +84,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) if (stream == NULL) { ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->syncDefaultStream(true/*waitOnSelf*/); + device->locked_syncDefaultStream(true/*waitOnSelf*/); } else { stream->wait(); e = hipSuccess; @@ -106,7 +108,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) //--- Drain the stream: if (stream == NULL) { ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->syncDefaultStream(true/*waitOnSelf*/); + device->locked_syncDefaultStream(true/*waitOnSelf*/); } else { stream->wait(); e = hipSuccess; @@ -115,7 +117,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) ihipDevice_t *device = stream->getDevice(); if (device) { - device->_streams.remove(stream); + device->locked_removeStream(stream); delete stream; } else { e = hipErrorInvalidResourceHandle; diff --git a/projects/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt index acd2060647..4288624ccf 100644 --- a/projects/clr/hipamd/tests/src/CMakeLists.txt +++ b/projects/clr/hipamd/tests/src/CMakeLists.txt @@ -9,7 +9,6 @@ set (HIP_Unit_Test_VERSION_MAJOR 1) set (HIP_Unit_Test_VERSION_MINOR 0) set(HIP_PATH $ENV{HIP_PATH}) -MESSAGE("HIP_PATH=" ${HIP_PATH}) if (NOT DEFINED HIP_PATH) set (HIP_PATH ../..) endif() @@ -39,12 +38,13 @@ if (${HIP_PLATFORM} STREQUAL "hcc") #These includes are used for all files. #Include HIP and HC since the tests need both of these: - #Note below HSA path is surgically included only where necessary. include_directories(${HIP_PATH}/include) - include_directories(${HSA_PATH}/include) # This will create a subdir "hip_hcc" in the test build directory # Any changes to hip_hcc source will be detected and force the library and then the tests to be rebuilt. + add_subdirectory(${HIP_PATH} build.hip_hcc) + link_directories(${CMAKE_CURRENT_BINARY_DIR}/build.hip_hcc) # search the local hip_hcc for libhip_hcc.a + elseif (${HIP_PLATFORM} STREQUAL "nvcc") MESSAGE ("HIP_PLATFORM=nvcc") @@ -156,6 +156,8 @@ make_hip_executable (hipFuncGetDevice hipFuncGetDevice.cpp) make_hip_executable (hipFuncSetDevice hipFuncSetDevice.cpp) make_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp) +make_hip_executable (hipThreadSafeDevice hipThreadSafeDevice.cpp) + make_test(hip_ballot " " ) make_test(hip_anyall " " ) make_test(hip_popc " " ) @@ -196,5 +198,6 @@ make_test(hipFuncSetDeviceFlags " ") make_test(hipFuncGetDevice " ") make_test(hipFuncSetDevice " ") make_test(hipFuncDeviceSynchronize " ") +make_test (hipThreadSafeDevice " ") make_hipify_test(specialFunc.cu ) diff --git a/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp b/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp new file mode 100644 index 0000000000..fca994cba8 --- /dev/null +++ b/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp @@ -0,0 +1,137 @@ +#include +#include "test_common.h" + + +// Create a lot of streams and then destroy 'em. +void createThenDestroyStreams(int iterations, int burstSize) +{ + hipStream_t *streams = new hipStream_t[burstSize]; + + for (int i=0; i Date: Sat, 26 Mar 2016 11:45:25 -0500 Subject: [PATCH 06/16] Protect _stream_id as well. - move lockedaccessor - clean up device class. - add simple ihipDevice constructor. [ROCm/clr commit: c47b5b04ef3f1d897c345270f78df351056d380b] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 84 ++++++++++--------- projects/clr/hipamd/src/hip_hcc.cpp | 22 +++-- projects/clr/hipamd/src/hip_stream.cpp | 2 +- .../hipamd/tests/src/hipThreadSafeDevice.cpp | 8 +- 4 files changed, 60 insertions(+), 56 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 409573d2dc..4e8b513f60 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -275,23 +275,47 @@ class FakeMutex #if STREAM_THREAD_SAFE typedef std::mutex StreamMutex; #else +#warning "Stream thread-safe disabled" typedef FakeMutex StreamMutex; #endif +// +//--- +// Protects access to the member _data with a lock acquired on contruction/destruction. +// T must contain a _mutex field which meets the BasicLockable requirements (lock/unlock) +template +class LockedAccessor +{ +public: + LockedAccessor(T &criticalData) : _criticalData(&criticalData) + { + _criticalData->_mutex.lock(); + }; + + ~LockedAccessor() + { + _criticalData->_mutex.unlock(); + } + + // Syntactic sugar so -> can be used to get the underlying type. + T *operator->() { return _criticalData; }; + +private: + T *_criticalData; +}; + + + + // TODO - move async copy code into stream? Stream->async-copy. // Add PreCopy / PostCopy to manage locks? -// - - - - // Internal stream structure. class ihipStream_t { public: typedef uint64_t SeqNum_t ; - ihipStream_t(unsigned device_index, hc::accelerator_view av, SeqNum_t id, unsigned int flags); + ihipStream_t(unsigned device_index, hc::accelerator_view av, unsigned int flags); ~ihipStream_t(); // kind is hipMemcpyKind @@ -379,29 +403,6 @@ struct ihipEvent_t { } ; -//--- -// Protects access to the member _data with a lock acquired on contruction/destruction. -// T must contain a _mutex field which meets the BasicLockable requirements (lock/unlock) -template -class LockedAccessor -{ -public: - LockedAccessor(T &data) : _data(&data) - { - _data->_mutex.lock(); - }; - - ~LockedAccessor() - { - _data->_mutex.unlock(); - } - - // Syntactic sugar so -> can be used to get the underlying type. - T *operator->() { return _data; }; - -private: - T *_data; -}; @@ -415,17 +416,22 @@ template class ihipDeviceCriticalBase_t { public: + ihipDeviceCriticalBase_t() : _stream_id(0) {}; friend class LockedAccessor; 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++; }; + + private: std::list _streams; // streams associated with this device. + ihipStream_t::SeqNum_t _stream_id; MUTEX_TYPE _mutex; }; -// Typedefs for common definitions. #if DEVICE_THREAD_SAFE typedef ihipDeviceCriticalBase_t ihipDeviceCritical_t; // Use real mutex #else @@ -445,8 +451,15 @@ typedef LockedAccessor Locked_ihipDeviceCritical_t; class ihipDevice_t { public: // Functions: + ihipDevice_t() {}; // note: calls constructor for _criticalData + void init(unsigned device_index, hc::accelerator &acc, unsigned flags); + ~ihipDevice_t(); + void locked_addStream(ihipStream_t *s); void locked_removeStream(ihipStream_t *s); + void locked_reset(); + void locked_waitAllStreams(); + void locked_syncDefaultStream(bool waitOnSelf); public: // Data, set at initialization: unsigned _device_index; // index into g_devices. @@ -464,20 +477,13 @@ public: // Data, set at initialization: StagingBuffer *_staging_buffer[2]; // one buffer for each direction. - ihipStream_t::SeqNum_t _stream_id; unsigned _device_flags; -public: - void locked_init(unsigned device_index, hc::accelerator acc, unsigned flags); - ~ihipDevice_t(); +private: hipError_t getProperties(hipDeviceProp_t* prop); - void locked_reset(); - void locked_waitAllStreams(); - void locked_syncDefaultStream(bool waitOnSelf); - -private: +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; diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index cb216919e2..ce136addd9 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -124,8 +124,8 @@ ihipSignal_t::~ihipSignal_t() // ihipStream_t: //================================================================================================= //--- -ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, SeqNum_t id, unsigned int flags) : - _id(id), +ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, unsigned int flags) : + _id(0), // will be set by add function. _av(av), _flags(flags), _device_index(device_index), @@ -402,13 +402,12 @@ void ihipDevice_t::locked_reset() //--- -void ihipDevice_t::locked_init(unsigned device_index, hc::accelerator acc, unsigned flags) +void ihipDevice_t::init(unsigned device_index, hc::accelerator &acc, unsigned flags) { - _stream_id = 0; - _device_index = device_index; _device_flags = flags; _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); @@ -423,15 +422,11 @@ void ihipDevice_t::locked_init(unsigned device_index, hc::accelerator acc, unsig getProperties(&_props); - { - Locked_ihipDeviceCritical_t l(_criticalData); - _default_stream = new ihipStream_t(device_index, acc.get_default_view(), _stream_id++, hipStreamDefault); - l->streams().push_back(_default_stream); - } + _default_stream = new ihipStream_t(device_index, acc.get_default_view(), hipStreamDefault); + locked_addStream(_default_stream); tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream); - hsa_region_t *pinnedHostRegion; pinnedHostRegion = static_cast(_acc.get_hsa_am_system_region()); _staging_buffer[0] = new StagingBuffer(_hsa_agent, *pinnedHostRegion, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS); @@ -440,6 +435,8 @@ void ihipDevice_t::locked_init(unsigned device_index, hc::accelerator acc, unsig }; + + ihipDevice_t::~ihipDevice_t() { if (_default_stream) { @@ -709,6 +706,7 @@ void ihipDevice_t::locked_addStream(ihipStream_t *s) Locked_ihipDeviceCritical_t l(_criticalData); l->streams().push_back(s); + s->_id = l->incStreamId(); } //--- @@ -895,7 +893,7 @@ void ihipInit() //If device is not in visible devices list, ignore continue; } - g_devices[g_deviceCnt].locked_init(g_deviceCnt, accs[i], hipDeviceMapHost); + g_devices[g_deviceCnt].init(g_deviceCnt, accs[i], hipDeviceMapHost); g_deviceCnt++; } } diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 784fa9224c..9e0d32e971 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -41,7 +41,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) //Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete: //This matches CUDA stream behavior: - auto istream = new ihipStream_t(device->_device_index, acc.create_view(), device->_stream_id, flags); + auto istream = new ihipStream_t(device->_device_index, acc.create_view(), flags); device->locked_addStream(istream); diff --git a/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp b/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp index fca994cba8..0534d6fbae 100644 --- a/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp +++ b/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp @@ -109,17 +109,17 @@ int main(int argc, char *argv[]) // Serial version, just call once: if (p_tests & 0x1) { - printf ("test 0x1 : serial createThenDestroyStreams(10) \n"); + printf ("\ntest 0x1 : serial createThenDestroyStreams(10) \n"); createThenDestroyStreams(10, 10); }; if (p_tests & 0x2) { - printf ("test 0x2 : serialized multiThread_1(1) \n"); + printf ("\ntest 0x2 : serialized multiThread_pyramid(1) \n"); multiThread_pyramid(true, 10); } if (p_tests & 0x4) { - printf ("test 0x4 : multiThread_pyramid(1) \n"); + printf ("\ntest 0x4 : parallel multiThread_pyramid(1) \n"); multiThread_pyramid(false, 10); } @@ -129,7 +129,7 @@ int main(int argc, char *argv[]) // } if (p_tests & 0x10) { - printf ("test 0x10 : multiThread_tiny(1000) \n"); + printf ("\ntest 0x10 : parallel multiThread_tiny(1000) \n"); multiThread_tiny(false, 1000); } From ef1314a91c44c854f5fc6aa03543250fe879423c Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 26 Mar 2016 12:35:04 -0500 Subject: [PATCH 07/16] Improve debug print messages. - Remove "call-to-call" for hipStreamCreate and hipEventCreate. These now call an internal functions rather than calling through hipStreamCreateWithFalgs and hipEventCreateWithFlags. - Add HIP_INIT_API for more functions so they trace correctly. - Use stream#DEVICE.STREAMID in debug messages via new specialization in tace_helper. [ROCm/clr commit: 82f57ca6105f5dd7a1ed630f96f986d92fd264fe] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 13 +++++++- .../include/hcc_detail/hip_runtime_api.h | 12 ++----- .../hipamd/include/hcc_detail/trace_helper.h | 28 ++++++++++++++-- projects/clr/hipamd/src/hip_event.cpp | 30 ++++++++++++----- projects/clr/hipamd/src/hip_stream.cpp | 33 ++++++++++++++----- projects/clr/hipamd/util/vim/hip.vim | 1 + 6 files changed, 87 insertions(+), 30 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 4e8b513f60..f19e5fd361 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -341,7 +341,7 @@ 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) - ihipDevice_t * getDevice() const; + ihipDevice_t * getDevice() const; StreamMutex & mutex() {return _mutex;}; //--- @@ -375,9 +375,20 @@ private: std::deque _signalPool; // Pool of signals for use by this stream. StreamMutex _mutex; + + friend std::ostream& operator<<(std::ostream& os, const ihipStream_t& s); }; +inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) +{ + os << "stream#"; + os << s._device_index; + os << '.'; + os << s._id; + return os; +} + //---- // Internal event structure: diff --git a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h index 22247af4b5..13716371d7 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h @@ -443,10 +443,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags); * @see hipStreamDestroy * */ -static inline hipError_t hipStreamCreate(hipStream_t *stream) -{ - return hipStreamCreateWithFlags(stream, hipStreamDefault); -} +hipError_t hipStreamCreate(hipStream_t *stream); /** @@ -544,13 +541,10 @@ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags); /** * Create an event * - * @param[in] event Creates an event + * @param[in,out] event Returns the newly created event. * */ -static inline hipError_t hipEventCreate(hipEvent_t* event) -{ - return hipEventCreateWithFlags(event, 0); -} +hipError_t hipEventCreate(hipEvent_t* event); /** diff --git a/projects/clr/hipamd/include/hcc_detail/trace_helper.h b/projects/clr/hipamd/include/hcc_detail/trace_helper.h index 1275016188..db7cc07073 100644 --- a/projects/clr/hipamd/include/hcc_detail/trace_helper.h +++ b/projects/clr/hipamd/include/hcc_detail/trace_helper.h @@ -44,16 +44,38 @@ std::string ToHexString(T v) //--- -// Template overloads for ToString to handle various types: -// Note these use C++11 variadic templates +// Template overloads for ToString to handle specific types + +// This is the default which works for most types: template -std::string ToString(T v) { +std::string ToString(T v) +{ std::ostringstream ss; ss << v; return ss.str(); }; +// hipEvent_t specialization. TODO - maybe add an event ID for debug? +template <> +std::string ToString(hipEvent_t v) +{ + return ToString(&v); +}; + + + +// hipStream_t +template <> +std::string ToString(hipStream_t v) +{ + std::ostringstream ss; + ss << *v; + + return ss.str(); +}; + +// hipMemcpyKind specialization template <> std::string ToString(hipMemcpyKind v) { switch(v) { diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 66a7c7c7fc..74e9ee9eb2 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -25,16 +25,13 @@ THE SOFTWARE. //------------------------------------------------------------------------------------------------- // Events //--- -/** - * @warning : flags must be 0. - */ -hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) -{ - // TODO - support hipEventDefault, hipEventBlockingSync, hipEventDisableTiming - std::call_once(hip_initialized, ihipInit); + +hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) +{ hipError_t e = hipSuccess; + // TODO - support hipEventDefault, hipEventBlockingSync, hipEventDisableTiming if (flags == 0) { ihipEvent_t *eh = event->_handle = new ihipEvent_t(); @@ -47,8 +44,25 @@ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) e = hipErrorInvalidValue; } + return e; +} - return ihipLogStatus(e); +/** + * @warning : flags must be 0. + */ +hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) +{ + HIP_INIT_API(event, flags); + + return ihipLogStatus(ihipEventCreate(event, flags)); +} + + +hipError_t hipEventCreate(hipEvent_t* event) +{ + HIP_INIT_API(event); + + return ihipLogStatus(ihipEventCreate(event, 0)); } diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 9e0d32e971..607da3e3ca 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -28,10 +28,8 @@ THE SOFTWARE. // //--- -hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) +hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags) { - std::call_once(hip_initialized, ihipInit); - ihipDevice_t *device = ihipGetTlsDefaultDevice(); hc::accelerator acc = device->_acc; @@ -48,7 +46,25 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) *stream = istream; tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream); - return ihipLogStatus(hipSuccess); + return hipSuccess; +} + + +//--- +hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) +{ + HIP_INIT_API(stream, flags); + + return ihipLogStatus(ihipStreamCreate(stream, flags)); + +} + +//--- +hipError_t hipStreamCreate(hipStream_t *stream) +{ + HIP_INIT_API(stream); + + return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault)); } @@ -58,8 +74,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) */ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { - - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(stream, event, flags); hipError_t e = hipSuccess; @@ -78,7 +93,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int //--- hipError_t hipStreamSynchronize(hipStream_t stream) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(stream); hipError_t e = hipSuccess; @@ -101,7 +116,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) */ hipError_t hipStreamDestroy(hipStream_t stream) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(stream); hipError_t e = hipSuccess; @@ -130,7 +145,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) //--- hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(stream, flags); if (flags == NULL) { return ihipLogStatus(hipErrorInvalidValue); diff --git a/projects/clr/hipamd/util/vim/hip.vim b/projects/clr/hipamd/util/vim/hip.vim index 73cf71eb73..b530a3f9e5 100644 --- a/projects/clr/hipamd/util/vim/hip.vim +++ b/projects/clr/hipamd/util/vim/hip.vim @@ -150,6 +150,7 @@ syn keyword hipFunctionName hipMemset2D syn keyword hipFunctionName hipMemset3D syn keyword hipFunctionName hipSetDevice syn keyword hipFunctionName hipSetupArgument +syn keyword hipFunctionName hipStreamCreateWithFlags syn keyword hipFunctionName hipStreamCreate syn keyword hipFunctionName hipStreamDestroy syn keyword hipFunctionName hipStreamQuery From 77a720aafc3dce58dddbf68a872ffc25e829d80d Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 28 Mar 2016 04:22:20 -0500 Subject: [PATCH 08/16] Stream thread-safe checkpoint. [ROCm/clr commit: ecd56e1400780ffeb350a7bc74d941a17a0ef147] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 33 +++++++++++++++---- projects/clr/hipamd/src/hip_hcc.cpp | 4 +-- 2 files changed, 29 insertions(+), 8 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index f19e5fd361..38f279a2d3 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -306,10 +306,29 @@ private: }; +template +struct LockedBase { + + // Experts-only interface for explicit locking. + // Most uses should use the lock-accessor. + void lock() { _mutex.lock(); } + void unlock() { _mutex.unlock(); } + + MUTEX_TYPE _mutex; +}; + + +template +class ihipStreamCriticalBase_t : public LockedBase +{ +private: +}; +typedef ihipStreamCriticalBase_t ihipStreamCritical_t; + +typedef LockedAccessor Locked_ihipStreamCritical_t; + -// TODO - move async copy code into stream? Stream->async-copy. -// Add PreCopy / PostCopy to manage locks? // Internal stream structure. class ihipStream_t { public: @@ -342,13 +361,15 @@ 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) ihipDevice_t * getDevice() const; - StreamMutex & mutex() {return _mutex;}; //--- //Member vars - these are set at initialization: SeqNum_t _id; // monotonic sequence ID hc::accelerator_view _av; unsigned _flags; +private: + ihipStreamCritical_t _criticalData; + private: void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal); void waitCopy(ihipSignal_t *signal); @@ -360,6 +381,8 @@ private: //--- unsigned _device_index; + + // Critical Data: ihipCommand_t _last_command_type; // type of the last command // signal of last copy command sent to the stream. @@ -374,7 +397,6 @@ private: SIGSEQNUM _oldest_live_sig_id; // oldest live seq_id, anything < this can be allocated. std::deque _signalPool; // Pool of signals for use by this stream. - StreamMutex _mutex; friend std::ostream& operator<<(std::ostream& os, const ihipStream_t& s); }; @@ -424,7 +446,7 @@ struct ihipEvent_t { // // MUTEX_TYPE is template argument so can easily convert to FakeMutex for performance or stress testing. template -class ihipDeviceCriticalBase_t +class ihipDeviceCriticalBase_t : LockedBase { public: ihipDeviceCriticalBase_t() : _stream_id(0) {}; @@ -440,7 +462,6 @@ public: private: std::list _streams; // streams associated with this device. ihipStream_t::SeqNum_t _stream_id; - MUTEX_TYPE _mutex; }; #if DEVICE_THREAD_SAFE diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index ce136addd9..5662e90cc3 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -288,7 +288,7 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal) // bool ihipStream_t::preKernelCommand() { - _mutex.lock(); // will be unlocked in postKernelCommand + _criticalData.lock();// will be unlocked in postKernelCommand bool addedSync = false; // If switching command types, we need to add a barrier packet to synchronize things. @@ -323,7 +323,7 @@ void ihipStream_t::postKernelCommand(hc::completion_future &kernelFuture) { _last_kernel_future = kernelFuture; - _mutex.unlock(); + _criticalData.unlock(); // paired with lock from preKernelCommand }; From 03b0976e5435bcc82e528f241079db652c01467e Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 28 Mar 2016 05:23:53 -0500 Subject: [PATCH 09/16] fix ihipLogStatus so status arg only evaluated once [ROCm/clr commit: 4f8786c66aca922ef1ccafe9e2d9ca610cb4527c] --- projects/clr/hipamd/include/hcc_detail/hip_hcc.h | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 38f279a2d3..fecefe27b9 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -160,12 +160,13 @@ class ihipDevice_t; #define ihipLogStatus(_hip_status) \ ({\ - tls_lastHipError = _hip_status;\ + hipError_t _local_hip_status = _hip_status; /*local copy so _hip_status only evaluated once*/ \ + tls_lastHipError = _local_hip_status;\ \ if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API) {\ - fprintf(stderr, " %ship-api: %-30s ret=%2d (%s)>>\n" KNRM, (_hip_status == 0) ? API_COLOR:KRED, __func__, _hip_status, ihipErrorString(_hip_status));\ + fprintf(stderr, " %ship-api: %-30s ret=%2d (%s)>>\n" KNRM, (_local_hip_status == 0) ? API_COLOR:KRED, __func__, _local_hip_status, ihipErrorString(_local_hip_status));\ }\ - _hip_status;\ + _local_hip_status;\ }) From 96412bdea24004c794d37c10e898da327aec39b4 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 28 Mar 2016 09:46:40 -0500 Subject: [PATCH 10/16] Stream thread-safe checkpoint. Moving data structures to critical / protected section. [ROCm/clr commit: 6cab7862aea2f29270a8fb2ab94d75287ef97216] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 102 ++++++---- projects/clr/hipamd/src/hip_event.cpp | 5 +- projects/clr/hipamd/src/hip_hcc.cpp | 180 ++++++++++-------- projects/clr/hipamd/src/hip_memory.cpp | 7 +- projects/clr/hipamd/src/hip_stream.cpp | 6 +- 5 files changed, 171 insertions(+), 129 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index fecefe27b9..a64cdba6ef 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -280,6 +280,12 @@ typedef std::mutex StreamMutex; typedef FakeMutex StreamMutex; #endif +#if DEVICE_THREAD_SAFE +typedef std::mutex DeviceMutex; +#else +typedef FakeMutex DeviceMutex; +#warning "Device thread-safe disabled" +#endif // //--- @@ -322,10 +328,46 @@ struct LockedBase { template class ihipStreamCriticalBase_t : public LockedBase { -private: -}; -typedef ihipStreamCriticalBase_t ihipStreamCritical_t; +public: + ihipStreamCriticalBase_t() : + _last_command_type(ihipCommandCopyH2H), + _last_copy_signal(NULL), + _signalCursor(0), + _oldest_live_sig_id(1), + _stream_sig_id(0) + { + _signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1); + }; + ~ihipStreamCriticalBase_t() { + _signalPool.clear(); + } + + ihipStreamCriticalBase_t * mlock() { LockedBase::lock(); return this;}; + + +public: + // Critical Data: + ihipCommand_t _last_command_type; // type of the last command + + // signal of last copy command sent to the stream. + // May be NULL, indicating the previous command has completley finished and future commands don't need to create a dependency. + // Copy can be either H2D or D2H. + ihipSignal_t *_last_copy_signal; + + hc::completion_future _last_kernel_future; // Completion future of last kernel command sent to GPU. + + // Signal pool: + int _signalCursor; + SIGSEQNUM _oldest_live_sig_id; // oldest live seq_id, anything < this can be allocated. + std::deque _signalPool; // Pool of signals for use by this stream. + + + SIGSEQNUM _stream_sig_id; // Monotonically increasing unique signal id. +}; + + +typedef ihipStreamCriticalBase_t ihipStreamCritical_t; typedef LockedAccessor Locked_ihipStreamCritical_t; @@ -339,7 +381,9 @@ typedef uint64_t SeqNum_t ; ~ihipStream_t(); // kind is hipMemcpyKind - void copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind); + void copySync (Locked_ihipStreamCritical_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind); + void locked_copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind); + void copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind); //--- @@ -347,57 +391,46 @@ typedef uint64_t SeqNum_t ; bool preKernelCommand(); void postKernelCommand(hc::completion_future &kernel_future); - int preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType); + int preCopyCommand(Locked_ihipStreamCritical_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType); - void reclaimSignals_ts(SIGSEQNUM sigNum); - void wait(bool assertQueueEmpty=false); + void locked_reclaimSignals(SIGSEQNUM sigNum); + void locked_wait(bool assertQueueEmpty=false); + + // Use this if we already have the stream critical data mutex: + void wait(Locked_ihipStreamCritical_t &crit, bool assertQueueEmpty=false); + SIGSEQNUM locked_lastCopySeqId() {Locked_ihipStreamCritical_t crit(_criticalData); return lastCopySeqId(crit); }; - // Non-threadsafe accessors - must be protected by high-level stream lock: - SIGSEQNUM lastCopySeqId() { return _last_copy_signal ? _last_copy_signal->_sig_id : 0; }; - ihipSignal_t * allocSignal(); + // Non-threadsafe accessors - must be protected by high-level stream lock with accessor passed to function. + SIGSEQNUM lastCopySeqId(Locked_ihipStreamCritical_t &crit) { return crit->_last_copy_signal ? crit->_last_copy_signal->_sig_id : 0; }; + ihipSignal_t * allocSignal(Locked_ihipStreamCritical_t &crit); //-- Non-racy accessors: // These functions access fields set at initialization time and are non-racy (so do not acquire mutex) ihipDevice_t * getDevice() const; + + //--- //Member vars - these are set at initialization: SeqNum_t _id; // monotonic sequence ID hc::accelerator_view _av; unsigned _flags; + private: ihipStreamCritical_t _criticalData; private: void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal); - void waitCopy(ihipSignal_t *signal); + void waitCopy(Locked_ihipStreamCritical_t &crit, ihipSignal_t *signal); // The unsigned return is hipMemcpyKind unsigned resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem); void setCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent); - //--- - - unsigned _device_index; - - // Critical Data: - ihipCommand_t _last_command_type; // type of the last command - - // signal of last copy command sent to the stream. - // May be NULL, indicating the previous command has completley finished and future commands don't need to create a dependency. - // Copy can be either H2D or D2H. - ihipSignal_t *_last_copy_signal; - hc::completion_future _last_kernel_future; // Completion future of last kernel command sent to GPU. - - int _signalCursor; - - SIGSEQNUM _stream_sig_id; // Monotonically increasing unique signal id. - SIGSEQNUM _oldest_live_sig_id; // oldest live seq_id, anything < this can be allocated. - std::deque _signalPool; // Pool of signals for use by this stream. - + unsigned _device_index; // index into the g_device array friend std::ostream& operator<<(std::ostream& os, const ihipStream_t& s); }; @@ -465,13 +498,10 @@ private: ihipStream_t::SeqNum_t _stream_id; }; -#if DEVICE_THREAD_SAFE -typedef ihipDeviceCriticalBase_t ihipDeviceCritical_t; // Use real mutex -#else -#warning "Device thread-safe disabled" -typedef ihipDeviceCriticalBase_t ihipDeviceCritical_t; // Fake mutex, for testing -#endif +// Note Mutex selected based on DeviceMutex +typedef ihipDeviceCriticalBase_t ihipDeviceCritical_t; +// This type is used by functions that need access to the critical device structures. typedef LockedAccessor Locked_ihipDeviceCritical_t; diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 74e9ee9eb2..f3c28e176e 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -91,7 +91,8 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) // Clear timestamps eh->_timestamp = 0; eh->_marker = stream->_av.create_marker(); - eh->_copy_seq_id = stream->lastCopySeqId(); + + eh->_copy_seq_id = stream->locked_lastCopySeqId(); return ihipLogStatus(hipSuccess); } @@ -139,7 +140,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) #else eh->_marker.wait(); #endif - eh->_stream->reclaimSignals_ts(eh->_copy_seq_id); + eh->_stream->locked_reclaimSignals(eh->_copy_seq_id); return ihipLogStatus(hipSuccess); } diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 5662e90cc3..fccec004c9 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -128,40 +128,32 @@ ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, unsig _id(0), // will be set by add function. _av(av), _flags(flags), - _device_index(device_index), - _last_command_type(ihipCommandCopyH2D), - _last_copy_signal(NULL), - _signalCursor(0), - _stream_sig_id(0), - _oldest_live_sig_id(1) + _device_index(device_index) { tprintf(DB_SYNC, " streamCreate: stream=%p\n", this); - _signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1); - }; //--- ihipStream_t::~ihipStream_t() -{ - _signalPool.clear(); - // Hack to catch memory issues, in particular accesses to _acc after it has been destroyed. - //memset (&_av, 0x0, sizeof(hc::accelerator_view&)); +{ } //--- -void ihipStream_t::reclaimSignals_ts(SIGSEQNUM sigNum) +void ihipStream_t::locked_reclaimSignals(SIGSEQNUM sigNum) { + Locked_ihipStreamCritical_t crit(_criticalData); + tprintf(DB_SIGNAL, "reclaim signal #%lu\n", sigNum); - // Mark all signals older and including this one as available for - _oldest_live_sig_id = sigNum+1; + // Mark all signals older and including this one as available for re-allocation. + crit->_oldest_live_sig_id = sigNum+1; } //--- -void ihipStream_t::waitCopy(ihipSignal_t *signal) +void ihipStream_t::waitCopy(Locked_ihipStreamCritical_t &crit, ihipSignal_t *signal) { hsa_signal_wait_acquire(signal->_hsa_signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); @@ -169,32 +161,43 @@ void ihipStream_t::waitCopy(ihipSignal_t *signal) tprintf(DB_SIGNAL, "waitCopy reclaim signal #%lu\n", sigNum); // Mark all signals older and including this one as available for reclaim - if (sigNum > _oldest_live_sig_id) { - _oldest_live_sig_id = sigNum+1; // TODO, +1 here seems dangerous. + if (sigNum > crit->_oldest_live_sig_id) { + crit->_oldest_live_sig_id = sigNum+1; // TODO, +1 here seems dangerous. } } - -//--- //Wait for all kernel and data copy commands in this stream to complete. -void ihipStream_t::wait(bool assertQueueEmpty) +//This signature should be used in routines that already have locked the stream mutex +void ihipStream_t::wait(Locked_ihipStreamCritical_t &crit, bool assertQueueEmpty) { if (! assertQueueEmpty) { tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this); _av.wait(); } - if (_last_copy_signal) { - tprintf (DB_SYNC, "stream %p wait for lastCopy:#%lu...\n", this, _last_copy_signal ? _last_copy_signal->_sig_id: 0x0 ); - this->waitCopy(_last_copy_signal); + if (crit->_last_copy_signal) { + tprintf (DB_SYNC, "stream %p wait for lastCopy:#%lu...\n", this, lastCopySeqId(crit) ); + this->waitCopy(crit, crit->_last_copy_signal); } // Reset the stream to "empty" - next command will not set up an inpute dependency on any older signal. - _last_command_type = ihipCommandCopyH2D; - _last_copy_signal = NULL; + crit->_last_command_type = ihipCommandCopyH2D; + crit->_last_copy_signal = NULL; +} + + +//--- +//Wait for all kernel and data copy commands in this stream to complete. +void ihipStream_t::locked_wait(bool assertQueueEmpty) +{ + Locked_ihipStreamCritical_t crit(_criticalData); + + wait(crit, assertQueueEmpty); + }; + //--- ihipDevice_t * ihipStream_t::getDevice() const { @@ -210,26 +213,26 @@ ihipDevice_t * ihipStream_t::getDevice() const // Allocate a new signal from the signal pool. // Returned signals have value of 0. // Signals are intended for use in this stream and are always reclaimed "in-order". -ihipSignal_t *ihipStream_t::allocSignal() +ihipSignal_t *ihipStream_t::allocSignal(Locked_ihipStreamCritical_t &crit) { - int numToScan = _signalPool.size(); + int numToScan = crit->_signalPool.size(); do { - auto thisCursor = _signalCursor; - if (++_signalCursor == _signalPool.size()) { - _signalCursor = 0; + auto thisCursor = crit->_signalCursor; + if (++crit->_signalCursor == crit->_signalPool.size()) { + crit->_signalCursor = 0; } - if (_signalPool[thisCursor]._sig_id < _oldest_live_sig_id) { - SIGSEQNUM oldSigId = _signalPool[thisCursor]._sig_id; - _signalPool[thisCursor]._index = thisCursor; - _signalPool[thisCursor]._sig_id = ++_stream_sig_id; // allocate it. + if (crit->_signalPool[thisCursor]._sig_id < crit->_oldest_live_sig_id) { + SIGSEQNUM oldSigId = crit->_signalPool[thisCursor]._sig_id; + crit->_signalPool[thisCursor]._index = thisCursor; + crit->_signalPool[thisCursor]._sig_id = ++crit->_stream_sig_id; // allocate it. tprintf(DB_SIGNAL, "allocatSignal #%lu at pos:%i (old sigId:%lu < oldest_live:%lu)\n", - _signalPool[thisCursor]._sig_id, - thisCursor, oldSigId, _oldest_live_sig_id); + crit->_signalPool[thisCursor]._sig_id, + thisCursor, oldSigId, crit->_oldest_live_sig_id); - return &_signalPool[thisCursor]; + return &crit->_signalPool[thisCursor]; } } while (--numToScan) ; @@ -237,13 +240,13 @@ ihipSignal_t *ihipStream_t::allocSignal() assert(numToScan == 0); // Have to grow the pool: - _signalCursor = _signalPool.size(); // set to the beginning of the new entries: - if (_signalCursor > 10000) { - fprintf (stderr, "warning: signal pool size=%d, may indicate runaway number of inflight commands\n", _signalCursor); + crit->_signalCursor = crit->_signalPool.size(); // set to the beginning of the new entries: + if (crit->_signalCursor > 10000) { + fprintf (stderr, "warning: signal pool size=%d, may indicate runaway number of inflight commands\n", crit->_signalCursor); } - _signalPool.resize(_signalPool.size() * 2); - tprintf (DB_SIGNAL, "grow signal pool to %zu entries, cursor=%d\n", _signalPool.size(), _signalCursor); - return allocSignal(); // try again, + crit->_signalPool.resize(crit->_signalPool.size() * 2); + tprintf (DB_SIGNAL, "grow signal pool to %zu entries, cursor=%d\n", crit->_signalPool.size(), crit->_signalCursor); + return allocSignal(crit); // try again, // Should never reach here. assert(0); @@ -288,30 +291,31 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal) // bool ihipStream_t::preKernelCommand() { - _criticalData.lock();// will be unlocked in postKernelCommand + ihipStreamCritical_t *critData =_criticalData.mlock();// will be unlocked in postKernelCommand bool addedSync = false; // If switching command types, we need to add a barrier packet to synchronize things. - if (_last_command_type != ihipCommandKernel) { - if (_last_copy_signal) { + if (critData->_last_command_type != ihipCommandKernel) { + if (critData->_last_copy_signal) { addedSync = true; hsa_queue_t * q = (hsa_queue_t*)_av.get_hsa_queue(); if (HIP_DISABLE_HW_KERNEL_DEP == 0) { - this->enqueueBarrier(q, _last_copy_signal); + this->enqueueBarrier(q, critData->_last_copy_signal); tprintf (DB_SYNC, "stream %p switch %s to %s (barrier pkt inserted with wait on #%lu)\n", - this, ihipCommandName[_last_command_type], ihipCommandName[ihipCommandKernel], _last_copy_signal->_sig_id) + this, ihipCommandName[critData->_last_command_type], ihipCommandName[ihipCommandKernel], critData->_last_copy_signal->_sig_id) } else if (HIP_DISABLE_HW_KERNEL_DEP>0) { tprintf (DB_SYNC, "stream %p switch %s to %s (HOST wait for previous...)\n", - this, ihipCommandName[_last_command_type], ihipCommandName[ihipCommandKernel]); - this->waitCopy(_last_copy_signal); + this, ihipCommandName[critData->_last_command_type], ihipCommandName[ihipCommandKernel]); + assert(0); // Fix/enable next line. TODO + //this->waitCopy(critData, critData->_last_copy_signal); } else if (HIP_DISABLE_HW_KERNEL_DEP==-1) { tprintf (DB_SYNC, "stream %p switch %s to %s (IGNORE dependency)\n", - this, ihipCommandName[_last_command_type], ihipCommandName[ihipCommandKernel]); + this, ihipCommandName[critData->_last_command_type], ihipCommandName[ihipCommandKernel]); } } - _last_command_type = ihipCommandKernel; + critData->_last_command_type = ihipCommandKernel; } return addedSync; @@ -321,7 +325,8 @@ bool ihipStream_t::preKernelCommand() //--- void ihipStream_t::postKernelCommand(hc::completion_future &kernelFuture) { - _last_kernel_future = kernelFuture; + // We locked _criticalData in the preKernelCommand() so OK to access here: + _criticalData._last_kernel_future = kernelFuture; _criticalData.unlock(); // paired with lock from preKernelCommand }; @@ -331,7 +336,7 @@ void ihipStream_t::postKernelCommand(hc::completion_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. -int ihipStream_t::preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType) +int ihipStream_t::preCopyCommand(Locked_ihipStreamCritical_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType) { int needSync = 0; @@ -340,24 +345,24 @@ int ihipStream_t::preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSigna //_mutex.lock(); // will be unlocked in postCopyCommand // If switching command types, we need to add a barrier packet to synchronize things. - if (FORCE_SAMEDIR_COPY_DEP || (_last_command_type != copyType)) { + if (FORCE_SAMEDIR_COPY_DEP || (crit->_last_command_type != copyType)) { - if (_last_command_type == ihipCommandKernel) { + if (crit->_last_command_type == ihipCommandKernel) { tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on prev kernel)\n", - this, ihipCommandName[_last_command_type], ihipCommandName[copyType]); + this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType]); needSync = 1; - hsa_signal_t *hsaSignal = (static_cast (_last_kernel_future.get_native_handle())); + hsa_signal_t *hsaSignal = (static_cast (crit->_last_kernel_future.get_native_handle())); if (hsaSignal) { *waitSignal = * hsaSignal; } else { assert(0); // if NULL signal, and we return 1, hsa_amd_memory_copy_async will fail. Confirm this never happens. } - } else if (_last_copy_signal) { + } else if (crit->_last_copy_signal) { needSync = 1; tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on other copy #%lu)\n", - this, ihipCommandName[_last_command_type], ihipCommandName[copyType], _last_copy_signal->_sig_id); - *waitSignal = _last_copy_signal->_hsa_signal; + this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType], crit->_last_copy_signal->_sig_id); + *waitSignal = crit->_last_copy_signal->_hsa_signal; } if (HIP_DISABLE_HW_COPY_DEP && needSync) { @@ -372,10 +377,10 @@ int ihipStream_t::preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSigna } } - _last_command_type = copyType; + crit->_last_command_type = copyType; } - _last_copy_signal = lastCopy; + crit->_last_copy_signal = lastCopy; return needSync; } @@ -694,7 +699,7 @@ void ihipDevice_t::locked_syncDefaultStream(bool waitOnSelf) if (waitOnSelf || (stream != _default_stream)) { // TODO-hcc - use blocking or active wait here? // TODO-sync - cudaDeviceBlockingSync - stream->wait(); + stream->locked_wait(); } } } @@ -726,7 +731,7 @@ void ihipDevice_t::locked_waitAllStreams() tprintf(DB_SYNC, "waitAllStream\n"); for (auto streamI=l->const_streams().begin(); streamI!=l->const_streams().end(); streamI++) { - (*streamI)->wait(); + (*streamI)->locked_wait(); } } @@ -969,7 +974,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream) // Have to wait for legacy default stream to be empty: if (!(stream->_flags & hipStreamNonBlocking)) { tprintf(DB_SYNC, "stream %p wait default stream\n", stream); - stream->getDevice()->_default_stream->wait(); + stream->getDevice()->_default_stream->locked_wait(); } return stream; @@ -1101,7 +1106,7 @@ void ihipStream_t::setCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_ } -void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind) +void ihipStream_t::copySync(Locked_ihipStreamCritical_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind) { ihipDevice_t *device = this->getDevice(); @@ -1127,7 +1132,7 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign hsa_signal_t depSignal; if ((kind == hipMemcpyHostToDevice) && (!srcTracked)) { - int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyH2D); + int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyH2D); if (HIP_STAGING_BUFFERS) { tprintf(DB_COPY1, "D2H && !dstTracked: staged copy H2D dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); @@ -1138,7 +1143,7 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign } // The copy waits for inputs and then completes before returning so can reset queue to empty: - this->wait(true); + this->wait(crit, true); } else { // TODO - remove, slow path. tprintf(DB_COPY1, "H2D && ! srcTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); @@ -1149,14 +1154,14 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign #endif } } else if ((kind == hipMemcpyDeviceToHost) && (!dstTracked)) { - int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyD2H); + int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyD2H); if (HIP_STAGING_BUFFERS) { tprintf(DB_COPY1, "D2H && !dstTracked: staged copy D2H dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); //printf ("staged-copy- read dep signals\n"); device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); - // The copy waits for inputs and then completes before returning so can reset queue to empty: - this->wait(true); + // The copy completes before returning so can reset queue to empty: + this->wait(crit, true); } else { // TODO - remove, slow path. @@ -1168,7 +1173,7 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign #endif } } else if (kind == hipMemcpyHostToHost) { - int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyH2H); + int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyH2H); if (depSignalCnt) { // host waits before doing host memory copy. @@ -1183,10 +1188,10 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign hsa_agent_t srcAgent, dstAgent; setCopyAgents(kind, &commandType, &srcAgent, &dstAgent); - int depSignalCnt = preCopyCommand(NULL, &depSignal, commandType); + int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, commandType); // Get a completion signal: - ihipSignal_t *ihipSignal = allocSignal(); + ihipSignal_t *ihipSignal = allocSignal(crit); hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal; hsa_signal_store_relaxed(copyCompleteSignal, 1); @@ -1197,7 +1202,7 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign // This is sync copy, so let's wait for copy right here: if (hsa_status == HSA_STATUS_SUCCESS) { - waitCopy(ihipSignal); // wait for copy, and return to pool. + waitCopy(crit, ihipSignal); // wait for copy, and return to pool. } else { throw ihipException(hipErrorInvalidValue); } @@ -1205,10 +1210,19 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign } +// Sync copy that acquires lock: +void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind) +{ + Locked_ihipStreamCritical_t crit (_criticalData); + copySync(crit, dst, src, sizeBytes, kind); +} + void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind) { + Locked_ihipStreamCritical_t crit(_criticalData); + ihipDevice_t *device = this->getDevice(); if (device == NULL) { @@ -1222,7 +1236,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig /* As this is a CPU op, we need to wait until all the commands in current stream are finished. */ - this->wait(); + this->wait(crit); memcpy(dst, src, sizeBytes); @@ -1249,8 +1263,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig } - - ihipSignal_t *ihip_signal = allocSignal(); + ihipSignal_t *ihip_signal = allocSignal(crit); hsa_signal_store_relaxed(ihip_signal->_hsa_signal, 1); @@ -1261,7 +1274,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig setCopyAgents(kind, &commandType, &srcAgent, &dstAgent); hsa_signal_t depSignal; - int depSignalCnt = preCopyCommand(ihip_signal, &depSignal, commandType); + int depSignalCnt = preCopyCommand(crit, ihip_signal, &depSignal, commandType); tprintf (DB_SYNC, " copy-async, waitFor=%lu completion=#%lu(%lu)\n", depSignalCnt? depSignal.handle:0x0, ihip_signal->_sig_id, ihip_signal->_hsa_signal.handle); @@ -1271,7 +1284,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig if (hsa_status == HSA_STATUS_SUCCESS) { if (HIP_LAUNCH_BLOCKING) { tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes); - this->wait(); + this->wait(crit); } } else { // This path can be hit if src or dst point to unpinned host memory. @@ -1279,7 +1292,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig throw ihipException(hipErrorInvalidValue); } } else { - copySync(dst, src, sizeBytes, kind); + copySync(crit, dst, src, sizeBytes, kind); } } } @@ -1332,6 +1345,7 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a // Caps convention _ or camelCase // if { } // Should use ihip* data structures inside code rather than app-facing hip. For example, use ihipDevice_t (rather than hipDevice_t), ihipStream_t (rather than hipStream_t). +// locked_ // TODO - describe MT strategy // //// TODO - add identifier numbers for streams and devices to help with debugging. diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index c633d23611..12ad9451da 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -307,17 +307,14 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind hipError_t e = hipSuccess; try { - stream->copySync(dst, src, sizeBytes, kind); + + stream->locked_copySync(dst, src, sizeBytes, kind); } catch (ihipException ex) { e = ex._code; } - if (HIP_LAUNCH_BLOCKING) { - tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpy\n"); - stream->wait(); - } return ihipLogStatus(e); } diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 607da3e3ca..d62abc49e2 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -82,7 +82,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int // TODO-hcc Convert to use create_blocking_marker(...) functionality. // Currently we have a super-conservative version of this - block on host, and drain the queue. // This should create a barrier packet in the target queue. - stream->wait(); + stream->locked_wait(); e = hipSuccess; } @@ -101,7 +101,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) ihipDevice_t *device = ihipGetTlsDefaultDevice(); device->locked_syncDefaultStream(true/*waitOnSelf*/); } else { - stream->wait(); + stream->locked_wait(); e = hipSuccess; } @@ -125,7 +125,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) ihipDevice_t *device = ihipGetTlsDefaultDevice(); device->locked_syncDefaultStream(true/*waitOnSelf*/); } else { - stream->wait(); + stream->locked_wait(); e = hipSuccess; } From 34f21e1343080a57edf004a08f32bee9871e39ed Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 28 Mar 2016 21:41:47 -0500 Subject: [PATCH 11/16] Tweak thread-safe implementation. introduce LockedAccessor option so destructor does not unlock. Allows locks to exist across function boundaries, required for hipLaunchKernel macro which has several unusual requirements. (including C comppatibility, must use variadic macro, more). [ROCm/clr commit: 1b2ab173c1e15d30c1ef71b4ca467b948c92bcd8] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 43 +++++++---- projects/clr/hipamd/include/hip_runtime_api.h | 2 +- projects/clr/hipamd/src/hip_hcc.cpp | 73 ++++++++++--------- projects/clr/hipamd/src/hip_memory.cpp | 4 +- 4 files changed, 67 insertions(+), 55 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index a64cdba6ef..ee510efe47 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -295,14 +295,24 @@ template class LockedAccessor { public: - LockedAccessor(T &criticalData) : _criticalData(&criticalData) + LockedAccessor(T &criticalData, bool autoUnlock=true) : + _criticalData(&criticalData), + _autoUnlock(autoUnlock) + { _criticalData->_mutex.lock(); }; ~LockedAccessor() { - _criticalData->_mutex.unlock(); + if (_autoUnlock) { + _criticalData->_mutex.unlock(); + } + } + + void unlock() + { + _criticalData->_mutex.unlock(); } // Syntactic sugar so -> can be used to get the underlying type. @@ -310,6 +320,7 @@ public: private: T *_criticalData; + bool _autoUnlock; }; @@ -368,7 +379,7 @@ public: typedef ihipStreamCriticalBase_t ihipStreamCritical_t; -typedef LockedAccessor Locked_ihipStreamCritical_t; +typedef LockedAccessor LockedAccessor_StreamCrit_t; @@ -381,30 +392,30 @@ typedef uint64_t SeqNum_t ; ~ihipStream_t(); // kind is hipMemcpyKind - void copySync (Locked_ihipStreamCritical_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind); + void copySync (LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind); void locked_copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind); void copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind); //--- // Thread-safe accessors - these acquire / release mutex: - bool preKernelCommand(); - void postKernelCommand(hc::completion_future &kernel_future); + bool lockopen_preKernelCommand(); + void lockclose_postKernelCommand(hc::completion_future &kernel_future); - int preCopyCommand(Locked_ihipStreamCritical_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType); + int preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType); void locked_reclaimSignals(SIGSEQNUM sigNum); void locked_wait(bool assertQueueEmpty=false); + SIGSEQNUM locked_lastCopySeqId() {LockedAccessor_StreamCrit_t crit(_criticalData); return lastCopySeqId(crit); }; // Use this if we already have the stream critical data mutex: - void wait(Locked_ihipStreamCritical_t &crit, bool assertQueueEmpty=false); + void wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty=false); - SIGSEQNUM locked_lastCopySeqId() {Locked_ihipStreamCritical_t crit(_criticalData); return lastCopySeqId(crit); }; // Non-threadsafe accessors - must be protected by high-level stream lock with accessor passed to function. - SIGSEQNUM lastCopySeqId(Locked_ihipStreamCritical_t &crit) { return crit->_last_copy_signal ? crit->_last_copy_signal->_sig_id : 0; }; - ihipSignal_t * allocSignal(Locked_ihipStreamCritical_t &crit); + SIGSEQNUM lastCopySeqId (LockedAccessor_StreamCrit_t &crit) { return crit->_last_copy_signal ? crit->_last_copy_signal->_sig_id : 0; }; + ihipSignal_t * allocSignal (LockedAccessor_StreamCrit_t &crit); //-- Non-racy accessors: @@ -412,19 +423,19 @@ typedef uint64_t SeqNum_t ; ihipDevice_t * getDevice() const; - +public: //--- - //Member vars - these are set at initialization: + //Public member vars - these are set at initialization and never change: SeqNum_t _id; // monotonic sequence ID hc::accelerator_view _av; unsigned _flags; -private: +private: // Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t ihipStreamCritical_t _criticalData; private: void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal); - void waitCopy(Locked_ihipStreamCritical_t &crit, ihipSignal_t *signal); + void waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal); // The unsigned return is hipMemcpyKind unsigned resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem); @@ -502,7 +513,7 @@ private: typedef ihipDeviceCriticalBase_t ihipDeviceCritical_t; // This type is used by functions that need access to the critical device structures. -typedef LockedAccessor Locked_ihipDeviceCritical_t; +typedef LockedAccessor LockedAccessor_DeviceCrit_t; diff --git a/projects/clr/hipamd/include/hip_runtime_api.h b/projects/clr/hipamd/include/hip_runtime_api.h index cd0d841d97..3a0a4b399a 100644 --- a/projects/clr/hipamd/include/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip_runtime_api.h @@ -222,7 +222,7 @@ static inline hipError_t hipMalloc ( T** devPtr, size_t size) // Provide an override to automatically typecast the pointer type from void**, and also provide a default for the flags. template -static inline hipError_t hipHostMalloc( T** ptr, size_t size, unsigned int flags = 0) +static inline hipError_t hipHostMalloc( T** ptr, size_t size, unsigned int flags = hipHostMallocDefault) { return hipHostMalloc((void**)ptr, size, flags); } diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index fccec004c9..cf1cb3953a 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -144,7 +144,7 @@ ihipStream_t::~ihipStream_t() //--- void ihipStream_t::locked_reclaimSignals(SIGSEQNUM sigNum) { - Locked_ihipStreamCritical_t crit(_criticalData); + LockedAccessor_StreamCrit_t crit(_criticalData); tprintf(DB_SIGNAL, "reclaim signal #%lu\n", sigNum); // Mark all signals older and including this one as available for re-allocation. @@ -153,7 +153,7 @@ void ihipStream_t::locked_reclaimSignals(SIGSEQNUM sigNum) //--- -void ihipStream_t::waitCopy(Locked_ihipStreamCritical_t &crit, ihipSignal_t *signal) +void ihipStream_t::waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal) { hsa_signal_wait_acquire(signal->_hsa_signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); @@ -169,7 +169,7 @@ void ihipStream_t::waitCopy(Locked_ihipStreamCritical_t &crit, ihipSignal_t *sig //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(Locked_ihipStreamCritical_t &crit, bool assertQueueEmpty) +void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty) { if (! assertQueueEmpty) { tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this); @@ -190,7 +190,7 @@ void ihipStream_t::wait(Locked_ihipStreamCritical_t &crit, bool assertQueueEmpty //Wait for all kernel and data copy commands in this stream to complete. void ihipStream_t::locked_wait(bool assertQueueEmpty) { - Locked_ihipStreamCritical_t crit(_criticalData); + LockedAccessor_StreamCrit_t crit(_criticalData); wait(crit, assertQueueEmpty); @@ -213,7 +213,7 @@ ihipDevice_t * ihipStream_t::getDevice() const // Allocate a new signal from the signal pool. // Returned signals have value of 0. // Signals are intended for use in this stream and are always reclaimed "in-order". -ihipSignal_t *ihipStream_t::allocSignal(Locked_ihipStreamCritical_t &crit) +ihipSignal_t *ihipStream_t::allocSignal(LockedAccessor_StreamCrit_t &crit) { int numToScan = crit->_signalPool.size(); do { @@ -289,33 +289,32 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal) //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). // -bool ihipStream_t::preKernelCommand() +bool ihipStream_t::lockopen_preKernelCommand() { - ihipStreamCritical_t *critData =_criticalData.mlock();// will be unlocked in postKernelCommand + LockedAccessor_StreamCrit_t crit(_criticalData, false/*no unlock at destruction*/); bool addedSync = false; // If switching command types, we need to add a barrier packet to synchronize things. - if (critData->_last_command_type != ihipCommandKernel) { - if (critData->_last_copy_signal) { + if (crit->_last_command_type != ihipCommandKernel) { + if (crit->_last_copy_signal) { addedSync = true; hsa_queue_t * q = (hsa_queue_t*)_av.get_hsa_queue(); if (HIP_DISABLE_HW_KERNEL_DEP == 0) { - this->enqueueBarrier(q, critData->_last_copy_signal); + this->enqueueBarrier(q, crit->_last_copy_signal); tprintf (DB_SYNC, "stream %p switch %s to %s (barrier pkt inserted with wait on #%lu)\n", - this, ihipCommandName[critData->_last_command_type], ihipCommandName[ihipCommandKernel], critData->_last_copy_signal->_sig_id) + this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel], crit->_last_copy_signal->_sig_id) } else if (HIP_DISABLE_HW_KERNEL_DEP>0) { tprintf (DB_SYNC, "stream %p switch %s to %s (HOST wait for previous...)\n", - this, ihipCommandName[critData->_last_command_type], ihipCommandName[ihipCommandKernel]); - assert(0); // Fix/enable next line. TODO - //this->waitCopy(critData, critData->_last_copy_signal); + this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel]); + this->waitCopy(crit, crit->_last_copy_signal); } else if (HIP_DISABLE_HW_KERNEL_DEP==-1) { tprintf (DB_SYNC, "stream %p switch %s to %s (IGNORE dependency)\n", - this, ihipCommandName[critData->_last_command_type], ihipCommandName[ihipCommandKernel]); + this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel]); } } - critData->_last_command_type = ihipCommandKernel; + crit->_last_command_type = ihipCommandKernel; } return addedSync; @@ -323,12 +322,13 @@ bool ihipStream_t::preKernelCommand() //--- -void ihipStream_t::postKernelCommand(hc::completion_future &kernelFuture) +// Must be called after kernel finishes, this releases the lock on the stream so other commands can submit. +void ihipStream_t::lockclose_postKernelCommand(hc::completion_future &kernelFuture) { - // We locked _criticalData in the preKernelCommand() so OK to access here: + // We locked _criticalData in the lockopen_preKernelCommand() so OK to access here: _criticalData._last_kernel_future = kernelFuture; - _criticalData.unlock(); // paired with lock from preKernelCommand + _criticalData.unlock(); // paired with lock from lockopen_preKernelCommand. }; @@ -336,7 +336,7 @@ void ihipStream_t::postKernelCommand(hc::completion_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. -int ihipStream_t::preCopyCommand(Locked_ihipStreamCritical_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType) +int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType) { int needSync = 0; @@ -395,10 +395,10 @@ int ihipStream_t::preCopyCommand(Locked_ihipStreamCritical_t &crit, ihipSignal_t void ihipDevice_t::locked_reset() { // Obtain mutex access to the device critical data, release by destructor - Locked_ihipDeviceCritical_t l(_criticalData); + LockedAccessor_DeviceCrit_t crit(_criticalData); // Reset and remove streams: - l->streams().clear(); + crit->streams().clear(); // Reset and release all memory stored in the tracker: am_memtracker_reset(_acc); @@ -685,11 +685,11 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) // If waitOnSelf is set, this additionally waits for the default stream to empty. void ihipDevice_t::locked_syncDefaultStream(bool waitOnSelf) { - Locked_ihipDeviceCritical_t l(_criticalData); + LockedAccessor_DeviceCrit_t crit(_criticalData); tprintf(DB_SYNC, "syncDefaultStream\n"); - for (auto streamI=l->const_streams().begin(); streamI!=l->const_streams().end(); streamI++) { + for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) { ihipStream_t *stream = *streamI; // Don't wait for streams that have "opted-out" of syncing with NULL stream. @@ -708,18 +708,18 @@ void ihipDevice_t::locked_syncDefaultStream(bool waitOnSelf) //--- void ihipDevice_t::locked_addStream(ihipStream_t *s) { - Locked_ihipDeviceCritical_t l(_criticalData); + LockedAccessor_DeviceCrit_t crit(_criticalData); - l->streams().push_back(s); - s->_id = l->incStreamId(); + crit->streams().push_back(s); + s->_id = crit->incStreamId(); } //--- void ihipDevice_t::locked_removeStream(ihipStream_t *s) { - Locked_ihipDeviceCritical_t l(_criticalData); + LockedAccessor_DeviceCrit_t crit(_criticalData); - l->streams().remove(s); + crit->streams().remove(s); } @@ -727,10 +727,10 @@ void ihipDevice_t::locked_removeStream(ihipStream_t *s) //Heavyweight synchronization that waits on all streams, ignoring hipStreamNonBlocking flag. void ihipDevice_t::locked_waitAllStreams() { - Locked_ihipDeviceCritical_t l(_criticalData); + LockedAccessor_DeviceCrit_t crit(_criticalData); tprintf(DB_SYNC, "waitAllStream\n"); - for (auto streamI=l->const_streams().begin(); streamI!=l->const_streams().end(); streamI++) { + for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) { (*streamI)->locked_wait(); } } @@ -990,7 +990,8 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av) std::call_once(hip_initialized, ihipInit); stream = ihipSyncAndResolveStream(stream); - stream->preKernelCommand(); + + stream->lockopen_preKernelCommand(); *av = &stream->_av; @@ -1002,7 +1003,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av) //Called after kernel finishes execution. void ihipPostLaunchKernel(hipStream_t stream, hc::completion_future &kernelFuture) { - stream->postKernelCommand(kernelFuture); + stream->lockclose_postKernelCommand(kernelFuture); if (HIP_LAUNCH_BLOCKING) { tprintf(DB_SYNC, " stream:%p LAUNCH_BLOCKING for kernel completion\n", stream); } @@ -1106,7 +1107,7 @@ void ihipStream_t::setCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_ } -void ihipStream_t::copySync(Locked_ihipStreamCritical_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind) +void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind) { ihipDevice_t *device = this->getDevice(); @@ -1213,7 +1214,7 @@ void ihipStream_t::copySync(Locked_ihipStreamCritical_t &crit, void* dst, const // Sync copy that acquires lock: void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind) { - Locked_ihipStreamCritical_t crit (_criticalData); + LockedAccessor_StreamCrit_t crit (_criticalData); copySync(crit, dst, src, sizeBytes, kind); } @@ -1221,7 +1222,7 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind) { - Locked_ihipStreamCritical_t crit(_criticalData); + LockedAccessor_StreamCrit_t crit(_criticalData); ihipDevice_t *device = this->getDevice(); diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 12ad9451da..4d75a1d7a1 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -363,9 +363,9 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s hipError_t e = hipSuccess; stream = ihipSyncAndResolveStream(stream); - stream->preKernelCommand(); if (stream) { + stream->lockopen_preKernelCommand(); hc::completion_future cf ; @@ -389,7 +389,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s } } - stream->postKernelCommand(cf); + stream->lockclose_postKernelCommand(cf); if (HIP_LAUNCH_BLOCKING) { From 829fee518f24893af069c628d0c3f27e77b70bc5 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Tue, 29 Mar 2016 17:26:44 -0500 Subject: [PATCH 12/16] include codexl marker path, if found [ROCm/clr commit: 57783aad2de45add6a260415059ad14bcf4ecc2e] --- projects/clr/hipamd/bin/hipcc | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index 44b1e85a33..3ef1ad1a8a 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/projects/clr/hipamd/bin/hipcc @@ -79,6 +79,11 @@ if ($HIP_PLATFORM eq "hcc") { $HIPLDFLAGS .= " -L$HSA_PATH/lib -lhsa-runtime64 -lhc_am"; # Add trace marker library: + # TODO - once we cleanly separate the HIP API headers from HIP library headers this logic should move to CMakebuild option - apps do not need to see the marker library. + $marker_inc_path = "$marker_path/include"; + if (-e $marker_inc_path) { + $HIPCXXFLAGS .= " -I$marker_inc_path"; + } $marker_lib_path = "$marker_path/bin/x86_64"; if (-e $marker_lib_path) { From 36f2f683db25348d1a72dc4ca5e8d258e80cdcfe Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Tue, 29 Mar 2016 17:27:30 -0500 Subject: [PATCH 13/16] Add runtime switch to control HIP_ATP_MARKER Only generate the function strings if requested at compile-time && runtime. [ROCm/clr commit: e22925be224dd43037f8594de4fa43676a736dbd] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 20 +++++++++++-------- projects/clr/hipamd/src/hip_hcc.cpp | 11 ++++++++-- 2 files changed, 21 insertions(+), 10 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 3c06295eaa..2cca8940e6 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -47,7 +47,9 @@ extern const int release; extern int HIP_LAUNCH_BLOCKING; extern int HIP_PRINT_ENV; +extern int HIP_ATP_MARKER; extern int HIP_TRACE_API; +extern int HIP_ATP; extern int HIP_DB; extern int HIP_STAGING_SIZE; /* size of staging buffers, in KB */ extern int HIP_STAGING_BUFFERS; // TODO - remove, two buffers should be enough. @@ -114,8 +116,8 @@ class ihipDevice_t; // Compile code that generates trace markers for CodeXL ATP at HIP function begin/end. // ATP is standard CodeXL format that includes timestamps for kernels, HSA RT APIs, and HIP APIs. -#ifndef COMPILE_TRACE_MARKER -#define COMPILE_TRACE_MARKER 0 +#ifndef COMPILE_HIP_ATP_MARKER +#define COMPILE_HIP_ATP_MARKER 0 #endif @@ -126,7 +128,7 @@ class ihipDevice_t; // Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary. // TODO - currently we print the trace message at the beginning. if we waited, we could also include return codes, and any values returned // through ptr-to-args (ie the pointers allocated by hipMalloc). -#if COMPILE_TRACE_MARKER +#if COMPILE_HIP_ATP_MARKER #include "AMDTActivityLogger.h" #define SCOPED_MARKER(markerName,group,userString) amdtScopedMarker(markerName, group, userString) #else @@ -135,14 +137,16 @@ class ihipDevice_t; #endif -#if COMPILE_TRACE_MARKER || (COMPILE_HIP_TRACE_API & 0x1) +#if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1) #define API_TRACE(...)\ {\ - std::string s = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')';\ - if (COMPILE_HIP_DB && HIP_TRACE_API) {\ - fprintf (stderr, API_COLOR "< Date: Tue, 29 Mar 2016 17:28:27 -0500 Subject: [PATCH 14/16] Inline some new string functions. [ROCm/clr commit: e7b15a53da6d3aeb27cfd0e170fc9c5ce63044ec] --- projects/clr/hipamd/include/hcc_detail/trace_helper.h | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/trace_helper.h b/projects/clr/hipamd/include/hcc_detail/trace_helper.h index 8ba877f7ec..13af1eab26 100644 --- a/projects/clr/hipamd/include/hcc_detail/trace_helper.h +++ b/projects/clr/hipamd/include/hcc_detail/trace_helper.h @@ -58,7 +58,7 @@ inline std::string ToString(T v) // hipEvent_t specialization. TODO - maybe add an event ID for debug? template <> -std::string ToString(hipEvent_t v) +inline std::string ToString(hipEvent_t v) { return ToString(&v); }; @@ -67,10 +67,14 @@ std::string ToString(hipEvent_t v) // hipStream_t template <> -std::string ToString(hipStream_t v) +inline std::string ToString(hipStream_t v) { std::ostringstream ss; - ss << *v; + if (v == NULL) { + ss << "stream:"; + } else { + ss << *v; + } return ss.str(); }; From 875a8a2c84fe75ce5ab3712d1d9dfdd6d43e6327 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Tue, 29 Mar 2016 17:29:31 -0500 Subject: [PATCH 15/16] Test improvements - partition hipThreadSafeDevice into smaller pieces. - Add debug to hipMultiThreadStream. - print more precision when mismatch detected. - enable more tests in CMakeFiles.txt. [ROCm/clr commit: 98a766ba75b26754610f22bf792487fc17f9b606] --- projects/clr/hipamd/tests/src/CMakeLists.txt | 25 +++++++++++++------ .../tests/src/hipMultiThreadStreams1.cpp | 18 +++++++------ .../hipamd/tests/src/hipThreadSafeDevice.cpp | 15 +++++------ projects/clr/hipamd/tests/src/test_common.h | 6 ++++- 4 files changed, 40 insertions(+), 24 deletions(-) diff --git a/projects/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt index 84a19815be..d47728759f 100644 --- a/projects/clr/hipamd/tests/src/CMakeLists.txt +++ b/projects/clr/hipamd/tests/src/CMakeLists.txt @@ -8,6 +8,8 @@ include_directories( ${PROJECT_SOURCE_DIR}/include ) set (HIP_Unit_Test_VERSION_MAJOR 1) set (HIP_Unit_Test_VERSION_MINOR 0) +set (HIP_BUILD_LOCAL 0) + set(HIP_PATH $ENV{HIP_PATH}) if (NOT DEFINED HIP_PATH) set (HIP_PATH ../..) @@ -42,8 +44,11 @@ if (${HIP_PLATFORM} STREQUAL "hcc") # This will create a subdir "hip_hcc" in the test build directory # Any changes to hip_hcc source will be detected and force the library and then the tests to be rebuilt. - add_subdirectory(${HIP_PATH} build.hip_hcc) - link_directories(${CMAKE_CURRENT_BINARY_DIR}/build.hip_hcc) # search the local hip_hcc for libhip_hcc.a + if (${HIP_BUILD_LOCAL}) + add_subdirectory(${HIP_PATH} build.hip_hcc) + #link_directories(${CMAKE_CURRENT_BINARY_DIR}/build.hip_hcc) # search the local hip_hcc for libhip_hcc.a + set (CMAKE_CXX_FLAGS --hipcc_explicit_lib) + endif() elseif (${HIP_PLATFORM} STREQUAL "nvcc") @@ -62,7 +67,6 @@ endif() set (HIPCC ${HIP_PATH}/bin/hipcc) set (CMAKE_CXX_COMPILER ${HIPCC}) -#set (CMAKE_CXX_FLAGS --hipcc_explicit_lib) add_library(test_common OBJECT test_common.cpp ) @@ -72,7 +76,9 @@ add_library(test_common OBJECT test_common.cpp ) macro (make_hip_executable exe cpp) if (${HIP_PLATFORM} STREQUAL "hcc") add_executable (${exe} ${cpp} ${ARGN} $ ) - target_link_libraries(${exe} hip_hcc) + if (${HIP_BUILD_LOCAL}) + target_link_libraries(${exe} hip_hcc) + endif() else() add_executable (${exe} ${cpp} ${ARGN} $ ) endif() @@ -142,7 +148,7 @@ make_hip_executable (hipMathFunctionsHost hipMathFunctions.cpp hipSinglePrecisio make_hip_executable (hipMathFunctionsDevice hipMathFunctions.cpp hipSinglePrecisionMathDevice.cpp hipDoublePrecisionMathDevice.cpp) make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrinsics.cpp hipDoublePrecisionIntrinsics.cpp hipIntegerIntrinsics.cpp) #TODO - re-enable. This uses the pointer add feature. -#make_hip_executable (hipPointerAttrib hipPointerAttrib.cpp) +make_hip_executable (hipPointerAttrib hipPointerAttrib.cpp) make_hip_executable (hipMultiThreadStreams1 hipMultiThreadStreams1.cpp) make_hip_executable (hipMultiThreadStreams2 hipMultiThreadStreams2.cpp) make_hip_executable (hipHostAlloc hipHostAlloc.cpp) @@ -173,8 +179,9 @@ make_test(hipGridLaunch " " ) make_test(hipEnvVarDriver " " ) #TODO -reenable #make_test(hipPointerAttrib " " ) -#make_test(hipMultiThreadStreams1 " " ) -#make_test(hipMultiThreadStreams2 " " ) +#make_test(hipMultiThreadStreams1 " " ) Fails if 0x3 specified, passes otherwise. + +make_test(hipMultiThreadStreams2 " " ) make_test(hipMemcpy_simple " " ) make_named_test(hipMemcpy "hipMemcpy-modes" --tests 0x1 ) make_named_test(hipMemcpy "hipMemcpy-size" --tests 0x6 ) @@ -198,6 +205,8 @@ make_test(hipFuncSetDeviceFlags " ") make_test(hipFuncGetDevice " ") make_test(hipFuncSetDevice " ") make_test(hipFuncDeviceSynchronize " ") -make_test (hipThreadSafeDevice " ") +make_named_test (hipThreadSafeDevice "hipThreadSafeDevice-serial" --tests 0x1) +make_named_test (hipThreadSafeDevice "hipThreadSafeDevice-pyramid" --tests 0x4) +make_named_test (hipThreadSafeDevice "hipThreadSafeDevice-nearzero" --tests 0x10) make_hipify_test(specialFunc.cu ) diff --git a/projects/clr/hipamd/tests/src/hipMultiThreadStreams1.cpp b/projects/clr/hipamd/tests/src/hipMultiThreadStreams1.cpp index a3dd94e077..2dc8cb27e7 100644 --- a/projects/clr/hipamd/tests/src/hipMultiThreadStreams1.cpp +++ b/projects/clr/hipamd/tests/src/hipMultiThreadStreams1.cpp @@ -84,11 +84,11 @@ void test_multiThread_1(std::string testName, hipStream_t stream0, hipStream_t s std::cout << testName << std::endl; // Test 2 threads operating on same stream: - std::thread t1 (simpleVectorCopy, 2000000/*mb*/, 1000, stream0); + std::thread t1 (simpleVectorCopy, 2000000/*mb*/, 100/*iters*/, stream0); if (serialize) { t1.join(); } - std::thread t2 (simpleVectorCopy, 2000000/*mb*/, 1000, stream1); + std::thread t2 (simpleVectorCopy, 2000000/*mb*/, 100/*iters*/, stream1); if (serialize) { t2.join(); } @@ -119,19 +119,21 @@ int main(int argc, char *argv[]) simpleVectorCopy (2000000/*mb*/, 10/*iters*/, stream); simpleVectorCopy (2000000/*mb*/, 10/*iters*/, stream); - //HIPCHECK(hipStreamDestroy(stream)); + HIPCHECK(hipStreamDestroy(stream)); } - if (p_tests & 0x2) { - hipStream_t stream0, stream1; - HIPCHECK (hipStreamCreate(&stream0)); - HIPCHECK (hipStreamCreate(&stream1)); + hipStream_t stream0, stream1; + HIPCHECK (hipStreamCreate(&stream0)); + HIPCHECK (hipStreamCreate(&stream1)); + if (p_tests & 0x2) { // Easy tests to verify the test works - these don't allow overlap between the threads: test_multiThread_1 ("Multithread NULL with serialized", NULL, NULL, true); - test_multiThread_1 ("Multithread with serialized", stream0, stream1, true); + test_multiThread_1 ("Multithread two streams serialized", stream0, stream1, true); + } + if (p_tests & 0x4) { test_multiThread_1 ("Multithread with NULL stream", NULL, NULL, false); test_multiThread_1 ("Multithread with two streams", stream0, stream1, false); test_multiThread_1 ("Multithread with one stream", stream0, stream0, false); diff --git a/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp b/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp index 0534d6fbae..a1f64aceb3 100644 --- a/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp +++ b/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp @@ -74,7 +74,7 @@ void multiThread_pyramid(bool serialize, int iters) // Create 3 streams, all creating and destroying streams on the same device. // Try to keep number of streams near zero, to cause problems. -void multiThread_tiny(bool serialize, int iters) +void multiThread_nearzero(bool serialize, int iters) { printf ("%s creating %d streams x 3 threads\n", __func__, iters); std::thread t1 (createThenDestroyStreams, iters, 1); @@ -96,9 +96,9 @@ void multiThread_tiny(bool serialize, int iters) } if (!serialize) { - t1.join(); - t2.join(); - t3.join(); + t1.join(); printf ("t1 done\n"); + t2.join(); printf ("t2 done\n"); + t3.join(); printf ("t3 done\n"); } } @@ -113,7 +113,8 @@ int main(int argc, char *argv[]) createThenDestroyStreams(10, 10); }; - if (p_tests & 0x2) { + /*disable, this takess a while and if the next one works then no need to run serial*/ + if (1 && (p_tests & 0x2)) { printf ("\ntest 0x2 : serialized multiThread_pyramid(1) \n"); multiThread_pyramid(true, 10); } @@ -129,8 +130,8 @@ int main(int argc, char *argv[]) // } if (p_tests & 0x10) { - printf ("\ntest 0x10 : parallel multiThread_tiny(1000) \n"); - multiThread_tiny(false, 1000); + printf ("\ntest 0x10 : parallel multiThread_nearzero(1000) \n"); + multiThread_nearzero(false, 1000); } passed(); diff --git a/projects/clr/hipamd/tests/src/test_common.h b/projects/clr/hipamd/tests/src/test_common.h index d92692c978..ce85b3898a 100644 --- a/projects/clr/hipamd/tests/src/test_common.h +++ b/projects/clr/hipamd/tests/src/test_common.h @@ -18,6 +18,7 @@ THE SOFTWARE. */ #include +#include #include #include @@ -232,7 +233,10 @@ void checkVectorADD(T* A_h, T* B_h, T* result_H, size_t N, bool expectMatch=true } mismatchCount++; if ((mismatchCount <= mismatchesToPrint) && expectMatch) { - std::cout << "At " << i << " Computed:" << result_H[i] << ", expected:" << expected << std::endl; + std::cout << std::fixed << std::setprecision(32); + std::cout << "At " << i << std::endl; + std::cout << " Computed:" << result_H[i] << std::endl; + std::cout << " Expected:" << expected << std::endl; } } } From b8f2f4c45ce8cfe3df554df125b3228e5dd79cfc Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Tue, 29 Mar 2016 17:33:29 -0500 Subject: [PATCH 16/16] rename to hipMultiThreadDevice for consistency [ROCm/clr commit: 3b1cd2d626f8c4e8d004274f8942a8e59a28a3a0] --- projects/clr/hipamd/tests/src/CMakeLists.txt | 8 ++++---- .../{hipThreadSafeDevice.cpp => hipMultiThreadDevice.cpp} | 0 2 files changed, 4 insertions(+), 4 deletions(-) rename projects/clr/hipamd/tests/src/{hipThreadSafeDevice.cpp => hipMultiThreadDevice.cpp} (100%) diff --git a/projects/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt index d47728759f..7e5974c5c3 100644 --- a/projects/clr/hipamd/tests/src/CMakeLists.txt +++ b/projects/clr/hipamd/tests/src/CMakeLists.txt @@ -162,7 +162,7 @@ make_hip_executable (hipFuncGetDevice hipFuncGetDevice.cpp) make_hip_executable (hipFuncSetDevice hipFuncSetDevice.cpp) make_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp) -make_hip_executable (hipThreadSafeDevice hipThreadSafeDevice.cpp) +make_hip_executable (hipMultiThreadDevice hipMultiThreadDevice.cpp) make_test(hip_ballot " " ) make_test(hip_anyall " " ) @@ -205,8 +205,8 @@ make_test(hipFuncSetDeviceFlags " ") make_test(hipFuncGetDevice " ") make_test(hipFuncSetDevice " ") make_test(hipFuncDeviceSynchronize " ") -make_named_test (hipThreadSafeDevice "hipThreadSafeDevice-serial" --tests 0x1) -make_named_test (hipThreadSafeDevice "hipThreadSafeDevice-pyramid" --tests 0x4) -make_named_test (hipThreadSafeDevice "hipThreadSafeDevice-nearzero" --tests 0x10) +make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-serial" --tests 0x1) +make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-pyramid" --tests 0x4) +make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-nearzero" --tests 0x10) make_hipify_test(specialFunc.cu ) diff --git a/projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp b/projects/clr/hipamd/tests/src/hipMultiThreadDevice.cpp similarity index 100% rename from projects/clr/hipamd/tests/src/hipThreadSafeDevice.cpp rename to projects/clr/hipamd/tests/src/hipMultiThreadDevice.cpp