Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into privatestaging
[ROCm/hip commit: c567189cb8]
이 커밋은 다음에 포함됨:
Ben Sander
@@ -23,7 +23,14 @@ set(CMAKE_C_COMPILER "${HCC_HOME}/bin/hcc")
|
||||
set(CMAKE_CXX_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I${HIP_PATH}/include -stdlib=libc++ ")
|
||||
set(CMAKE_C_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I${HIP_PATH}/include -stdlib=libc++ ")
|
||||
|
||||
set(SOURCE_FILES src/hip_hcc.cpp src/staging_buffer.cpp)
|
||||
set(SOURCE_FILES src/hip_hcc.cpp
|
||||
src/hip_device.cpp
|
||||
src/hip_error.cpp
|
||||
src/hip_event.cpp
|
||||
src/hip_memory.cpp
|
||||
src/hip_peer.cpp
|
||||
src/hip_stream.cpp
|
||||
src/staging_buffer.cpp)
|
||||
|
||||
add_library(hip_hcc STATIC ${SOURCE_FILES})
|
||||
|
||||
|
||||
@@ -2,8 +2,9 @@
|
||||
#define HIP_HCC_H
|
||||
|
||||
#include <hc.hpp>
|
||||
#include "hip_util.h"
|
||||
#include "staging_buffer.h"
|
||||
#include "hcc_detail/hip_util.h"
|
||||
#include "hcc_detail/staging_buffer.h"
|
||||
|
||||
#define HIP_HCC
|
||||
|
||||
#if defined(__HCC__) && (__hcc_workweek__ < 1502)
|
||||
@@ -22,29 +23,30 @@
|
||||
|
||||
// Intended to distinguish whether an environment variable should be visible only in debug mode, or in debug+release.
|
||||
//static const int debug = 0;
|
||||
static const int release = 1;
|
||||
extern const int release;
|
||||
|
||||
extern int HIP_LAUNCH_BLOCKING;
|
||||
|
||||
static int HIP_LAUNCH_BLOCKING = 0;
|
||||
|
||||
static int HIP_PRINT_ENV = 0;
|
||||
static int HIP_TRACE_API= 0;
|
||||
static int HIP_DB= 0;
|
||||
static int HIP_STAGING_SIZE = 64; /* size of staging buffers, in KB */
|
||||
static int HIP_STAGING_BUFFERS = 2; // TODO - remove, two buffers should be enough.
|
||||
static int HIP_PININPLACE = 0;
|
||||
static int HIP_STREAM_SIGNALS = 2; /* number of signals to allocate at stream creation */
|
||||
static int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU identifiers */
|
||||
|
||||
extern int HIP_PRINT_ENV;
|
||||
extern int HIP_TRACE_API;
|
||||
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.
|
||||
extern int HIP_PININPLACE;
|
||||
extern int HIP_STREAM_SIGNALS; /* number of signals to allocate at stream creation */
|
||||
extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU identifiers */
|
||||
|
||||
|
||||
//---
|
||||
// Chicken bits for disabling functionality to work around potential issues:
|
||||
static int HIP_DISABLE_HW_KERNEL_DEP = 1;
|
||||
static int HIP_DISABLE_HW_COPY_DEP = 1;
|
||||
extern int HIP_DISABLE_HW_KERNEL_DEP;
|
||||
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;
|
||||
|
||||
static thread_local int tls_defaultDevice = 0;
|
||||
static thread_local hipError_t tls_lastHipError = hipSuccess;
|
||||
|
||||
// Color defs for debug messages:
|
||||
#define KNRM "\x1B[0m"
|
||||
@@ -96,7 +98,7 @@ static thread_local hipError_t tls_lastHipError = hipSuccess;
|
||||
|
||||
|
||||
// #include CPP files to produce one object file
|
||||
//#define ONE_OBJECT_FILE 1
|
||||
#define ONE_OBJECT_FILE 1
|
||||
|
||||
|
||||
// Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary.
|
||||
@@ -134,6 +136,16 @@ static thread_local hipError_t tls_lastHipError = hipSuccess;
|
||||
std::call_once(hip_initialized, ihipInit);\
|
||||
API_TRACE(__VA_ARGS__);
|
||||
|
||||
#define ihipLogStatus(_hip_status) \
|
||||
({\
|
||||
tls_lastHipError = _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));\
|
||||
}\
|
||||
_hip_status;\
|
||||
})
|
||||
|
||||
|
||||
|
||||
|
||||
@@ -179,14 +191,11 @@ public:
|
||||
};
|
||||
|
||||
|
||||
struct ihipStream_t;
|
||||
struct ihipDevice_t;
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
|
||||
typedef class ihipStream_t *hipStream_t;
|
||||
typedef class ihipStream_t* hipStream_t;
|
||||
typedef struct hipEvent_t {
|
||||
struct ihipEvent_t *_handle;
|
||||
} hipEvent_t;
|
||||
@@ -387,13 +396,95 @@ private:
|
||||
|
||||
|
||||
// Global initialization.
|
||||
static std::once_flag hip_initialized;
|
||||
static ihipDevice_t *g_devices; // Array of all non-emulated (ie GPU) accelerators in the system.
|
||||
static bool g_visible_device = false; // Set the flag when HIP_VISIBLE_DEVICES is set
|
||||
static unsigned g_deviceCnt;
|
||||
static std::vector<int> g_hip_visible_devices; /* vector of integers that contains the visible device IDs */
|
||||
static hsa_agent_t g_cpu_agent ; // the CPU agent.
|
||||
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
|
||||
extern unsigned g_deviceCnt;
|
||||
extern std::vector<int> g_hip_visible_devices; /* vector of integers that contains the visible device IDs */
|
||||
extern hsa_agent_t g_cpu_agent ; // the CPU agent.
|
||||
//=================================================================================================
|
||||
void ihipInit();
|
||||
const char *ihipErrorString(hipError_t);
|
||||
ihipDevice_t *ihipGetTlsDefaultDevice();
|
||||
ihipDevice_t *ihipGetDevice(int);
|
||||
void ihipSetTs(hipEvent_t e);
|
||||
|
||||
template<typename T>
|
||||
hc::completion_future ihipMemcpyKernel(hipStream_t, T*, const T*, size_t);
|
||||
|
||||
template<typename T>
|
||||
hc::completion_future ihipMemsetKernel(hipStream_t, T*, T, size_t);
|
||||
|
||||
hipStream_t ihipSyncAndResolveStream(hipStream_t);
|
||||
template <typename T>
|
||||
|
||||
hc::completion_future
|
||||
ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes)
|
||||
{
|
||||
int wg = std::min((unsigned)8, stream->getDevice()->_compute_units);
|
||||
const int threads_per_wg = 256;
|
||||
|
||||
int threads = wg * threads_per_wg;
|
||||
if (threads > sizeBytes) {
|
||||
threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;
|
||||
}
|
||||
|
||||
|
||||
hc::extent<1> ext(threads);
|
||||
auto ext_tile = ext.tile(threads_per_wg);
|
||||
|
||||
hc::completion_future cf =
|
||||
hc::parallel_for_each(
|
||||
stream->_av,
|
||||
ext_tile,
|
||||
[=] (hc::tiled_index<1> idx)
|
||||
__attribute__((hc))
|
||||
{
|
||||
int offset = amp_get_global_id(0);
|
||||
// TODO-HCC - change to hc_get_local_size()
|
||||
int stride = amp_get_local_size(0) * hc_get_num_groups(0) ;
|
||||
|
||||
for (int i=offset; i<sizeBytes; i+=stride) {
|
||||
ptr[i] = val;
|
||||
}
|
||||
});
|
||||
|
||||
return cf;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
hc::completion_future
|
||||
ihipMemcpyKernel(hipStream_t stream, T * c, const T * a, size_t sizeBytes)
|
||||
{
|
||||
int wg = std::min((unsigned)8, stream->getDevice()->_compute_units);
|
||||
const int threads_per_wg = 256;
|
||||
|
||||
int threads = wg * threads_per_wg;
|
||||
if (threads > sizeBytes) {
|
||||
threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;
|
||||
}
|
||||
|
||||
|
||||
hc::extent<1> ext(threads);
|
||||
auto ext_tile = ext.tile(threads_per_wg);
|
||||
|
||||
hc::completion_future cf =
|
||||
hc::parallel_for_each(
|
||||
stream->_av,
|
||||
ext_tile,
|
||||
[=] (hc::tiled_index<1> idx)
|
||||
__attribute__((hc))
|
||||
{
|
||||
int offset = amp_get_global_id(0);
|
||||
// TODO-HCC - change to hc_get_local_size()
|
||||
int stride = amp_get_local_size(0) * hc_get_num_groups(0) ;
|
||||
|
||||
for (int i=offset; i<sizeBytes; i+=stride) {
|
||||
c[i] = a[i];
|
||||
}
|
||||
});
|
||||
|
||||
return cf;
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
@@ -39,7 +39,7 @@ THE SOFTWARE.
|
||||
#define CUDA_SUCCESS hipSuccess
|
||||
|
||||
#include <hip_runtime_api.h>
|
||||
|
||||
#include "hcc_detail/hip_hcc.h"
|
||||
//---
|
||||
// Remainder of this file only compiles with HCC
|
||||
#ifdef __HCC__
|
||||
@@ -52,7 +52,6 @@ THE SOFTWARE.
|
||||
|
||||
#include <hcc_detail/hip_texture.h>
|
||||
#include <hcc_detail/host_defines.h>
|
||||
|
||||
// TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
|
||||
#if defined (__KALMAR_ACCELERATOR__) && not defined (__HCC_ACCELERATOR__)
|
||||
#define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__
|
||||
@@ -497,8 +496,6 @@ __device__ inline float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); };
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
#define HIP_KERNEL_NAME(...) __VA_ARGS__
|
||||
|
||||
|
||||
|
||||
@@ -30,7 +30,8 @@ THE SOFTWARE.
|
||||
#include <stddef.h>
|
||||
|
||||
#include <hcc_detail/host_defines.h>
|
||||
#include <hcc_detail/hip_hcc.h>
|
||||
#include <hip_runtime_api.h>
|
||||
#include "hip_hcc.h"
|
||||
|
||||
#if defined (__HCC__) && (__hcc_workweek__ < 1602)
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
|
||||
@@ -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 = hipHostMallocDefault)
|
||||
static inline hipError_t hipHostMalloc( T** ptr, size_t size, unsigned int flags = 0)
|
||||
{
|
||||
return hipHostMalloc((void**)ptr, size, flags);
|
||||
}
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
#include "hip_runtime.h"
|
||||
#include "hcc_detail/hip_hcc.h"
|
||||
#include "hcc_detail/trace_helper.h"
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//Devices
|
||||
@@ -111,7 +114,6 @@ hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig )
|
||||
hipError_t hipSetDevice(int device)
|
||||
{
|
||||
HIP_INIT_API(device);
|
||||
|
||||
if ((device < 0) || (device >= g_deviceCnt)) {
|
||||
return ihipLogStatus(hipErrorInvalidDevice);
|
||||
} else {
|
||||
@@ -263,3 +265,6 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device)
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
@@ -1,3 +1,7 @@
|
||||
#include "hip_runtime.h"
|
||||
#include "hcc_detail/hip_hcc.h"
|
||||
#include "hcc_detail/trace_helper.h"
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
// Error Handling
|
||||
@@ -26,31 +30,6 @@ hipError_t hipPeakAtLastError()
|
||||
return ihipLogStatus(tls_lastHipError);
|
||||
}
|
||||
|
||||
|
||||
const char *ihipErrorString(hipError_t hip_error)
|
||||
{
|
||||
switch (hip_error) {
|
||||
case hipSuccess : return "hipSuccess";
|
||||
case hipErrorMemoryAllocation : return "hipErrorMemoryAllocation";
|
||||
case hipErrorMemoryFree : return "hipErrorMemoryFree";
|
||||
case hipErrorUnknownSymbol : return "hipErrorUnknownSymbol";
|
||||
case hipErrorOutOfResources : return "hipErrorOutOfResources";
|
||||
case hipErrorInvalidValue : return "hipErrorInvalidValue";
|
||||
case hipErrorInvalidResourceHandle : return "hipErrorInvalidResourceHandle";
|
||||
case hipErrorInvalidDevice : return "hipErrorInvalidDevice";
|
||||
case hipErrorInvalidMemcpyDirection : return "hipErrorInvalidMemcpyDirection";
|
||||
case hipErrorNoDevice : return "hipErrorNoDevice";
|
||||
case hipErrorNotReady : return "hipErrorNotReady";
|
||||
case hipErrorRuntimeMemory : return "hipErrorRuntimeMemory";
|
||||
case hipErrorRuntimeOther : return "hipErrorRuntimeOther";
|
||||
case hipErrorUnknown : return "hipErrorUnknown";
|
||||
case hipErrorTbd : return "hipErrorTbd";
|
||||
default : return "hipErrorUnknown";
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
|
||||
//---
|
||||
const char *hipGetErrorName(hipError_t hip_error)
|
||||
{
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
#include "hip_runtime.h"
|
||||
#include "hcc_detail/hip_hcc.h"
|
||||
#include "hcc_detail/trace_helper.h"
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
@@ -113,25 +116,6 @@ hipError_t hipEventSynchronize(hipEvent_t event)
|
||||
}
|
||||
|
||||
|
||||
void ihipSetTs(hipEvent_t e)
|
||||
{
|
||||
ihipEvent_t *eh = e._handle;
|
||||
if (eh->_state == hipEventStatusRecorded) {
|
||||
// already recorded, done:
|
||||
return;
|
||||
} else {
|
||||
// TODO - use completion-future functions to obtain ticks and timestamps:
|
||||
hsa_signal_t *sig = static_cast<hsa_signal_t*> (eh->_marker.get_native_handle());
|
||||
if (sig) {
|
||||
if (hsa_signal_load_acquire(*sig) == 0) {
|
||||
eh->_timestamp = eh->_marker.get_end_tick();
|
||||
eh->_state = hipEventStatusRecorded;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop)
|
||||
{
|
||||
|
||||
+341
-20
@@ -40,7 +40,7 @@ THE SOFTWARE.
|
||||
#include <hc_am.hpp>
|
||||
|
||||
#include "hip_runtime.h"
|
||||
|
||||
#include "hcc_detail/hip_hcc.h"
|
||||
#include "hsa_ext_amd.h"
|
||||
|
||||
// HIP includes:
|
||||
@@ -52,12 +52,45 @@ THE SOFTWARE.
|
||||
extern const char *ihipErrorString(hipError_t hip_error);
|
||||
#include "hcc_detail/trace_helper.h"
|
||||
|
||||
const int release = 1;
|
||||
|
||||
|
||||
int HIP_LAUNCH_BLOCKING = 0;
|
||||
|
||||
int HIP_PRINT_ENV = 0;
|
||||
int HIP_TRACE_API= 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.
|
||||
int HIP_PININPLACE = 0;
|
||||
int HIP_STREAM_SIGNALS = 2; /* number of signals to allocate at stream creation */
|
||||
int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU identifiers */
|
||||
|
||||
|
||||
//---
|
||||
// Chicken bits for disabling functionality to work around potential issues:
|
||||
int HIP_DISABLE_HW_KERNEL_DEP = 1;
|
||||
int HIP_DISABLE_HW_COPY_DEP = 1;
|
||||
|
||||
thread_local int tls_defaultDevice = 0;
|
||||
thread_local hipError_t tls_lastHipError = hipSuccess;
|
||||
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
//Forward Declarations:
|
||||
//=================================================================================================
|
||||
bool ihipIsValidDevice(unsigned deviceIndex);
|
||||
|
||||
std::once_flag hip_initialized;
|
||||
ihipDevice_t *g_devices;
|
||||
bool g_visible_device = false;
|
||||
unsigned g_deviceCnt;
|
||||
std::vector<int> g_hip_visible_devices;
|
||||
hsa_agent_t g_cpu_agent;
|
||||
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
// Implementation:
|
||||
//=================================================================================================
|
||||
@@ -664,18 +697,6 @@ void ihipDevice_t::waitAllStreams()
|
||||
|
||||
|
||||
|
||||
#define ihipLogStatus(_hip_status) \
|
||||
({\
|
||||
tls_lastHipError = _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));\
|
||||
}\
|
||||
_hip_status;\
|
||||
})
|
||||
|
||||
|
||||
|
||||
// Read environment variables.
|
||||
void ihipReadEnv_I(int *var_ptr, const char *var_name1, const char *var_name2, const char *description)
|
||||
{
|
||||
@@ -809,6 +830,7 @@ void ihipInit()
|
||||
* Build a table of valid compute devices.
|
||||
*/
|
||||
auto accs = hc::accelerator::get_all();
|
||||
|
||||
int deviceCnt = 0;
|
||||
for (int i=0; i<accs.size(); i++) {
|
||||
if (! accs[i].get_is_emulated()) {
|
||||
@@ -960,16 +982,315 @@ void ihipPostLaunchKernel(hipStream_t stream, hc::completion_future &kernelFutur
|
||||
//
|
||||
//---
|
||||
|
||||
#include "hip_device.cpp"
|
||||
#include "hip_error.cpp"
|
||||
#include "hip_stream.cpp"
|
||||
#include "hip_event.cpp"
|
||||
#include "hip_memory.cpp"
|
||||
#include "hip_p2p.cpp"
|
||||
#include "hip_misc.cpp"
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
|
||||
|
||||
|
||||
const char *ihipErrorString(hipError_t hip_error)
|
||||
{
|
||||
switch (hip_error) {
|
||||
case hipSuccess : return "hipSuccess";
|
||||
case hipErrorMemoryAllocation : return "hipErrorMemoryAllocation";
|
||||
case hipErrorMemoryFree : return "hipErrorMemoryFree";
|
||||
case hipErrorUnknownSymbol : return "hipErrorUnknownSymbol";
|
||||
case hipErrorOutOfResources : return "hipErrorOutOfResources";
|
||||
case hipErrorInvalidValue : return "hipErrorInvalidValue";
|
||||
case hipErrorInvalidResourceHandle : return "hipErrorInvalidResourceHandle";
|
||||
case hipErrorInvalidDevice : return "hipErrorInvalidDevice";
|
||||
case hipErrorInvalidMemcpyDirection : return "hipErrorInvalidMemcpyDirection";
|
||||
case hipErrorNoDevice : return "hipErrorNoDevice";
|
||||
case hipErrorNotReady : return "hipErrorNotReady";
|
||||
case hipErrorRuntimeMemory : return "hipErrorRuntimeMemory";
|
||||
case hipErrorRuntimeOther : return "hipErrorRuntimeOther";
|
||||
case hipErrorUnknown : return "hipErrorUnknown";
|
||||
case hipErrorTbd : return "hipErrorTbd";
|
||||
default : return "hipErrorUnknown";
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
void ihipSetTs(hipEvent_t e)
|
||||
{
|
||||
ihipEvent_t *eh = e._handle;
|
||||
if (eh->_state == hipEventStatusRecorded) {
|
||||
// already recorded, done:
|
||||
return;
|
||||
} else {
|
||||
// TODO - use completion-future functions to obtain ticks and timestamps:
|
||||
hsa_signal_t *sig = static_cast<hsa_signal_t*> (eh->_marker.get_native_handle());
|
||||
if (sig) {
|
||||
if (hsa_signal_load_acquire(*sig) == 0) {
|
||||
eh->_timestamp = eh->_marker.get_end_tick();
|
||||
eh->_state = hipEventStatusRecorded;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
// Resolve hipMemcpyDefault to a known type.
|
||||
unsigned ihipStream_t::resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem)
|
||||
{
|
||||
hipMemcpyKind kind = hipMemcpyDefault;
|
||||
|
||||
if (!srcInDeviceMem && !dstInDeviceMem) {
|
||||
kind = hipMemcpyHostToHost;
|
||||
} else if (!srcInDeviceMem && dstInDeviceMem) {
|
||||
kind = hipMemcpyHostToDevice;
|
||||
} else if (srcInDeviceMem && !dstInDeviceMem) {
|
||||
kind = hipMemcpyDeviceToHost;
|
||||
} else if (srcInDeviceMem && dstInDeviceMem) {
|
||||
kind = hipMemcpyDeviceToDevice;
|
||||
}
|
||||
|
||||
assert (kind != hipMemcpyDefault);
|
||||
|
||||
return kind;
|
||||
}
|
||||
|
||||
|
||||
// Setup the copyCommandType and the copy agents (for hsa_amd_memory_async_copy)
|
||||
void ihipStream_t::setCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent)
|
||||
{
|
||||
ihipDevice_t *device = this->getDevice();
|
||||
hsa_agent_t deviceAgent = device->_hsa_agent;
|
||||
|
||||
switch (kind) {
|
||||
case hipMemcpyHostToHost : *commandType = ihipCommandCopyH2H; *srcAgent=g_cpu_agent; *dstAgent=g_cpu_agent; break;
|
||||
case hipMemcpyHostToDevice : *commandType = ihipCommandCopyH2D; *srcAgent=g_cpu_agent; *dstAgent=deviceAgent; break;
|
||||
case hipMemcpyDeviceToHost : *commandType = ihipCommandCopyD2H; *srcAgent=deviceAgent; *dstAgent=g_cpu_agent; break;
|
||||
case hipMemcpyDeviceToDevice : *commandType = ihipCommandCopyD2D; *srcAgent=deviceAgent; *dstAgent=deviceAgent; break;
|
||||
default: throw ihipException(hipErrorInvalidMemcpyDirection);
|
||||
};
|
||||
}
|
||||
|
||||
|
||||
void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
|
||||
{
|
||||
ihipDevice_t *device = this->getDevice();
|
||||
|
||||
if (device == NULL) {
|
||||
throw ihipException(hipErrorInvalidDevice);
|
||||
}
|
||||
|
||||
hc::accelerator acc;
|
||||
hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
|
||||
bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
|
||||
bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
|
||||
|
||||
|
||||
// Resolve default to a specific Kind so we know which algorithm to use:
|
||||
if (kind == hipMemcpyDefault) {
|
||||
bool srcInDeviceMem = (srcTracked && srcPtrInfo._isInDeviceMem);
|
||||
bool dstInDeviceMem = (dstTracked && dstPtrInfo._isInDeviceMem);
|
||||
kind = resolveMemcpyDirection(srcInDeviceMem, dstInDeviceMem);
|
||||
};
|
||||
|
||||
hsa_signal_t depSignal;
|
||||
|
||||
if ((kind == hipMemcpyHostToDevice) && (!srcTracked)) {
|
||||
int depSignalCnt = preCopyCommand(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);
|
||||
|
||||
if (HIP_PININPLACE) {
|
||||
device->_staging_buffer[0]->CopyHostToDevicePinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
|
||||
} else {
|
||||
device->_staging_buffer[0]->CopyHostToDevice(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);
|
||||
} else {
|
||||
// TODO - remove, slow path.
|
||||
tprintf(DB_COPY1, "H2D && ! srcTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
#if USE_AV_COPY
|
||||
_av.copy(src,dst,sizeBytes);
|
||||
#else
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
#endif
|
||||
}
|
||||
} else if ((kind == hipMemcpyDeviceToHost) && (!dstTracked)) {
|
||||
int depSignalCnt = preCopyCommand(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);
|
||||
|
||||
} else {
|
||||
// TODO - remove, slow path.
|
||||
tprintf(DB_COPY1, "D2H && !dstTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
#if USE_AV_COPY
|
||||
_av.copy(src, dst, sizeBytes);
|
||||
#else
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
#endif
|
||||
}
|
||||
} else if (kind == hipMemcpyHostToHost) {
|
||||
int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyH2H);
|
||||
|
||||
if (depSignalCnt) {
|
||||
// host waits before doing host memory copy.
|
||||
hsa_signal_wait_acquire(depSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
|
||||
}
|
||||
tprintf(DB_COPY1, "H2H memcpy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
memcpy(dst, src, sizeBytes);
|
||||
|
||||
} else {
|
||||
// If not special case - these can all be handled by the hsa async copy:
|
||||
ihipCommand_t commandType;
|
||||
hsa_agent_t srcAgent, dstAgent;
|
||||
setCopyAgents(kind, &commandType, &srcAgent, &dstAgent);
|
||||
|
||||
int depSignalCnt = preCopyCommand(NULL, &depSignal, commandType);
|
||||
|
||||
// Get a completion signal:
|
||||
ihipSignal_t *ihipSignal = allocSignal();
|
||||
hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal;
|
||||
|
||||
hsa_signal_store_relaxed(copyCompleteSignal, 1);
|
||||
|
||||
tprintf(DB_COPY1, "HSA Async_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
|
||||
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, copyCompleteSignal);
|
||||
|
||||
// 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.
|
||||
} else {
|
||||
throw ihipException(hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
|
||||
{
|
||||
ihipDevice_t *device = this->getDevice();
|
||||
|
||||
if (device == NULL) {
|
||||
throw ihipException(hipErrorInvalidDevice);
|
||||
}
|
||||
|
||||
if (kind == hipMemcpyHostToHost) {
|
||||
tprintf (DB_COPY2, "Asyc: H2H with memcpy");
|
||||
|
||||
// TODO - consider if we want to perhaps use the GPU SDMA engines anyway, to avoid the host-side sync here and keep everything flowing on the GPU.
|
||||
/* As this is a CPU op, we need to wait until all
|
||||
the commands in current stream are finished.
|
||||
*/
|
||||
this->wait();
|
||||
|
||||
memcpy(dst, src, sizeBytes);
|
||||
|
||||
} else {
|
||||
bool trueAsync = true;
|
||||
|
||||
hc::accelerator acc;
|
||||
hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
|
||||
bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
|
||||
|
||||
|
||||
// "tracked" really indicates if the pointer's virtual address is available in the GPU address space.
|
||||
// If both pointers are not tracked, we need to fall back to a sync copy.
|
||||
if (!dstTracked || !srcTracked) {
|
||||
trueAsync = false;
|
||||
}
|
||||
|
||||
if (kind == hipMemcpyDefault) {
|
||||
bool srcInDeviceMem = (srcTracked && srcPtrInfo._isInDeviceMem);
|
||||
bool dstInDeviceMem = (dstTracked && dstPtrInfo._isInDeviceMem);
|
||||
kind = resolveMemcpyDirection(srcInDeviceMem, dstInDeviceMem);
|
||||
}
|
||||
|
||||
|
||||
|
||||
ihipSignal_t *ihip_signal = allocSignal();
|
||||
hsa_signal_store_relaxed(ihip_signal->_hsa_signal, 1);
|
||||
|
||||
|
||||
if(trueAsync == true){
|
||||
|
||||
ihipCommand_t commandType;
|
||||
hsa_agent_t srcAgent, dstAgent;
|
||||
setCopyAgents(kind, &commandType, &srcAgent, &dstAgent);
|
||||
|
||||
hsa_signal_t depSignal;
|
||||
int depSignalCnt = preCopyCommand(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);
|
||||
|
||||
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsa_signal);
|
||||
|
||||
|
||||
if (hsa_status == HSA_STATUS_SUCCESS) {
|
||||
if (HIP_LAUNCH_BLOCKING) {
|
||||
tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
|
||||
this->wait();
|
||||
}
|
||||
} else {
|
||||
// This path can be hit if src or dst point to unpinned host memory.
|
||||
// TODO-stream - does async-copy fall back to sync if input pointers are not pinned?
|
||||
throw ihipException(hipErrorInvalidValue);
|
||||
}
|
||||
} else {
|
||||
copySync(dst, src, sizeBytes, kind);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
// HCC-specific accessor functions:
|
||||
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidDevice
|
||||
*/
|
||||
//---
|
||||
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
ihipDevice_t *d = ihipGetDevice(deviceId);
|
||||
hipError_t err;
|
||||
if (d == NULL) {
|
||||
err = hipErrorInvalidDevice;
|
||||
} else {
|
||||
*acc = d->_acc;
|
||||
err = hipSuccess;
|
||||
}
|
||||
return ihipLogStatus(err);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
if (stream == hipStreamNull ) {
|
||||
ihipDevice_t *device = ihipGetTlsDefaultDevice();
|
||||
stream = device->_default_stream;
|
||||
}
|
||||
|
||||
*av = &(stream->_av);
|
||||
|
||||
hipError_t err = hipSuccess;
|
||||
return ihipLogStatus(err);
|
||||
}
|
||||
|
||||
// TODO - review signal / error reporting code.
|
||||
// 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.
|
||||
|
||||
파일 크기가 너무 크기때문에 변경 상태를 표시하지 않습니다.
Diff 로드
+7
-296
@@ -1,4 +1,9 @@
|
||||
|
||||
#include "hip_runtime.h"
|
||||
#include "hcc_detail/hip_hcc.h"
|
||||
#include "hcc_detail/trace_helper.h"
|
||||
#include <hsa.h>
|
||||
#include <hc_am.hpp>
|
||||
#include <hsa_ext_amd.h>
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
// Memory
|
||||
@@ -85,79 +90,6 @@ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsi
|
||||
|
||||
|
||||
|
||||
// kernel for launching memcpy operations:
|
||||
template <typename T>
|
||||
hc::completion_future
|
||||
ihipMemcpyKernel(hipStream_t stream, T * c, const T * a, size_t sizeBytes)
|
||||
{
|
||||
int wg = std::min((unsigned)8, stream->getDevice()->_compute_units);
|
||||
const int threads_per_wg = 256;
|
||||
|
||||
int threads = wg * threads_per_wg;
|
||||
if (threads > sizeBytes) {
|
||||
threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;
|
||||
}
|
||||
|
||||
|
||||
hc::extent<1> ext(threads);
|
||||
auto ext_tile = ext.tile(threads_per_wg);
|
||||
|
||||
hc::completion_future cf =
|
||||
hc::parallel_for_each(
|
||||
stream->_av,
|
||||
ext_tile,
|
||||
[=] (hc::tiled_index<1> idx)
|
||||
__attribute__((hc))
|
||||
{
|
||||
int offset = amp_get_global_id(0);
|
||||
// TODO-HCC - change to hc_get_local_size()
|
||||
int stride = amp_get_local_size(0) * hc_get_num_groups(0) ;
|
||||
|
||||
for (int i=offset; i<sizeBytes; i+=stride) {
|
||||
c[i] = a[i];
|
||||
}
|
||||
});
|
||||
|
||||
return cf;
|
||||
}
|
||||
|
||||
|
||||
// kernel for launching memset operations:
|
||||
template <typename T>
|
||||
hc::completion_future
|
||||
ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes)
|
||||
{
|
||||
int wg = std::min((unsigned)8, stream->getDevice()->_compute_units);
|
||||
const int threads_per_wg = 256;
|
||||
|
||||
int threads = wg * threads_per_wg;
|
||||
if (threads > sizeBytes) {
|
||||
threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;
|
||||
}
|
||||
|
||||
|
||||
hc::extent<1> ext(threads);
|
||||
auto ext_tile = ext.tile(threads_per_wg);
|
||||
|
||||
hc::completion_future cf =
|
||||
hc::parallel_for_each(
|
||||
stream->_av,
|
||||
ext_tile,
|
||||
[=] (hc::tiled_index<1> idx)
|
||||
__attribute__((hc))
|
||||
{
|
||||
int offset = amp_get_global_id(0);
|
||||
// TODO-HCC - change to hc_get_local_size()
|
||||
int stride = amp_get_local_size(0) * hc_get_num_groups(0) ;
|
||||
|
||||
for (int i=offset; i<sizeBytes; i+=stride) {
|
||||
ptr[i] = val;
|
||||
}
|
||||
});
|
||||
|
||||
return cf;
|
||||
|
||||
}
|
||||
|
||||
//---
|
||||
/**
|
||||
@@ -344,228 +276,6 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
// Resolve hipMemcpyDefault to a known type.
|
||||
unsigned ihipStream_t::resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem)
|
||||
{
|
||||
hipMemcpyKind kind = hipMemcpyDefault;
|
||||
|
||||
if (!srcInDeviceMem && !dstInDeviceMem) {
|
||||
kind = hipMemcpyHostToHost;
|
||||
} else if (!srcInDeviceMem && dstInDeviceMem) {
|
||||
kind = hipMemcpyHostToDevice;
|
||||
} else if (srcInDeviceMem && !dstInDeviceMem) {
|
||||
kind = hipMemcpyDeviceToHost;
|
||||
} else if (srcInDeviceMem && dstInDeviceMem) {
|
||||
kind = hipMemcpyDeviceToDevice;
|
||||
}
|
||||
|
||||
assert (kind != hipMemcpyDefault);
|
||||
|
||||
return kind;
|
||||
}
|
||||
|
||||
|
||||
// Setup the copyCommandType and the copy agents (for hsa_amd_memory_async_copy)
|
||||
void ihipStream_t::setCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent)
|
||||
{
|
||||
ihipDevice_t *device = this->getDevice();
|
||||
hsa_agent_t deviceAgent = device->_hsa_agent;
|
||||
|
||||
switch (kind) {
|
||||
case hipMemcpyHostToHost : *commandType = ihipCommandCopyH2H; *srcAgent=g_cpu_agent; *dstAgent=g_cpu_agent; break;
|
||||
case hipMemcpyHostToDevice : *commandType = ihipCommandCopyH2D; *srcAgent=g_cpu_agent; *dstAgent=deviceAgent; break;
|
||||
case hipMemcpyDeviceToHost : *commandType = ihipCommandCopyD2H; *srcAgent=deviceAgent; *dstAgent=g_cpu_agent; break;
|
||||
case hipMemcpyDeviceToDevice : *commandType = ihipCommandCopyD2D; *srcAgent=deviceAgent; *dstAgent=deviceAgent; break;
|
||||
default: throw ihipException(hipErrorInvalidMemcpyDirection);
|
||||
};
|
||||
}
|
||||
|
||||
|
||||
void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
|
||||
{
|
||||
ihipDevice_t *device = this->getDevice();
|
||||
|
||||
if (device == NULL) {
|
||||
throw ihipException(hipErrorInvalidDevice);
|
||||
}
|
||||
|
||||
hc::accelerator acc;
|
||||
hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
|
||||
bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
|
||||
bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
|
||||
|
||||
|
||||
// Resolve default to a specific Kind so we know which algorithm to use:
|
||||
if (kind == hipMemcpyDefault) {
|
||||
bool srcInDeviceMem = (srcTracked && srcPtrInfo._isInDeviceMem);
|
||||
bool dstInDeviceMem = (dstTracked && dstPtrInfo._isInDeviceMem);
|
||||
kind = resolveMemcpyDirection(srcInDeviceMem, dstInDeviceMem);
|
||||
};
|
||||
|
||||
hsa_signal_t depSignal;
|
||||
|
||||
if ((kind == hipMemcpyHostToDevice) && (!srcTracked)) {
|
||||
int depSignalCnt = preCopyCommand(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);
|
||||
|
||||
if (HIP_PININPLACE) {
|
||||
device->_staging_buffer[0]->CopyHostToDevicePinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
|
||||
} else {
|
||||
device->_staging_buffer[0]->CopyHostToDevice(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);
|
||||
} else {
|
||||
// TODO - remove, slow path.
|
||||
tprintf(DB_COPY1, "H2D && ! srcTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
#if USE_AV_COPY
|
||||
_av.copy(src,dst,sizeBytes);
|
||||
#else
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
#endif
|
||||
}
|
||||
} else if ((kind == hipMemcpyDeviceToHost) && (!dstTracked)) {
|
||||
int depSignalCnt = preCopyCommand(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);
|
||||
|
||||
} else {
|
||||
// TODO - remove, slow path.
|
||||
tprintf(DB_COPY1, "D2H && !dstTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
#if USE_AV_COPY
|
||||
_av.copy(src, dst, sizeBytes);
|
||||
#else
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
#endif
|
||||
}
|
||||
} else if (kind == hipMemcpyHostToHost) {
|
||||
int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyH2H);
|
||||
|
||||
if (depSignalCnt) {
|
||||
// host waits before doing host memory copy.
|
||||
hsa_signal_wait_acquire(depSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
|
||||
}
|
||||
tprintf(DB_COPY1, "H2H memcpy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
memcpy(dst, src, sizeBytes);
|
||||
|
||||
} else {
|
||||
// If not special case - these can all be handled by the hsa async copy:
|
||||
ihipCommand_t commandType;
|
||||
hsa_agent_t srcAgent, dstAgent;
|
||||
setCopyAgents(kind, &commandType, &srcAgent, &dstAgent);
|
||||
|
||||
int depSignalCnt = preCopyCommand(NULL, &depSignal, commandType);
|
||||
|
||||
// Get a completion signal:
|
||||
ihipSignal_t *ihipSignal = allocSignal();
|
||||
hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal;
|
||||
|
||||
hsa_signal_store_relaxed(copyCompleteSignal, 1);
|
||||
|
||||
tprintf(DB_COPY1, "HSA Async_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
|
||||
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, copyCompleteSignal);
|
||||
|
||||
// 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.
|
||||
} else {
|
||||
throw ihipException(hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
|
||||
{
|
||||
ihipDevice_t *device = this->getDevice();
|
||||
|
||||
if (device == NULL) {
|
||||
throw ihipException(hipErrorInvalidDevice);
|
||||
}
|
||||
|
||||
if (kind == hipMemcpyHostToHost) {
|
||||
tprintf (DB_COPY2, "Asyc: H2H with memcpy");
|
||||
|
||||
// TODO - consider if we want to perhaps use the GPU SDMA engines anyway, to avoid the host-side sync here and keep everything flowing on the GPU.
|
||||
/* As this is a CPU op, we need to wait until all
|
||||
the commands in current stream are finished.
|
||||
*/
|
||||
this->wait();
|
||||
|
||||
memcpy(dst, src, sizeBytes);
|
||||
|
||||
} else {
|
||||
bool trueAsync = true;
|
||||
|
||||
hc::accelerator acc;
|
||||
hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0);
|
||||
bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
|
||||
bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
|
||||
|
||||
|
||||
// "tracked" really indicates if the pointer's virtual address is available in the GPU address space.
|
||||
// If both pointers are not tracked, we need to fall back to a sync copy.
|
||||
if (!dstTracked || !srcTracked) {
|
||||
trueAsync = false;
|
||||
}
|
||||
|
||||
if (kind == hipMemcpyDefault) {
|
||||
bool srcInDeviceMem = (srcTracked && srcPtrInfo._isInDeviceMem);
|
||||
bool dstInDeviceMem = (dstTracked && dstPtrInfo._isInDeviceMem);
|
||||
kind = resolveMemcpyDirection(srcInDeviceMem, dstInDeviceMem);
|
||||
}
|
||||
|
||||
|
||||
|
||||
ihipSignal_t *ihip_signal = allocSignal();
|
||||
hsa_signal_store_relaxed(ihip_signal->_hsa_signal, 1);
|
||||
|
||||
|
||||
if(trueAsync == true){
|
||||
|
||||
ihipCommand_t commandType;
|
||||
hsa_agent_t srcAgent, dstAgent;
|
||||
setCopyAgents(kind, &commandType, &srcAgent, &dstAgent);
|
||||
|
||||
hsa_signal_t depSignal;
|
||||
int depSignalCnt = preCopyCommand(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);
|
||||
|
||||
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsa_signal);
|
||||
|
||||
|
||||
if (hsa_status == HSA_STATUS_SUCCESS) {
|
||||
if (HIP_LAUNCH_BLOCKING) {
|
||||
tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
|
||||
this->wait();
|
||||
}
|
||||
} else {
|
||||
// This path can be hit if src or dst point to unpinned host memory.
|
||||
// TODO-stream - does async-copy fall back to sync if input pointers are not pinned?
|
||||
throw ihipException(hipErrorInvalidValue);
|
||||
}
|
||||
} else {
|
||||
copySync(dst, src, sizeBytes, kind);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind)
|
||||
{
|
||||
@@ -779,3 +489,4 @@ hipError_t hipFreeHost(void* ptr)
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
@@ -1,60 +0,0 @@
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipDriverGetVersion(int *driverVersion)
|
||||
{
|
||||
HIP_INIT_API(driverVersion);
|
||||
|
||||
if (driverVersion) {
|
||||
*driverVersion = 4;
|
||||
}
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
// HCC-specific accessor functions:
|
||||
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidDevice
|
||||
*/
|
||||
//---
|
||||
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
ihipDevice_t *d = ihipGetDevice(deviceId);
|
||||
hipError_t err;
|
||||
if (d == NULL) {
|
||||
err = hipErrorInvalidDevice;
|
||||
} else {
|
||||
*acc = d->_acc;
|
||||
err = hipSuccess;
|
||||
}
|
||||
return ihipLogStatus(err);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
if (stream == hipStreamNull ) {
|
||||
ihipDevice_t *device = ihipGetTlsDefaultDevice();
|
||||
stream = device->_default_stream;
|
||||
}
|
||||
|
||||
*av = &(stream->_av);
|
||||
|
||||
hipError_t err = hipSuccess;
|
||||
return ihipLogStatus(err);
|
||||
}
|
||||
@@ -1,3 +1,6 @@
|
||||
#include "hip_runtime.h"
|
||||
#include "hcc_detail/hip_hcc.h"
|
||||
#include "hcc_detail/trace_helper.h"
|
||||
|
||||
/**
|
||||
* @warning HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P
|
||||
@@ -56,3 +59,21 @@ hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int
|
||||
// HCC has a unified memory architecture so device specifiers are not required.
|
||||
return hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream);
|
||||
};
|
||||
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipDriverGetVersion(int *driverVersion)
|
||||
{
|
||||
HIP_INIT_API(driverVersion);
|
||||
|
||||
if (driverVersion) {
|
||||
*driverVersion = 4;
|
||||
}
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
|
||||
@@ -1,3 +1,7 @@
|
||||
#include "hip_runtime.h"
|
||||
#include "hcc_detail/hip_hcc.h"
|
||||
#include "hcc_detail/trace_helper.h"
|
||||
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
@@ -116,3 +120,6 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags)
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
@@ -40,12 +40,12 @@ if (${HIP_PLATFORM} STREQUAL "hcc")
|
||||
#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)
|
||||
include_directories(${HSA_PATH}/include)
|
||||
# hip_hcc.o:
|
||||
#add_library(hip_hcc STATIC ${HIP_PATH}/src/hip_hcc.cpp )
|
||||
add_library(hip_hcc STATIC ${HIP_PATH}/src/hip_hcc.cpp ${HIP_PATH}/src/hip_error.cpp ${HIP_PATH}/src/hip_device.cpp ${HIP_PATH}/src/hip_stream.cpp ${HIP_PATH}/src/hip_event.cpp ${HIP_PATH}/src/hip_memory.cpp ${HIP_PATH}/src/hip_peer.cpp ${HIP_PATH}/src/staging_buffer.cpp)
|
||||
target_include_directories(hip_hcc PRIVATE ${HSA_PATH}/include ${CODEXL_SDK_ATAL_PATH}/include)
|
||||
|
||||
# This will create a subdir "hip_hcc" in the test build directory
|
||||
# Any changes to hip_hcc source will be detected and force the library and then the tests to be rebuilt.
|
||||
add_subdirectory(${HIP_PATH} hip_hcc)
|
||||
link_directories(${CMAKE_CURRENT_BINARY_DIR}/hip_hcc) # search the local hip_hcc for libhip_hcc.a
|
||||
|
||||
elseif (${HIP_PLATFORM} STREQUAL "nvcc")
|
||||
MESSAGE ("HIP_PLATFORM=nvcc")
|
||||
@@ -73,7 +73,7 @@ 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)
|
||||
# target_link_libraries(${exe} hip_hcc)
|
||||
else()
|
||||
add_executable (${exe} ${cpp} ${ARGN} $<TARGET_OBJECTS:test_common> )
|
||||
endif()
|
||||
|
||||
새 이슈에서 참조
사용자 차단