diff --git a/bin/hipcc b/bin/hipcc index 44b1e85a33..3ef1ad1a8a 100755 --- a/bin/hipcc +++ b/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) { diff --git a/bin/hipconfig b/bin/hipconfig index d5f068c097..2d880f157d 100755 --- a/bin/hipconfig +++ b/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 @@ -116,6 +121,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; } diff --git a/docs/markdown/hip_faq.md b/docs/markdown/hip_faq.md index 734ced94ed..1a62784100 100644 --- a/docs/markdown/hip_faq.md +++ b/docs/markdown/hip_faq.md @@ -106,3 +106,10 @@ 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. + +### 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 +``` diff --git a/include/hcc_detail/hip_hcc.h b/include/hcc_detail/hip_hcc.h index a1e69ea02c..2cca8940e6 100644 --- a/include/hcc_detail/hip_hcc.h +++ b/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. @@ -63,8 +65,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 +88,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. @@ -111,8 +116,8 @@ struct 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 @@ -123,7 +128,7 @@ struct 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 @@ -132,14 +137,16 @@ struct 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 "<>\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;\ }) @@ -272,14 +280,110 @@ class FakeMutex #if STREAM_THREAD_SAFE typedef std::mutex StreamMutex; #else +#warning "Stream thread-safe disabled" typedef FakeMutex StreamMutex; #endif +#if DEVICE_THREAD_SAFE +typedef std::mutex DeviceMutex; +#else +typedef FakeMutex DeviceMutex; +#warning "Device thread-safe disabled" +#endif -// TODO - move async copy code into stream? Stream->async-copy. -// Add PreCopy / PostCopy to manage locks? // +//--- +// 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, bool autoUnlock=true) : + _criticalData(&criticalData), + _autoUnlock(autoUnlock) + { + _criticalData->_mutex.lock(); + }; + + ~LockedAccessor() + { + if (_autoUnlock) { + _criticalData->_mutex.unlock(); + } + } + + void unlock() + { + _criticalData->_mutex.unlock(); + } + + // Syntactic sugar so -> can be used to get the underlying type. + T *operator->() { return _criticalData; }; + +private: + T *_criticalData; + bool _autoUnlock; +}; + + +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 +{ +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 LockedAccessor_StreamCrit_t; @@ -288,69 +392,74 @@ 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 - void copySync (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(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 reclaimSignals_ts(SIGSEQNUM sigNum); - void wait(bool assertQueueEmpty=false); + 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(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty=false); - // 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 (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: // 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;}; + 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: // 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(ihipSignal_t *signal); + void waitCopy(LockedAccessor_StreamCrit_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; // index into the g_device array - unsigned _device_index; - 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. - - 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: @@ -379,9 +488,58 @@ struct ihipEvent_t { -//------------------------------------------------------------------------------------------------- -struct ihipDevice_t +//--- +// Data that must be protected with thread-safe access +// All members are private - this class must be accessed through friend LockedAccessor which +// will lock the mutex on construction and unlock on destruction. +// +// MUTEX_TYPE is template argument so can easily convert to FakeMutex for performance or stress testing. +template +class ihipDeviceCriticalBase_t : LockedBase { +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; +}; + +// 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 LockedAccessor_DeviceCrit_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: + 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. hipDeviceProp_t _props; // saved device properties. @@ -392,31 +550,27 @@ 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; StagingBuffer *_staging_buffer[2]; // one buffer for each direction. - ihipStream_t::SeqNum_t _stream_id; unsigned _device_flags; -public: - void init(unsigned device_index, hc::accelerator acc, unsigned flags); - ~ihipDevice_t(); - void reset(); +private: hipError_t getProperties(hipDeviceProp_t* prop); - void waitAllStreams(); - void 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; }; -// 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/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 38fbe4fc29..544c16cd83 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/include/hcc_detail/hip_runtime_api.h @@ -449,10 +449,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); /** @@ -550,13 +547,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/include/hcc_detail/trace_helper.h b/include/hcc_detail/trace_helper.h index 35a8b80c0c..13af1eab26 100644 --- a/include/hcc_detail/trace_helper.h +++ b/include/hcc_detail/trace_helper.h @@ -44,18 +44,45 @@ inline 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 -inline std::string ToString(T v) { +inline 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 <> -inline std::string ToString(hipMemcpyKind v) { +inline std::string ToString(hipEvent_t v) +{ + return ToString(&v); +}; + + + +// hipStream_t +template <> +inline std::string ToString(hipStream_t v) +{ + std::ostringstream ss; + if (v == NULL) { + ss << "stream:"; + } else { + ss << *v; + } + + return ss.str(); +}; + +// hipMemcpyKind specialization +template <> +inline std::string ToString(hipMemcpyKind v) +{ switch(v) { CASE_STR(hipMemcpyHostToHost); CASE_STR(hipMemcpyHostToDevice); @@ -68,13 +95,15 @@ inline std::string ToString(hipMemcpyKind v) { template <> -inline std::string ToString(hipError_t v) { +inline std::string ToString(hipError_t v) +{ return ihipErrorString(v); }; // Catch empty arguments case -inline std::string ToString() { +inline std::string ToString() +{ return (""); } @@ -83,6 +112,7 @@ inline std::string ToString() { // C++11 variadic template - peels off first argument, converts to string, and calls itself again to peel the next arg. // Strings are automatically separated by comma+space. template -inline std::string ToString(T first, Args... args) { +inline std::string ToString(T first, Args... args) +{ return ToString(first) + ", " + ToString(args...) ; } diff --git a/include/hip_runtime_api.h b/include/hip_runtime_api.h index cd0d841d97..3a0a4b399a 100644 --- a/include/hip_runtime_api.h +++ b/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/src/hip_device.cpp b/src/hip_device.cpp index b819673967..c935f20d34 100644 --- a/src/hip_device.cpp +++ b/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/src/hip_event.cpp b/src/hip_event.cpp index 27232875b9..f3c28e176e 100644 --- a/src/hip_event.cpp +++ b/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)); } @@ -67,7 +81,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; @@ -77,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); } @@ -117,7 +132,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 @@ -125,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/src/hip_hcc.cpp b/src/hip_hcc.cpp index 1c93ae48dd..6c80b84e0e 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -59,6 +59,7 @@ int HIP_LAUNCH_BLOCKING = 0; int HIP_PRINT_ENV = 0; int HIP_TRACE_API= 0; +int HIP_ATP_MARKER= 0; int HIP_DB= 0; int HIP_STAGING_SIZE = 64; /* size of staging buffers, in KB */ int HIP_STAGING_BUFFERS = 2; // TODO - remove, two buffers should be enough. @@ -124,42 +125,36 @@ 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), - _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(); +{ } //--- -void ihipStream_t::reclaimSignals_ts(SIGSEQNUM sigNum) +void ihipStream_t::locked_reclaimSignals(SIGSEQNUM sigNum) { + LockedAccessor_StreamCrit_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(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); @@ -167,32 +162,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(LockedAccessor_StreamCrit_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) +{ + LockedAccessor_StreamCrit_t crit(_criticalData); + + wait(crit, assertQueueEmpty); + }; + //--- ihipDevice_t * ihipStream_t::getDevice() const { @@ -208,26 +214,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(LockedAccessor_StreamCrit_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) ; @@ -235,13 +241,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); @@ -284,32 +290,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() { - _mutex.lock(); // 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 (_last_command_type != ihipCommandKernel) { - if (_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, _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[_last_command_type], ihipCommandName[ihipCommandKernel], _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[_last_command_type], ihipCommandName[ihipCommandKernel]); - this->waitCopy(_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[_last_command_type], ihipCommandName[ihipCommandKernel]); + this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel]); } } - _last_command_type = ihipCommandKernel; + crit->_last_command_type = ihipCommandKernel; } return addedSync; @@ -317,11 +323,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) { - _last_kernel_future = kernelFuture; + // We locked _criticalData in the lockopen_preKernelCommand() so OK to access here: + _criticalData._last_kernel_future = kernelFuture; - _mutex.unlock(); + _criticalData.unlock(); // paired with lock from lockopen_preKernelCommand. }; @@ -329,7 +337,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(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType) { int needSync = 0; @@ -338,22 +346,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) { @@ -368,10 +378,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; } @@ -383,10 +393,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 + LockedAccessor_DeviceCrit_t crit(_criticalData); + // Reset and remove streams: - _streams.clear(); + crit->streams().clear(); // Reset and release all memory stored in the tracker: am_memtracker_reset(_acc); @@ -395,13 +408,12 @@ void ihipDevice_t::reset() //--- -void ihipDevice_t::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); @@ -416,10 +428,10 @@ 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); - tprintf(DB_SYNC, "created device with default_stream=%p\n", _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()); @@ -429,6 +441,8 @@ void ihipDevice_t::init(unsigned device_index, hc::accelerator acc, unsigned fla }; + + ihipDevice_t::~ihipDevice_t() { if (_default_stream) { @@ -670,11 +684,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) { + LockedAccessor_DeviceCrit_t crit(_criticalData); + tprintf(DB_SYNC, "syncDefaultStream\n"); - for (auto streamI=_streams.begin(); streamI!=_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. @@ -684,20 +700,39 @@ void ihipDevice_t::syncDefaultStream(bool waitOnSelf) if (waitOnSelf || (stream != _default_stream)) { // TODO-hcc - use blocking or active wait here? // TODO-sync - cudaDeviceBlockingSync - stream->wait(); + stream->locked_wait(); } } } } +//--- +void ihipDevice_t::locked_addStream(ihipStream_t *s) +{ + LockedAccessor_DeviceCrit_t crit(_criticalData); + + crit->streams().push_back(s); + s->_id = crit->incStreamId(); +} + +//--- +void ihipDevice_t::locked_removeStream(ihipStream_t *s) +{ + LockedAccessor_DeviceCrit_t crit(_criticalData); + + crit->streams().remove(s); +} + //--- //Heavyweight synchronization that waits on all streams, ignoring hipStreamNonBlocking flag. -void ihipDevice_t::waitAllStreams() +void ihipDevice_t::locked_waitAllStreams() { + LockedAccessor_DeviceCrit_t crit(_criticalData); + tprintf(DB_SYNC, "waitAllStream\n"); - for (auto streamI=_streams.begin(); streamI!=_streams.end(); streamI++) { - (*streamI)->wait(); + for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) { + (*streamI)->locked_wait(); } } @@ -808,12 +843,12 @@ void ihipInit() READ_ENV_I(release, HIP_LAUNCH_BLOCKING, CUDA_LAUNCH_BLOCKING, "Make HIP APIs 'host-synchronous', so they block until any kernel launches or data copy commands complete. Alias: CUDA_LAUNCH_BLOCKING." ); READ_ENV_I(release, HIP_DB, 0, "Print various debug info. Bitmask, see hip_hcc.cpp for more information."); if ((HIP_DB & DB_API) && (HIP_TRACE_API == 0)) { - // Set HIP_TRACE_API before we read it, so it is printed correctly. + // Set HIP_TRACE_API default before we read it, so it is printed correctly. HIP_TRACE_API = 1; } - READ_ENV_I(release, HIP_TRACE_API, 0, "Trace each HIP API call. Print function name and return code to stderr as program executes."); + READ_ENV_I(release, HIP_ATP_MARKER, 0, "Add HIP function begin/end to ATP file generated with CodeXL"); READ_ENV_I(release, HIP_STAGING_SIZE, 0, "Size of each staging buffer (in KB)" ); READ_ENV_I(release, HIP_STAGING_BUFFERS, 0, "Number of staging buffers to use in each direction. 0=use hsa_memory_copy."); READ_ENV_I(release, HIP_PININPLACE, 0, "For unpinned transfers, pin the memory in-place in chunks before doing the copy. Under development."); @@ -823,6 +858,8 @@ void ihipInit() READ_ENV_I(release, HIP_DISABLE_HW_KERNEL_DEP, 0, "Disable HW dependencies before kernel commands - instead wait for dependency on host. -1 means ignore these dependencies. (debug mode)"); READ_ENV_I(release, HIP_DISABLE_HW_COPY_DEP, 0, "Disable HW dependencies before copy commands - instead wait for dependency on host. -1 means ifnore these dependencies (debug mode)"); + + // Some flags have both compile-time and runtime flags - generate a warning if user enables the runtime flag but the compile-time flag is disabled. if (HIP_DB && !COMPILE_HIP_DB) { fprintf (stderr, "warning: env var HIP_DB=0x%x but COMPILE_HIP_DB=0. (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_DB); } @@ -831,6 +868,10 @@ void ihipInit() fprintf (stderr, "warning: env var HIP_TRACE_API=0x%x but COMPILE_HIP_TRACE_API=0. (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_DB); } + if (HIP_ATP_MARKER && !COMPILE_HIP_ATP_MARKER) { + fprintf (stderr, "warning: env var HIP_ATP_MARKER=0x%x but COMPILE_HIP_ATP_MARKER=0. (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_ATP_MARKER); + } + /* * Build a table of valid compute devices. @@ -933,14 +974,14 @@ 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 { // 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; @@ -956,7 +997,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; @@ -968,7 +1010,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); } @@ -1072,7 +1114,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(LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind) { ihipDevice_t *device = this->getDevice(); @@ -1098,7 +1140,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); @@ -1109,7 +1151,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); @@ -1120,14 +1162,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. @@ -1139,7 +1181,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. @@ -1154,10 +1196,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); @@ -1168,7 +1210,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); } @@ -1176,10 +1218,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) +{ + LockedAccessor_StreamCrit_t crit (_criticalData); + copySync(crit, dst, src, sizeBytes, kind); +} + void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind) { + LockedAccessor_StreamCrit_t crit(_criticalData); + ihipDevice_t *device = this->getDevice(); if (device == NULL) { @@ -1193,7 +1244,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); @@ -1220,8 +1271,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); @@ -1232,7 +1282,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); @@ -1242,7 +1292,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. @@ -1250,7 +1300,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); } } } @@ -1302,6 +1352,8 @@ 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). +// locked_ // TODO - describe MT strategy // //// TODO - add identifier numbers for streams and devices to help with debugging. diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 46fbd38fd5..4d75a1d7a1 100644 --- a/src/hip_memory.cpp +++ b/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); } @@ -366,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 ; @@ -392,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) { @@ -459,7 +456,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/src/hip_stream.cpp b/src/hip_stream.cpp index 6dbf7458ce..d62abc49e2 100644 --- a/src/hip_stream.cpp +++ b/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; @@ -41,12 +39,32 @@ 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); - device->_streams.push_back(istream); + auto istream = new ihipStream_t(device->_device_index, acc.create_view(), flags); + + device->locked_addStream(istream); + *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)); } @@ -56,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; @@ -65,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; } @@ -76,15 +93,15 @@ 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; if (stream == NULL) { ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->syncDefaultStream(true/*waitOnSelf*/); + device->locked_syncDefaultStream(true/*waitOnSelf*/); } else { - stream->wait(); + stream->locked_wait(); e = hipSuccess; } @@ -99,23 +116,23 @@ 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; //--- Drain the stream: if (stream == NULL) { ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->syncDefaultStream(true/*waitOnSelf*/); + device->locked_syncDefaultStream(true/*waitOnSelf*/); } else { - stream->wait(); + stream->locked_wait(); e = hipSuccess; } ihipDevice_t *device = stream->getDevice(); if (device) { - device->_streams.remove(stream); + device->locked_removeStream(stream); delete stream; } else { e = hipErrorInvalidResourceHandle; @@ -128,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/tests/src/CMakeLists.txt b/tests/src/CMakeLists.txt index 5cc280cec7..7e5974c5c3 100644 --- a/tests/src/CMakeLists.txt +++ b/tests/src/CMakeLists.txt @@ -8,8 +8,9 @@ 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}) -MESSAGE("HIP_PATH=" ${HIP_PATH}) if (NOT DEFINED HIP_PATH) set (HIP_PATH ../..) endif() @@ -39,12 +40,16 @@ 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. + 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") MESSAGE ("HIP_PLATFORM=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) @@ -156,6 +162,8 @@ make_hip_executable (hipFuncGetDevice hipFuncGetDevice.cpp) make_hip_executable (hipFuncSetDevice hipFuncSetDevice.cpp) make_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp) +make_hip_executable (hipMultiThreadDevice hipMultiThreadDevice.cpp) + make_test(hip_ballot " " ) make_test(hip_anyall " " ) make_test(hip_popc " " ) @@ -171,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 ) @@ -196,5 +205,8 @@ make_test(hipFuncSetDeviceFlags " ") make_test(hipFuncGetDevice " ") make_test(hipFuncSetDevice " ") make_test(hipFuncDeviceSynchronize " ") +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/tests/src/hipMultiThreadDevice.cpp b/tests/src/hipMultiThreadDevice.cpp new file mode 100644 index 0000000000..a1f64aceb3 --- /dev/null +++ b/tests/src/hipMultiThreadDevice.cpp @@ -0,0 +1,138 @@ +#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, 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/tests/src/test_common.h b/tests/src/test_common.h index d92692c978..ce85b3898a 100644 --- a/tests/src/test_common.h +++ b/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; } } } diff --git a/util/vim/hip.vim b/util/vim/hip.vim index 73cf71eb73..b530a3f9e5 100644 --- a/util/vim/hip.vim +++ b/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