From 2c3d1498d4e3b03d5a7a5aa763855dcd755ea738 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 23 May 2018 14:43:47 +0530 Subject: [PATCH 1/5] Optimize memcpy2D kernel use --- hipamd/src/hip_memory.cpp | 16 ++++++++++++---- 1 file changed, 12 insertions(+), 4 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index fc60dffe86..b76f7f0524 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -1495,10 +1495,18 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, size_t idx = blockIdx.x * blockDim.x + threadIdx.x; size_t idy = blockIdx.y * blockDim.y + threadIdx.y; - if((idx < width) && (idy < height)){ + size_t floorWidth = (width/sizeof(T)); + if((idx < floorWidth)){ T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); dstPtr[idx] = srcPtr[idx]; + } else { + size_t bytesToCopy = width - floorWidth; + uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch); + uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch); + for(int i =0 ; i < bytesToCopy ; i++) { + dstPtr[idx+i]= srcPtr[idx+i]; + } } } } // namespace @@ -1516,7 +1524,7 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { template 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( (width+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); + uint32_t grid_dim_x = clamp_integer( ((width/sizeof(T))+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); uint32_t grid_dim_y = clamp_integer( (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); @@ -1622,7 +1630,7 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch stream->locked_copySync((unsigned char*)dst + i * dpitch, (unsigned char*)src + i * spitch, width, kind); } else { - ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); + ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); stream->locked_wait(); } } catch (ihipException& ex) { @@ -1661,7 +1669,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch, (unsigned char*)src + i * spitch, width, kind, stream); } else{ - ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); + ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); } } catch (ihipException& ex) { e = ex._code; From 96b4618d26acfb7e21537ed8b48a77feb0305209 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 24 May 2018 08:27:24 +0530 Subject: [PATCH 2/5] Correct remaining bytes in copy 2d kernel --- hipamd/src/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index b76f7f0524..a1bc1416a0 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -1501,7 +1501,7 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); dstPtr[idx] = srcPtr[idx]; } else { - size_t bytesToCopy = width - floorWidth; + size_t bytesToCopy = width - (floorWidth * sizeof(T)); uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch); uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch); for(int i =0 ; i < bytesToCopy ; i++) { From 53a7c61e9b93d74067f363de5959b82712cee5b9 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 24 May 2018 17:00:12 +0530 Subject: [PATCH 3/5] Fix memcpy2d kernel dims --- hipamd/src/hip_memory.cpp | 26 ++++++++++++++------------ 1 file changed, 14 insertions(+), 12 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index a1bc1416a0..2f0c793033 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -1496,17 +1496,19 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, size_t idx = blockIdx.x * blockDim.x + threadIdx.x; size_t idy = blockIdx.y * blockDim.y + threadIdx.y; size_t floorWidth = (width/sizeof(T)); - if((idx < floorWidth)){ - T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); - T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); - dstPtr[idx] = srcPtr[idx]; - } else { - size_t bytesToCopy = width - (floorWidth * sizeof(T)); - uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch); - uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch); - for(int i =0 ; i < bytesToCopy ; i++) { - dstPtr[idx+i]= srcPtr[idx+i]; - } + if((idx < width) && (idy < height)) { + if((idx < floorWidth)){ + T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); + T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); + dstPtr[idx] = srcPtr[idx]; + } else { + size_t bytesToCopy = width - (floorWidth * sizeof(T)); + uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch); + uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch); + for(int i =0 ; i < bytesToCopy ; i++) { + dstPtr[idx+i]= srcPtr[idx+i]; + } + } } } } // namespace @@ -1524,7 +1526,7 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { template 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( ((width/sizeof(T))+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); + uint32_t grid_dim_x = clamp_integer( (width+(threadsPerBlock*sizeof(T)-1)) / (threadsPerBlock*sizeof(T)), 1, UINT32_MAX); uint32_t grid_dim_y = clamp_integer( (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); From aed26538575d4a501134a17ba59002be58a5b442 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 24 May 2018 23:30:27 +0530 Subject: [PATCH 4/5] Clean up and fix remaining bytes copy --- hipamd/src/hip_memory.cpp | 22 +++++++++------------- 1 file changed, 9 insertions(+), 13 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 2f0c793033..fa68d7cfb1 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -1496,19 +1496,15 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, size_t idx = blockIdx.x * blockDim.x + threadIdx.x; size_t idy = blockIdx.y * blockDim.y + threadIdx.y; size_t floorWidth = (width/sizeof(T)); - if((idx < width) && (idy < height)) { - if((idx < floorWidth)){ - T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); - T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); - dstPtr[idx] = srcPtr[idx]; - } else { - size_t bytesToCopy = width - (floorWidth * sizeof(T)); - uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch); - uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch); - for(int i =0 ; i < bytesToCopy ; i++) { - dstPtr[idx+i]= srcPtr[idx+i]; - } - } + T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); + T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); + if((idx < floorWidth) && (idy < height)){ + dstPtr[idx] = srcPtr[idx]; + } else if((idx < width) && (idy < height)){ + size_t bytesToCopy = width - (floorWidth * sizeof(T)); + dstPtr += floorWidth; + srcPtr += floorWidth; + __builtin_memcpy(reinterpret_cast(dstPtr), reinterpret_cast(srcPtr),bytesToCopy); } } } // namespace From 27e0566c3ac9864ead727d8311dcbc5bd585cbdd Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 24 May 2018 23:51:52 +0530 Subject: [PATCH 5/5] Use 64x4 grid dims --- hipamd/src/hip_memory.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index fa68d7cfb1..9b6758ddf4 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -1521,10 +1521,11 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { template 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( (width+(threadsPerBlock*sizeof(T)-1)) / (threadsPerBlock*sizeof(T)), 1, UINT32_MAX); - uint32_t grid_dim_y = clamp_integer( (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, + size_t threadsPerBlock_x = 64; + size_t threadsPerBlock_y = 4; + uint32_t grid_dim_x = clamp_integer( (width+(threadsPerBlock_x*sizeof(T)-1)) / (threadsPerBlock_x*sizeof(T)), 1, UINT32_MAX); + uint32_t grid_dim_y = clamp_integer( (height+(threadsPerBlock_y-1)) / threadsPerBlock_y, 1, UINT32_MAX); + hipLaunchKernelGGL(hip_copy2d_n, dim3(grid_dim_x,grid_dim_y), dim3(threadsPerBlock_x,threadsPerBlock_y), 0u, stream, dst, src, width, height, destPitch, srcPitch); }