P4 to Git Change 1083967 by gandryey@gera-dev-w7 on 2014/10/03 11:20:24

ECR #304775 - Fix for BUG#10330.
	- Add an optimized version for unaligned buffer copy

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/blitcl.cpp#7 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpublit.cpp#111 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/hsa/hsablit.cpp#9 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/hsa_foundation/hsablit.cpp#5 edit

rocclr/runtime/device/blitcl.cpp

@@ -316,11 +316,12 @@ __kernel void copyBufferRectAligned(
 }
 \n
 __kernel void copyBuffer(
-    __global   uchar*  src,
-    __global   uchar*  dst,
+    __global   uchar*  srcI,
+    __global   uchar*  dstI,
     ulong   srcOrigin,
     ulong   dstOrigin,
-    ulong   size)
+    ulong   size,
+    uint    remain)
 {
     ulong   id = get_global_id(0);
 
@@ -328,10 +329,24 @@ __kernel void copyBuffer(
         return;
     }
 
-    ulong   offsSrc = id + srcOrigin;
-    ulong   offsDst = id + dstOrigin;
+    __global uchar* src = srcI + srcOrigin;
+    __global uchar* dst = dstI + dstOrigin;
 
-    dst[offsDst] = src[offsSrc];
+    if (remain == 8) {
+        dst[id] = src[id];
+    }
+    else {
+        if (id < (size - 1)) {
+            __global uint* srcD = (__global uint*)(src);
+            __global uint* dstD = (__global uint*)(dst);
+            dstD[id] = srcD[id];
+        }
+        else {
+            for (uint i = 0; i < remain; ++i) {
+                dst[id * 4 + i] = src[id * 4 + i];
+            }
+        }
+    }
 }
 \n
 __kernel void copyBufferAligned(

rocclr/runtime/device/gpu/gpublit.cpp

@@ -2513,7 +2513,30 @@ KernelBlitManager::copyBuffer(
             }
         }
 
-        size.c[0] /= CopyBuffAlignment[i];
+        cl_uint remain;
+        if (blitType == BlitCopyBufferAligned) {
+            size.c[0] /= CopyBuffAlignment[i];
+        }
+        else {
+            if (dev().settings().ciPlus_) {
+                remain = size[0] % 4;
+                size.c[0] /= 4;
+                size.c[0] += 1;
+            }
+            else {
+                // Check if offsets are aligned
+                aligned = ((srcOrigin[0] % sizeof(uint32_t)) == 0);
+                aligned &= ((dstOrigin[0] % sizeof(uint32_t)) == 0);
+                if (aligned) {
+                    remain = size[0] % 4;
+                    size.c[0] /= 4;
+                    size.c[0] += 1;
+                }
+                else {
+                    remain = 8;
+                }
+            }
+        }
 
         // Program the dispatch dimensions
         localWorkSize = 256;
@@ -2539,6 +2562,9 @@ KernelBlitManager::copyBuffer(
             cl_int  alignment = CopyBuffAlignment[i];
             setArgument(kernels_[blitType], 5, sizeof(alignment), &alignment);
         }
+        else {
+            setArgument(kernels_[blitType], 5, sizeof(remain), &remain);
+        }
 
         // Create ND range object for the kernel's execution
         amd::NDRangeContainer ndrange(1,

rocclr/runtime/device/hsa/hsablit.cpp

@@ -1055,7 +1055,15 @@ KernelBlitManager::copyBuffer(
         }
     }
 
-    size.c[0] /= CopyBuffAlignment[i];
+    cl_uint remain;
+    if (blitType == BlitCopyBufferAligned) {
+        size.c[0] /= CopyBuffAlignment[i];
+    }
+    else {
+        remain = size[0] % 4;
+        size.c[0] /= 4;
+        size.c[0] += 1;
+    }
 
     // Program the dispatch dimensions
     localWorkSize = 256;
@@ -1081,6 +1089,9 @@ KernelBlitManager::copyBuffer(
         cl_int  alignment = CopyBuffAlignment[i];
         kernels_[blitType]->parameters().set(5, sizeof(alignment), &alignment);
     }
+    else {
+        kernels_[blitType]->parameters().set(5, sizeof(remain), &remain);
+    }
 
     // Create ND range object for the kernel's execution
     amd::NDRangeContainer ndrange(