merge amd-stg-open into amd-mainline-open

Change-Id: Ib9491ebc1df677a45c1edb6c870eb76f2ce8f350
ROCm · Jan 18, 2022 · 62809a5 · 62809a5
2 parents d8171e7 + ff17b35
commit 62809a5
Show file tree

Hide file tree

Showing 3 changed files with 54 additions and 0 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -84,8 +84,13 @@ if(CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
   set ( CPACK_PACKAGE_CONTACT "Advanced Micro Devices Inc." )
   set ( CPACK_PACKAGE_DESCRIPTION_SUMMARY "Radeon Open Compute - device libraries" )
   set ( CPACK_PACKAGE_DESCRIPTION "This package includes LLVM bitcode libraries." )
+  set ( CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE.TXT" )
   set ( CPACK_DEBIAN_PACKAGE_HOMEPAGE "https://github.com/RadeonOpenCompute/ROCm-Device-Libs" )
 
+  # Install License file
+  install ( FILES "${CPACK_RESOURCE_FILE_LICENSE}"
+          DESTINATION ${CMAKE_INSTALL_DOCDIR}/${CPACK_PACKAGE_NAME})
+
   set ( CPACK_GENERATOR "DEB;RPM" CACHE STRING "Default packaging generators." )
 
   ## ROCM version updates as per naming convention
@@ -107,6 +112,7 @@ if(CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
   endif()
   ## get distro for RPM package using dist
   message("device-libs CPACK_RPM_PACKAGE_RELEASE now is ${CPACK_RPM_PACKAGE_RELEASE}")
+  set( CPACK_RPM_PACKAGE_LICENSE "NCSA" )
 
   execute_process( COMMAND rpm --eval %{?dist}
                  RESULT_VARIABLE _result_var

diff --git a/LICENSE → LICENSE.TXT b/LICENSE → LICENSE.TXT
diff --git a/opencl/src/misc/amdblit.cl b/opencl/src/misc/amdblit.cl
@@ -471,6 +471,54 @@ __amd_fillBufferAligned(
     }
 }
 
+__attribute__((always_inline)) void
+__amd_fillBufferAligned2D(__global uchar* bufUChar,
+                          __global ushort* bufUShort,
+                          __global uint* bufUInt,
+                          __global ulong* bufULong,
+                          __constant uchar* pattern,
+                          uint patternSize,
+                          ulong origin,
+                          ulong width,
+                          ulong height,
+                          ulong pitch)
+{
+  ulong tid_x = get_global_id(0);
+  ulong tid_y = get_global_id(1);
+
+  if (tid_x >= width || tid_y >= height) {
+    return;
+  }
+
+  ulong offset = (tid_y * pitch + tid_x);
+
+  if (bufULong) {
+    __global ulong* element = &bufULong[origin + offset];
+    __constant ulong* pt = (__constant ulong*)pattern;
+    for (uint i = 0; i < patternSize; ++i) {
+      element[i] = pt[i];
+    }
+  } else if (bufUInt) {
+    __global uint* element = &bufUInt[origin + offset];
+    __constant uint* pt = (__constant uint*)pattern;
+    for (uint i = 0; i < patternSize; ++i) {
+      element[i] = pt[i];
+    }
+  } else if (bufUShort) {
+    __global ushort* element = &bufUShort[origin + offset];
+    __constant ushort* pt = (__constant ushort*)pattern;
+    for (uint i = 0; i < patternSize; ++i) {
+      element[i] = pt[i];
+    }
+  } else if (bufUChar) {
+    __global uchar* element = &bufUChar[origin + offset];
+    __constant uchar* pt = (__constant uchar*)pattern;
+    for (uint i = 0; i < patternSize; ++i) {
+      element[i] = pt[i];
+    }
+  }
+}
+
 __attribute__((always_inline)) void
 __amd_fillImage(
     __write_only image2d_array_t image,