AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

Merge pull request #464 from gargrahul/fix_memcpy2d_pinned_mem_case

Bläddra i källkod 
Fixed memcpy2D for pinned memory case using 2D kernel

[ROCm/clr commit: 7042fe6067]
This commit is contained in:
Maneesh Gupta
2018-05-22 10:42:28 +05:30
committad av GitHub
förälder e72516bd2f 54497d1e2b
incheckning 335dad2feb
1 ändrade filer med 73 tillägg och 70 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_memory.cpp
+73 -70
Visa fil
@@ -1212,7 +1212,6 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte
return ihipLogStatus(e);
}
hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
@@ -1231,7 +1230,6 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
return ihipLogStatus(e);
}
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind, stream);
@@ -1239,7 +1237,6 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, kind, stream));
}
hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
@@ -1262,51 +1259,6 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToHost, stream));
}
// TODO - review and optimize
hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind) {
if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue;
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hipError_t e = hipSuccess;
if((width == dpitch) && (width == spitch)) {
stream->locked_copySync((void*)dst, (void*)src, width*height, kind, false);
} else {
try {
for (int i = 0; i < height; ++i) {
stream->locked_copySync((unsigned char*)dst + i * dpitch,
(unsigned char*)src + i * spitch, width, kind);
}
} catch (ihipException& ex) {
e = ex._code;
}
}
return e;
}
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
hipError_t e = hipSuccess;
e = ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind);
return ihipLogStatus(e);
}
hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), pCopy);
hipError_t e = hipSuccess;
if (pCopy == nullptr) {
e = hipErrorInvalidValue;
}
e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch,
pCopy->widthInBytes, pCopy->height, hipMemcpyDefault);
return ihipLogStatus(e);
}
hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
size_t spitch, size_t width, size_t height, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind);
@@ -1538,16 +1490,15 @@ inline const T& clamp_integer(const T& x, const T& lower, const T& upper) {
return std::min(upper, std::max(x, lower));
}
template <uint32_t block_dim, typename T>
__global__ void hip_copy_n(T* dst, const T* src, size_t n) {
const uint32_t grid_dim = gridDim.x * blockDim.x;
template <typename T>
__global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, size_t destPitch, size_t srcPitch) {
size_t idx = blockIdx.x * block_dim + threadIdx.x;
while (idx < n) {
// __builtin_memcpy(reinterpret_cast<void*>(dst+idx), reinterpret_cast<const void*>(src+idx),
// sizeof(T));
dst[idx] = src[idx];
idx += grid_dim;
size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
size_t idy = blockIdx.y * blockDim.y + threadIdx.y;
if((idx < width) && (idy < height)){
T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch);
T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch);
dstPtr[idx] = srcPtr[idx];
}
}
} // namespace
@@ -1563,13 +1514,12 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) {
}
template <typename T>
void ihipMemcpyKernel(hipStream_t stream, T* dst, const T* src, size_t sizeBytes) {
static constexpr uint32_t block_dim_ = 256;
const uint32_t grid_dim = clamp_integer<size_t>(sizeBytes / block_dim_, 1, UINT32_MAX);
hipLaunchKernelGGL(hip_copy_n<block_dim_>, dim3(grid_dim), dim3{block_dim_}, 0u, stream, dst, src,
sizeBytes);
void ihipMemcpy2dKernel(hipStream_t stream, T* dst, const T* src, size_t width, size_t height, size_t destPitch, size_t srcPitch) {
size_t threadsPerBlock = 16;
uint32_t grid_dim_x = clamp_integer<size_t>( (width+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX);
uint32_t grid_dim_y = clamp_integer<size_t>( (height+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX);
hipLaunchKernelGGL(hip_copy2d_n, dim3(grid_dim_x,grid_dim_y), dim3(threadsPerBlock,threadsPerBlock), 0u, stream, dst, src,
width, height, destPitch, srcPitch);
}
typedef enum ihipMemsetDataType {
@@ -1647,6 +1597,50 @@ int isLockedPointer(const void *ptr)
return isLocked;
}
// TODO - review and optimize
hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind) {
if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue;
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
int isLocked = 0;
if(kind == hipMemcpyHostToDevice) {
isLocked = isLockedPointer(src);
} else if(kind == hipMemcpyDeviceToHost) {
isLocked = isLockedPointer(dst);
}
hc::completion_future marker;
hipError_t e = hipSuccess;
if((width == dpitch) && (width == spitch)) {
stream->locked_copySync((void*)dst, (void*)src, width*height, kind, false);
} else {
try {
if(isLocked) {
for (int i = 0; i < height; ++i)
stream->locked_copySync((unsigned char*)dst + i * dpitch,
(unsigned char*)src + i * spitch, width, kind);
} else {
ihipMemcpy2dKernel<uint8_t> (stream, static_cast<uint8_t*> (dst), static_cast<const uint8_t*> (src), width, height, dpitch, spitch);
stream->locked_wait();
}
} catch (ihipException& ex) {
e = ex._code;
}
}
return e;
}
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
hipError_t e = hipSuccess;
e = ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind);
return ihipLogStatus(e);
}
hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream);
@@ -1662,14 +1656,12 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
hip_internal::memcpyAsync(dst, src, width*height, kind, stream);
} else {
try {
for (int i = 0; i < height; ++i) {
if(!isLocked) {
if(!isLocked){
for (int i = 0; i < height; ++i)
e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch,
(unsigned char*)src + i * spitch, width, kind, stream);
} else{
size_t sizeBytes = width*height;
ihipMemcpyKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), static_cast<const uint32_t*> (src), sizeBytes/sizeof(uint32_t));
}
} else{
ihipMemcpy2dKernel<uint8_t> (stream, static_cast<uint8_t*> (dst), static_cast<const uint8_t*> (src), width, height, dpitch, spitch);
}
} catch (ihipException& ex) {
e = ex._code;
@@ -1679,6 +1671,17 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
return ihipLogStatus(e);
}
hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), pCopy);
hipError_t e = hipSuccess;
if (pCopy == nullptr) {
e = hipErrorInvalidValue;
}
e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch,
pCopy->widthInBytes, pCopy->height, hipMemcpyDefault);
return ihipLogStatus(e);
}
// TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes, stream);