AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Optimize memcpy2D kernel use

Browse Source 
[ROCm/clr commit: 2c3d1498d4]
This commit is contained in:
Rahul Garg
2018-05-23 14:43:47 +05:30
parent 335dad2feb
commit 432467d88b
1 changed files with 12 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_memory.cpp
+12 -4
View File
@@ -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 <typename T>
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_x = clamp_integer<size_t>( ((width/sizeof(T))+(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);
@@ -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<uint8_t> (stream, static_cast<uint8_t*> (dst), static_cast<const uint8_t*> (src), width, height, dpitch, spitch);
ihipMemcpy2dKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), static_cast<const uint32_t*> (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<uint8_t> (stream, static_cast<uint8_t*> (dst), static_cast<const uint8_t*> (src), width, height, dpitch, spitch);
ihipMemcpy2dKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), static_cast<const uint32_t*> (src), width, height, dpitch, spitch);
}
} catch (ihipException& ex) {
e = ex._code;