Merge branch 'privatestaging' into p2p
This commit is contained in:
Ben Sander
@@ -799,4 +799,412 @@ __device__ float __dsqrt_rn(double x) {return hc::fast_math::sqrt(x); };
|
||||
__device__ float __dsqrt_ru(double x) {return hc::fast_math::sqrt(x); };
|
||||
__device__ float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); };
|
||||
|
||||
__HIP_DEVICE__ char1 make_char1(signed char x)
|
||||
{
|
||||
char1 c1;
|
||||
c1.x = x;
|
||||
return c1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ char2 make_char2(signed char x, signed char y)
|
||||
{
|
||||
char2 c2;
|
||||
c2.x = x;
|
||||
c2.y = y;
|
||||
return c2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ char3 make_char3(signed char x, signed char y, signed char z)
|
||||
{
|
||||
char3 c3;
|
||||
c3.x = x;
|
||||
c3.y = y;
|
||||
c3.z = z;
|
||||
return c3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ char4 make_char4(signed char x, signed char y, signed char z, signed char w)
|
||||
{
|
||||
char4 c4;
|
||||
c4.x = x;
|
||||
c4.y = y;
|
||||
c4.z = z;
|
||||
c4.w = w;
|
||||
return c4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ short1 make_short1(short x)
|
||||
{
|
||||
short1 s1;
|
||||
s1.x = x;
|
||||
return s1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ short2 make_short2(short x, short y)
|
||||
{
|
||||
short2 s2;
|
||||
s2.x = x;
|
||||
s2.y = y;
|
||||
return s2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ short3 make_short3(short x, short y, short z)
|
||||
{
|
||||
short3 s3;
|
||||
s3.x = x;
|
||||
s3.y = y;
|
||||
s3.z = z;
|
||||
return s3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ short4 make_short4(short x, short y, short z, short w)
|
||||
{
|
||||
short4 s4;
|
||||
s4.x = x;
|
||||
s4.y = y;
|
||||
s4.z = z;
|
||||
s4.w = w;
|
||||
return s4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ int1 make_int1(int x)
|
||||
{
|
||||
int1 i1;
|
||||
i1.x = x;
|
||||
return i1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ int2 make_int2(int x, int y)
|
||||
{
|
||||
int2 i2;
|
||||
i2.x = x;
|
||||
i2.y = y;
|
||||
return i2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ int3 make_int3(int x, int y, int z)
|
||||
{
|
||||
int3 i3;
|
||||
i3.x = x;
|
||||
i3.y = y;
|
||||
i3.z = z;
|
||||
return i3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ int4 make_int4(int x, int y, int z, int w)
|
||||
{
|
||||
int4 i4;
|
||||
i4.x = x;
|
||||
i4.y = y;
|
||||
i4.z = z;
|
||||
i4.w = w;
|
||||
return i4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ long1 make_long1(long x)
|
||||
{
|
||||
long1 l1;
|
||||
l1.x = x;
|
||||
return l1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ long2 make_long2(long x, long y)
|
||||
{
|
||||
long2 l2;
|
||||
l2.x = x;
|
||||
l2.y = y;
|
||||
return l2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ long3 make_long3(long x, long y, long z)
|
||||
{
|
||||
long3 l3;
|
||||
l3.x = x;
|
||||
l3.y = y;
|
||||
l3.z = z;
|
||||
return l3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ long4 make_long4(long x, long y, long z, long w)
|
||||
{
|
||||
long4 l4;
|
||||
l4.x = x;
|
||||
l4.y = y;
|
||||
l4.z = z;
|
||||
l4.w = w;
|
||||
return l4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ longlong1 make_longlong1(long long x)
|
||||
{
|
||||
longlong1 l1;
|
||||
l1.x = x;
|
||||
return l1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ longlong2 make_longlong2(long long x, long long y)
|
||||
{
|
||||
longlong2 l2;
|
||||
l2.x = x;
|
||||
l2.y = y;
|
||||
return l2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ longlong3 make_longlong3(long long x, long long y, long long z)
|
||||
{
|
||||
longlong3 l3;
|
||||
l3.x = x;
|
||||
l3.y = y;
|
||||
l3.z = z;
|
||||
return l3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ longlong4 make_longlong4(long long x, long long y, long long z, long long w)
|
||||
{
|
||||
longlong4 l4;
|
||||
l4.x = x;
|
||||
l4.y = y;
|
||||
l4.z = z;
|
||||
l4.w = w;
|
||||
return l4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uchar1 make_uchar1(unsigned char x)
|
||||
{
|
||||
uchar1 c1;
|
||||
c1.x = x;
|
||||
return c1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uchar2 make_uchar2(unsigned char x, unsigned char y)
|
||||
{
|
||||
uchar2 c2;
|
||||
c2.x = x;
|
||||
c2.y = y;
|
||||
return c2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uchar3 make_uchar3(unsigned char x, unsigned char y, unsigned char z)
|
||||
{
|
||||
uchar3 c3;
|
||||
c3.x = x;
|
||||
c3.y = y;
|
||||
c3.z = z;
|
||||
return c3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uchar4 make_uchar4(unsigned char x, unsigned char y, unsigned char z, unsigned char w)
|
||||
{
|
||||
uchar4 c4;
|
||||
c4.x = x;
|
||||
c4.y = y;
|
||||
c4.z = z;
|
||||
c4.w = w;
|
||||
return c4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ushort1 make_ushort1(unsigned short x)
|
||||
{
|
||||
ushort1 s1;
|
||||
s1.x = x;
|
||||
return s1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ushort2 make_ushort2(unsigned short x, unsigned short y)
|
||||
{
|
||||
ushort2 s2;
|
||||
s2.x = x;
|
||||
s2.y = y;
|
||||
return s2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ushort3 make_ushort3(unsigned short x, unsigned short y, unsigned short z)
|
||||
{
|
||||
ushort3 s3;
|
||||
s3.x = x;
|
||||
s3.y = y;
|
||||
s3.z = z;
|
||||
return s3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ushort4 make_ushort4(unsigned short x, unsigned short y, unsigned short z, unsigned short w)
|
||||
{
|
||||
ushort4 s4;
|
||||
s4.x = x;
|
||||
s4.y = y;
|
||||
s4.z = z;
|
||||
s4.w = w;
|
||||
return s4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uint1 make_uint1(unsigned int x)
|
||||
{
|
||||
uint1 i1;
|
||||
i1.x = x;
|
||||
return i1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uint2 make_uint2(unsigned int x, unsigned int y)
|
||||
{
|
||||
uint2 i2;
|
||||
i2.x = x;
|
||||
i2.y = y;
|
||||
return i2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uint3 make_uint3(unsigned int x, unsigned int y, unsigned int z)
|
||||
{
|
||||
uint3 i3;
|
||||
i3.x = x;
|
||||
i3.y = y;
|
||||
i3.z = z;
|
||||
return i3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uint4 make_uint4(unsigned int x, unsigned int y, unsigned int z, unsigned int w)
|
||||
{
|
||||
uint4 i4;
|
||||
i4.x = x;
|
||||
i4.y = y;
|
||||
i4.z = z;
|
||||
i4.w = w;
|
||||
return i4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulong1 make_ulong1(unsigned long x)
|
||||
{
|
||||
ulong1 l1;
|
||||
l1.x = x;
|
||||
return l1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulong2 make_ulong2(unsigned long x, unsigned long y)
|
||||
{
|
||||
ulong2 l2;
|
||||
l2.x = x;
|
||||
l2.y = y;
|
||||
return l2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulong3 make_ulong3(unsigned long x, unsigned long y, unsigned long z)
|
||||
{
|
||||
ulong3 l3;
|
||||
l3.x = x;
|
||||
l3.y = y;
|
||||
l3.z = z;
|
||||
return l3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulong4 make_ulong4(unsigned long x, unsigned long y, unsigned long z, unsigned long w)
|
||||
{
|
||||
ulong4 l4;
|
||||
l4.x = x;
|
||||
l4.y = y;
|
||||
l4.z = z;
|
||||
l4.w = w;
|
||||
return l4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulonglong1 make_ulonglong1(unsigned long long x)
|
||||
{
|
||||
ulonglong1 l1;
|
||||
l1.x = x;
|
||||
return l1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulonglong2 make_ulonglong2(unsigned long long x, unsigned long long y)
|
||||
{
|
||||
ulonglong2 l2;
|
||||
l2.x = x;
|
||||
l2.y = y;
|
||||
return l2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulonglong3 make_ulonglong3(unsigned long long x, unsigned long long y, unsigned long long z)
|
||||
{
|
||||
ulonglong3 l3;
|
||||
l3.x = x;
|
||||
l3.y = y;
|
||||
l3.z = z;
|
||||
return l3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulonglong4 make_ulonglong4(unsigned long long x, unsigned long long y, unsigned long long z, unsigned long long w)
|
||||
{
|
||||
ulonglong4 l4;
|
||||
l4.x = x;
|
||||
l4.y = y;
|
||||
l4.z = z;
|
||||
l4.w = w;
|
||||
return l4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ float1 make_float1(float x)
|
||||
{
|
||||
float1 f1;
|
||||
f1.x = x;
|
||||
return f1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ float2 make_float2(float x, float y)
|
||||
{
|
||||
float2 f2;
|
||||
f2.x = x;
|
||||
f2.y = y;
|
||||
return f2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ float3 make_float3(float x, float y, float z)
|
||||
{
|
||||
float3 f3;
|
||||
f3.x = x;
|
||||
f3.y = y;
|
||||
f3.z = z;
|
||||
return f3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ float4 make_float4(float x, float y, float z, float w)
|
||||
{
|
||||
float4 f4;
|
||||
f4.x = x;
|
||||
f4.y = y;
|
||||
f4.z = z;
|
||||
f4.w = w;
|
||||
return f4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ double1 make_double1(double x)
|
||||
{
|
||||
double1 d1;
|
||||
d1.x = x;
|
||||
return d1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ double2 make_double2(double x, double y)
|
||||
{
|
||||
double2 d2;
|
||||
d2.x = x;
|
||||
d2.y = y;
|
||||
return d2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ double3 make_double3(double x, double y, double z)
|
||||
{
|
||||
double3 d3;
|
||||
d3.x = x;
|
||||
d3.y = y;
|
||||
d3.z = z;
|
||||
return d3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ double4 make_double4(double x, double y, double z, double w)
|
||||
{
|
||||
double4 d4;
|
||||
d4.x = x;
|
||||
d4.y = y;
|
||||
d4.z = z;
|
||||
d4.w = w;
|
||||
return d4;
|
||||
}
|
||||
|
||||
|
||||
|
||||
+78
-33
@@ -65,8 +65,8 @@ int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU ident
|
||||
|
||||
//---
|
||||
// Chicken bits for disabling functionality to work around potential issues:
|
||||
int HIP_DISABLE_HW_KERNEL_DEP = 1;
|
||||
int HIP_DISABLE_HW_COPY_DEP = 1;
|
||||
int HIP_DISABLE_HW_KERNEL_DEP = 0;
|
||||
int HIP_DISABLE_HW_COPY_DEP = 0;
|
||||
|
||||
thread_local int tls_defaultDevice = 0;
|
||||
thread_local hipError_t tls_lastHipError = hipSuccess;
|
||||
@@ -181,6 +181,8 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty
|
||||
// Reset the stream to "empty" - next command will not set up an inpute dependency on any older signal.
|
||||
crit->_last_command_type = ihipCommandCopyH2D;
|
||||
crit->_last_copy_signal = NULL;
|
||||
|
||||
_depFutures.clear();
|
||||
}
|
||||
|
||||
|
||||
@@ -428,6 +430,9 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t
|
||||
needSync = 1;
|
||||
hsa_signal_t *hsaSignal = (static_cast<hsa_signal_t*> (crit->_last_kernel_future.get_native_handle()));
|
||||
if (hsaSignal) {
|
||||
// Keep reference to the kernel future in order to keep the
|
||||
// dependent signal alive.
|
||||
_depFutures.push_back(crit->_last_kernel_future);
|
||||
*waitSignal = * hsaSignal;
|
||||
} else {
|
||||
assert(0); // if NULL signal, and we return 1, hsa_amd_memory_copy_async will fail. Confirm this never happens.
|
||||
@@ -529,7 +534,7 @@ void ihipDevice_t::init(unsigned device_index, unsigned deviceCnt, hc::accelerat
|
||||
|
||||
locked_reset();
|
||||
|
||||
|
||||
|
||||
tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream);
|
||||
|
||||
hsa_region_t *pinnedHostRegion;
|
||||
@@ -1236,12 +1241,12 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
|
||||
|
||||
bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
|
||||
bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
|
||||
bool srcInDeviceMem = srcPtrInfo._isInDeviceMem;
|
||||
bool dstInDeviceMem = dstPtrInfo._isInDeviceMem;
|
||||
|
||||
// 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(srcTracked, dstTracked, srcPtrInfo._isInDeviceMem, dstPtrInfo._isInDeviceMem);
|
||||
kind = resolveMemcpyDirection(srcTracked, dstTracked, srcInDeviceMem, dstInDeviceMem);
|
||||
};
|
||||
|
||||
hsa_signal_t depSignal;
|
||||
@@ -1259,44 +1264,86 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
|
||||
|
||||
if (kind == hipMemcpyHostToDevice) {
|
||||
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);
|
||||
if(!srcTracked){
|
||||
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);
|
||||
}
|
||||
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(crit, true);
|
||||
} else {
|
||||
// TODO - remove, slow path.
|
||||
tprintf(DB_COPY1, "H2D && ! srcTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
// The copy waits for inputs and then completes before returning so can reset queue to empty:
|
||||
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);
|
||||
#if USE_AV_COPY
|
||||
_av.copy(src,dst,sizeBytes);
|
||||
_av.copy(src,dst,sizeBytes);
|
||||
#else
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
#endif
|
||||
}
|
||||
}else{
|
||||
hsa_agent_t dstAgent = *(static_cast<hsa_agent_t*>(dstPtrInfo._acc.get_hsa_agent()));
|
||||
hsa_agent_t srcAgent = *(static_cast<hsa_agent_t*>(srcPtrInfo._acc.get_hsa_agent()));
|
||||
|
||||
ihipSignal_t *ihipSignal = allocSignal(crit);
|
||||
hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal;
|
||||
|
||||
hsa_signal_store_relaxed(copyCompleteSignal, 1);
|
||||
void *devPtrSrc = srcPtrInfo._devicePointer;
|
||||
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, devPtrSrc, 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(crit, ihipSignal); // wait for copy, and return to pool.
|
||||
} else {
|
||||
throw ihipException(hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
} else if (kind == hipMemcpyDeviceToHost) {
|
||||
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);
|
||||
if (!dstTracked){
|
||||
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 completes before returning so can reset queue to empty:
|
||||
this->wait(crit, true);
|
||||
|
||||
// The copy completes before returning so can reset queue to empty:
|
||||
this->wait(crit, true);
|
||||
|
||||
} else {
|
||||
} else {
|
||||
// TODO - remove, slow path.
|
||||
tprintf(DB_COPY1, "D2H && !dstTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
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);
|
||||
_av.copy(src, dst, sizeBytes);
|
||||
#else
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
#endif
|
||||
}
|
||||
}else{
|
||||
hsa_agent_t dstAgent = *(static_cast<hsa_agent_t*>(dstPtrInfo._acc.get_hsa_agent()));
|
||||
hsa_agent_t srcAgent = *(static_cast<hsa_agent_t*>(srcPtrInfo._acc.get_hsa_agent()));
|
||||
|
||||
ihipSignal_t *ihipSignal = allocSignal(crit);
|
||||
hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal;
|
||||
|
||||
hsa_signal_store_relaxed(copyCompleteSignal, 1);
|
||||
void *devPtrDst = dstPtrInfo._devicePointer;
|
||||
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(devPtrDst, 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(crit, ihipSignal); // wait for copy, and return to pool.
|
||||
} else {
|
||||
throw ihipException(hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
} else if (kind == hipMemcpyHostToHost) {
|
||||
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyH2H);
|
||||
@@ -1305,9 +1352,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
|
||||
// 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 ((kind == hipMemcpyDeviceToDevice) && !copyEngineCanSeeSrcAndDest) {
|
||||
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyP2P);
|
||||
if (HIP_STAGING_BUFFERS) {
|
||||
|
||||
@@ -0,0 +1,319 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include"hcc_detail/hip_ldg.h"
|
||||
#if __hcc_workweek__ >= 16164
|
||||
__device__ char __ldg(const char* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ char1 __ldg(const char1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ char2 __ldg(const char2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ char3 __ldg(const char3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ char4 __ldg(const char4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ signed char __ldg(const signed char* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ unsigned char __ldg(const unsigned char* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
__device__ short __ldg(const short* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ short1 __ldg(const short1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ short2 __ldg(const short2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ short3 __ldg(const short3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ short4 __ldg(const short4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ unsigned short __ldg(const unsigned short* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
__device__ int __ldg(const int* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ int1 __ldg(const int1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ int2 __ldg(const int2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ int3 __ldg(const int3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ int4 __ldg(const int4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ unsigned int __ldg(const unsigned int* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
__device__ long __ldg(const long* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ long1 __ldg(const long1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ long2 __ldg(const long2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ long3 __ldg(const long3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ long4 __ldg(const long4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ unsigned long __ldg(const unsigned long* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
__device__ long long __ldg(const long long* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ longlong1 __ldg(const longlong1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ longlong2 __ldg(const longlong2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ longlong3 __ldg(const longlong3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ longlong4 __ldg(const longlong4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ unsigned long long __ldg(const unsigned long long* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
__device__ uchar1 __ldg(const uchar1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ uchar2 __ldg(const uchar2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ uchar3 __ldg(const uchar3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ uchar4 __ldg(const uchar4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
__device__ ushort1 __ldg(const ushort1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ ushort2 __ldg(const ushort2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ ushort3 __ldg(const ushort3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ ushort4 __ldg(const ushort4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
__device__ uint1 __ldg(const uint1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ uint2 __ldg(const uint2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ uint3 __ldg(const uint3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ uint4 __ldg(const uint4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
__device__ ulonglong1 __ldg(const ulonglong1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ ulonglong2 __ldg(const ulonglong2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ ulonglong3 __ldg(const ulonglong3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ ulonglong4 __ldg(const ulonglong4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
__device__ float __ldg(const float* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ float1 __ldg(const float1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ float2 __ldg(const float2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ float3 __ldg(const float3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ float4 __ldg(const float4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
__device__ double __ldg(const double* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ double1 __ldg(const double1* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ double2 __ldg(const double2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ double3 __ldg(const double3* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
__device__ double4 __ldg(const double4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
@@ -253,7 +253,11 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags)
|
||||
for(int i=0;i<g_deviceCnt;i++){
|
||||
vecAcc.push_back(g_devices[i]._acc);
|
||||
}
|
||||
#if USE_HCC_LOCK_API
|
||||
am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0], vecAcc.size());
|
||||
#else
|
||||
am_status = AM_ERROR_MISC;
|
||||
#endif
|
||||
if(am_status == AM_SUCCESS){
|
||||
hip_status = hipSuccess;
|
||||
}else{
|
||||
@@ -276,7 +280,11 @@ hipError_t hipHostUnregister(void *hostPtr)
|
||||
if(hostPtr == NULL){
|
||||
hip_status = hipErrorInvalidValue;
|
||||
}else{
|
||||
#if USE_HCC_LOCK_API
|
||||
am_status_t am_status = hc::am_memory_host_unlock(device->_acc, hostPtr);
|
||||
#else
|
||||
am_status_t am_status = AM_ERROR_MISC;
|
||||
#endif
|
||||
if(am_status != AM_SUCCESS){
|
||||
hip_status = hipErrorHostMemoryNotRegistered;
|
||||
}
|
||||
|
||||
مرجع در شماره جدید
Block a user