SWDEV-434298 - Add the new copy kernel

The new copy kernel can work with a limited number
of launched workgroups

Change-Id: I7807b37e4feb4ae21f10542837a6b17d563f03b5

Author: gandryey

amd/device-libs/opencl/src/misc/amdblit.cl

@@ -389,6 +389,45 @@ __amd_copyBufferAligned(
     }
 }
 
+__attribute__((always_inline)) void
+__amd_copyBufferExt(
+    __global uchar* srcI,
+    __global uchar* dstI,
+    ulong srcOrigin,
+    ulong dstOrigin,
+    ulong size,
+    uint remainder,
+    uint aligned_size,
+    ulong end_ptr,
+    uint next_chunk) {
+  ulong id = get_global_id(0);
+  ulong id_remainder = id;
+
+  __global uchar* src = srcI + srcOrigin;
+  __global uchar* dst = dstI + dstOrigin;
+
+  if (aligned_size == sizeof(ulong2)) {
+    __global ulong2* srcD = (__global ulong2*)(src);
+    __global ulong2* dstD = (__global ulong2*)(dst);
+    while ((ulong)(&dstD[id]) < end_ptr) {
+      dstD[id] = srcD[id];
+      id += next_chunk;
+    }
+  } else {
+    __global uint* srcD = (__global uint*)(src);
+    __global uint* dstD = (__global uint*)(dst);
+    while ((ulong)(&dstD[id]) < end_ptr) {
+      dstD[id] = srcD[id];
+      id += next_chunk;
+    }
+  }
+  if ((remainder != 0) && (id_remainder == 0)) {
+    for (ulong i = size - remainder; i < size; ++i) {
+      dst[i] = src[i];
+    }
+  }
+}
+
 __attribute__((always_inline)) void
 __amd_fillBuffer(
     __global uchar* bufUChar,