SWDEV-301667 - Kernelarg gpuvm

Add aligned, nontemporal `memcpy` for kernarg.

Change-Id: I5d8ac76904feaf793b45ec2ea5fbd1069be20068

rocclr/device/rocm/rocvirtual.cpp

@@ -43,6 +43,7 @@
 #include <string>
 #include <thread>
 #include <vector>
+#include <immintrin.h>
 
 
 /**
@@ -2783,6 +2784,44 @@ bool VirtualGPU::createVirtualQueue(uint deviceQueueSize)
   return true;
 }
 
+// ================================================================================================
+__attribute__((optimize("unroll-all-loops"), always_inline))
+static void nontemporalMemcpy(void* __restrict dst, const void* __restrict src,
+                              uint16_t size) {
+  #if defined(__AVX512F__)
+    for (auto i = 0u; i != size / sizeof(__m512i); ++i) {
+      _mm512_stream_si512(reinterpret_cast<__m512i* __restrict&>(dst)++,
+                          *reinterpret_cast<const __m512i* __restrict&>(src)++);
+    }
+    size = size % sizeof(__m512i);
+  #endif
+
+  #if defined(__AVX__)
+    for (auto i = 0u; i != size / sizeof(__m256i); ++i) {
+      _mm256_stream_si256(reinterpret_cast<__m256i* __restrict&>(dst)++,
+                          *reinterpret_cast<const __m256i* __restrict&>(src)++);
+    }
+    size = size % sizeof(__m256i);
+  #endif
+
+  for (auto i = 0u; i != size / sizeof(__m128i); ++i) {
+    _mm_stream_si128(reinterpret_cast<__m128i* __restrict&>(dst)++,
+                     *(reinterpret_cast<const __m128i* __restrict&>(src)++));
+  }
+  size = size % sizeof(__m128i);
+
+  for (auto i = 0u; i != size / sizeof(long long); ++i) {
+    _mm_stream_si64(reinterpret_cast<long long* __restrict&>(dst)++,
+                    *reinterpret_cast<const long long* __restrict&>(src)++);
+  }
+  size = size % sizeof(long long);
+
+  for (auto i = 0u; i != size / sizeof(int); ++i) {
+    _mm_stream_si32(reinterpret_cast<int* __restrict&>(dst)++,
+                    *reinterpret_cast<const int* __restrict&>(src)++);
+  }
+}
+
 // ================================================================================================
 bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes,
     const amd::Kernel& kernel, const_address parameters, void* eventHandle,
@@ -3051,8 +3090,9 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes,
       argBuffer = reinterpret_cast<address>(allocKernArg(gpuKernel.KernargSegmentByteSize(),
                                             gpuKernel.KernargSegmentAlignment()));
       // Load all kernel arguments
-      memcpy(argBuffer, parameters, std::min(gpuKernel.KernargSegmentByteSize(),
-                                             signature.paramsSize()));
+      nontemporalMemcpy(argBuffer, parameters,
+                        std::min(gpuKernel.KernargSegmentByteSize(),
+                                 signature.paramsSize()));
     }
 
     // Check for group memory overflow

rocclr/utils/flags.hpp

@@ -52,7 +52,7 @@ debug(size_t, CPU_MEMORY_GUARD_PAGE_SIZE, 64,                                 \
         "Size in KB of CPU memory guard page")                                \
 debug(size_t, CPU_MEMORY_ALIGNMENT_SIZE, 256,                                 \
         "Size in bytes for the default alignment for guarded memory on CPU")  \
-debug(size_t, PARAMETERS_MIN_ALIGNMENT, 16,                                   \
+debug(size_t, PARAMETERS_MIN_ALIGNMENT, 64,                                   \
         "Minimum alignment required for the abstract parameters stack")       \
 debug(size_t, MEMOBJ_BASE_ADDR_ALIGN, 4*Ki,                                   \
         "Alignment of the base address of any allocate memory object")        \