Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into privatestaging
[ROCm/clr commit: 00bf37d28d]
This commit is contained in:
Aditya Atluri
@@ -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) {
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
|
||||
@@ -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
|
||||
```
|
||||
|
||||
@@ -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 "<<hip-api: %s\n" KNRM, s.c_str());\
|
||||
if (HIP_ATP_MARKER || (COMPILE_HIP_DB && HIP_TRACE_API)) {\
|
||||
std::string s = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')';\
|
||||
if (COMPILE_HIP_DB && HIP_TRACE_API) {\
|
||||
fprintf (stderr, API_COLOR "<<hip-api: %s\n" KNRM, s.c_str());\
|
||||
}\
|
||||
SCOPED_MARKER(s.c_str(), "HIP", NULL);\
|
||||
}\
|
||||
SCOPED_MARKER(s.c_str(), "HIP", NULL);\
|
||||
}
|
||||
#else
|
||||
// Swallow API_TRACE
|
||||
@@ -157,12 +164,13 @@ struct 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;\
|
||||
})
|
||||
|
||||
|
||||
@@ -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<typename T>
|
||||
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 <typename MUTEX_TYPE>
|
||||
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 <typename MUTEX_TYPE>
|
||||
class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE>
|
||||
{
|
||||
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<StreamMutex> * mlock() { LockedBase<MUTEX_TYPE>::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<ihipSignal_t> _signalPool; // Pool of signals for use by this stream.
|
||||
|
||||
|
||||
SIGSEQNUM _stream_sig_id; // Monotonically increasing unique signal id.
|
||||
};
|
||||
|
||||
|
||||
typedef ihipStreamCriticalBase_t<StreamMutex> ihipStreamCritical_t;
|
||||
typedef LockedAccessor<ihipStreamCritical_t> 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<ihipSignal_t> _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 <typename MUTEX_TYPE>
|
||||
class ihipDeviceCriticalBase_t : LockedBase<MUTEX_TYPE>
|
||||
{
|
||||
public:
|
||||
ihipDeviceCriticalBase_t() : _stream_id(0) {};
|
||||
friend class LockedAccessor<ihipDeviceCriticalBase_t>;
|
||||
|
||||
std::list<ihipStream_t*> &streams() { return _streams; };
|
||||
const std::list<ihipStream_t*> &const_streams() const { return _streams; };
|
||||
|
||||
// "Allocate" a stream ID:
|
||||
ihipStream_t::SeqNum_t incStreamId() { return _stream_id++; };
|
||||
|
||||
|
||||
private:
|
||||
std::list<ihipStream_t*> _streams; // streams associated with this device.
|
||||
ihipStream_t::SeqNum_t _stream_id;
|
||||
};
|
||||
|
||||
// Note Mutex selected based on DeviceMutex
|
||||
typedef ihipDeviceCriticalBase_t<DeviceMutex> ihipDeviceCritical_t;
|
||||
|
||||
// This type is used by functions that need access to the critical device structures.
|
||||
typedef LockedAccessor<ihipDeviceCritical_t> 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<ihipStream_t*> _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<ihipDeviceCritical_t> 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
|
||||
|
||||
@@ -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);
|
||||
|
||||
|
||||
/**
|
||||
|
||||
@@ -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 <typename T>
|
||||
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:<null>";
|
||||
} 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 <typename T, typename... Args>
|
||||
inline std::string ToString(T first, Args... args) {
|
||||
inline std::string ToString(T first, Args... args)
|
||||
{
|
||||
return ToString(first) + ", " + ToString(args...) ;
|
||||
}
|
||||
|
||||
@@ -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<class T>
|
||||
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);
|
||||
}
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
|
||||
+155
-103
@@ -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<hsa_signal_t*> (_last_kernel_future.get_native_handle()));
|
||||
hsa_signal_t *hsaSignal = (static_cast<hsa_signal_t*> (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<hsa_agent_t*> (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<hsa_region_t*>(_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.
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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_OBJECTS:test_common> )
|
||||
target_link_libraries(${exe} hip_hcc)
|
||||
if (${HIP_BUILD_LOCAL})
|
||||
target_link_libraries(${exe} hip_hcc)
|
||||
endif()
|
||||
else()
|
||||
add_executable (${exe} ${cpp} ${ARGN} $<TARGET_OBJECTS:test_common> )
|
||||
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 )
|
||||
|
||||
@@ -0,0 +1,138 @@
|
||||
#include <hip_runtime_api.h>
|
||||
#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<iterations; i++) {
|
||||
if (p_verbose & 0x1) {
|
||||
printf ("%s iter=%d, create %d then destroy %d\n", __func__, i, burstSize, burstSize);
|
||||
}
|
||||
for (int j=0; j<burstSize; j++) {
|
||||
if (p_verbose & 0x2) {
|
||||
printf (" %d.%d streamCreate\n", i, j);
|
||||
}
|
||||
HIPCHECK( hipStreamCreate(&streams[j]));
|
||||
}
|
||||
for (int j=0; j<burstSize; j++) {
|
||||
if (p_verbose & 0x2) {
|
||||
printf (" %d.%d streamDestroy\n", i, j);
|
||||
}
|
||||
HIPCHECK( hipStreamDestroy(streams[j]));
|
||||
}
|
||||
}
|
||||
|
||||
delete streams;
|
||||
}
|
||||
|
||||
|
||||
void waitStreams(int iterations)
|
||||
{
|
||||
// Repeatedly sync and wait for all streams to complete.
|
||||
// TO make this interesting, the test has other threads repeatedly adding and removing streams to the device.
|
||||
for (int i=0; i<iterations; i++) {
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
// Create 3 streams, all creating and destroying streams on the same device.
|
||||
// Some create many queue, some not many.
|
||||
//
|
||||
void multiThread_pyramid(bool serialize, int iters)
|
||||
{
|
||||
printf ("%s creating %d streams\n", __func__, iters*100);
|
||||
std::thread t1 (createThenDestroyStreams, iters*1, 100);
|
||||
if (serialize) {
|
||||
t1.join();
|
||||
printf("t1 done\n");
|
||||
}
|
||||
|
||||
std::thread t2 (createThenDestroyStreams, iters*10, 10);
|
||||
if (serialize) {
|
||||
t2.join();
|
||||
printf("t2 done\n");
|
||||
}
|
||||
|
||||
std::thread t3 (createThenDestroyStreams, iters*100, 1);
|
||||
if (serialize) {
|
||||
t3.join();
|
||||
printf("t3 done\n");
|
||||
}
|
||||
|
||||
if (!serialize) {
|
||||
t1.join();
|
||||
t2.join();
|
||||
t3.join();
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
|
||||
// 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_nearzero(bool serialize, int iters)
|
||||
{
|
||||
printf ("%s creating %d streams x 3 threads\n", __func__, iters);
|
||||
std::thread t1 (createThenDestroyStreams, iters, 1);
|
||||
if (serialize) {
|
||||
t1.join();
|
||||
printf("t1 done\n");
|
||||
}
|
||||
|
||||
std::thread t2 (createThenDestroyStreams, iters, 1);
|
||||
if (serialize) {
|
||||
t2.join();
|
||||
printf("t2 done\n");
|
||||
}
|
||||
|
||||
std::thread t3 (waitStreams, iters*50);
|
||||
if (serialize) {
|
||||
t3.join();
|
||||
printf("t3 done\n");
|
||||
}
|
||||
|
||||
if (!serialize) {
|
||||
t1.join(); printf ("t1 done\n");
|
||||
t2.join(); printf ("t2 done\n");
|
||||
t3.join(); printf ("t3 done\n");
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
// Serial version, just call once:
|
||||
if (p_tests & 0x1) {
|
||||
printf ("\ntest 0x1 : serial createThenDestroyStreams(10) \n");
|
||||
createThenDestroyStreams(10, 10);
|
||||
};
|
||||
|
||||
/*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);
|
||||
}
|
||||
|
||||
if (p_tests & 0x4) {
|
||||
printf ("\ntest 0x4 : parallel multiThread_pyramid(1) \n");
|
||||
multiThread_pyramid(false, 10);
|
||||
}
|
||||
|
||||
//if (p_tests & 0x8) {
|
||||
// printf ("test 0x8 : multiThread_pyramid(100) \n");
|
||||
// multiThread_pyramid(false, 100);
|
||||
// }
|
||||
|
||||
if (p_tests & 0x10) {
|
||||
printf ("\ntest 0x10 : parallel multiThread_nearzero(1000) \n");
|
||||
multiThread_nearzero(false, 1000);
|
||||
}
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -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<T, HipTest::Pinned, C>, 2000000/*mb*/, 1000, stream0);
|
||||
std::thread t1 (simpleVectorCopy<T, HipTest::Pinned, C>, 2000000/*mb*/, 100/*iters*/, stream0);
|
||||
if (serialize) {
|
||||
t1.join();
|
||||
}
|
||||
std::thread t2 (simpleVectorCopy<T, HipTest::Pinned, C>, 2000000/*mb*/, 1000, stream1);
|
||||
std::thread t2 (simpleVectorCopy<T, HipTest::Pinned, C>, 2000000/*mb*/, 100/*iters*/, stream1);
|
||||
if (serialize) {
|
||||
t2.join();
|
||||
}
|
||||
@@ -119,19 +119,21 @@ int main(int argc, char *argv[])
|
||||
simpleVectorCopy<float, HipTest::Pinned, HipTest::MemcpyAsync> (2000000/*mb*/, 10/*iters*/, stream);
|
||||
simpleVectorCopy<float, HipTest::Pinned, HipTest::Memcpy> (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<float, HipTest::MemcpyAsync> ("Multithread NULL with serialized", NULL, NULL, true);
|
||||
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with serialized", stream0, stream1, true);
|
||||
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread two streams serialized", stream0, stream1, true);
|
||||
}
|
||||
|
||||
if (p_tests & 0x4) {
|
||||
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with NULL stream", NULL, NULL, false);
|
||||
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with two streams", stream0, stream1, false);
|
||||
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with one stream", stream0, stream0, false);
|
||||
|
||||
@@ -18,6 +18,7 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <iostream>
|
||||
#include <iomanip>
|
||||
#include <sys/time.h>
|
||||
#include <stddef.h>
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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
|
||||
|
||||
Reference in New Issue
Block a user