AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

Split files based on functionality and changed cmake file

فهرست منبع 
[ROCm/clr commit: 1c55019322]
This commit is contained in:
Aditya Atluri
2016-03-24 04:57:30 -05:00
والد 539f1a419f
کامیت bd9979eda6
15فایلهای تغییر یافته به همراه527 افزوده شده و 3185 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
projects/clr/hipamd/CMakeLists.txt
+8 -1
مشاهده پرونده
@@ -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})
projects/clr/hipamd/include/hcc_detail/hip_hcc.h
+120 -29
مشاهده پرونده
@@ -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
projects/clr/hipamd/include/hcc_detail/hip_runtime.h
+1 -4
مشاهده پرونده
@@ -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__
projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h
+2 -1
مشاهده پرونده
@@ -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.");
projects/clr/hipamd/include/hip_runtime_api.h
+1 -1
مشاهده پرونده
@@ -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);
}
projects/clr/hipamd/src/hip_device.cpp
+6 -1
مشاهده پرونده
@@ -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);
}
projects/clr/hipamd/src/hip_error.cpp
+4 -25
مشاهده پرونده
@@ -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)
{
projects/clr/hipamd/src/hip_event.cpp
+3 -19
مشاهده پرونده
@@ -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)
{
projects/clr/hipamd/src/hip_hcc.cpp
+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.
projects/clr/hipamd/src/hip_hcc2.cpp
-2722
مشاهده پرونده
تفاوت فایلی نمایش داده نمی شود زیرا این فایل بسیار بزرگ است Diff را بارگزاری کن
projects/clr/hipamd/src/hip_memory.cpp
+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)
}
projects/clr/hipamd/src/hip_misc.cpp
-60
مشاهده پرونده
@@ -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);
}
projects/clr/hipamd/src/hip_p2p.cpp → projects/clr/hipamd/src/hip_peer.cpp
+21
مشاهده پرونده
@@ -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);
}
projects/clr/hipamd/src/hip_stream.cpp
+7
مشاهده پرونده
@@ -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);
}
}
projects/clr/hipamd/tests/src/CMakeLists.txt
+6 -6
مشاهده پرونده
@@ -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")
@@ -71,7 +71,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()