From e5b1930606796aec7cf8f8d4bd24c773324ab1ef Mon Sep 17 00:00:00 2001
From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com>
Date: Mon, 13 Jun 2022 19:27:08 +0530
Subject: [PATCH] SWDEV-332522 - array test cases for streamOpsWrite &
 streamOpsWait (#2659)

Change-Id: I746c3bae592e0251b8c9482d28bd21e1bd094256

[ROCm/hip commit: 1c7d42064b36154b59bae6a89bfdc480cdb14282]
---
 .../streamOperations/hipstream_operations.cpp | 32 +++++++++++++++++--
 1 file changed, 30 insertions(+), 2 deletions(-)

diff --git a/projects/hip/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp b/projects/hip/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp
index f47328d573..cd4d9fd425 100644
--- a/projects/hip/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp
+++ b/projects/hip/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp
@@ -58,6 +58,7 @@ constexpr unsigned int writeFlag = 0;
 constexpr float SLEEP_MS = 100;
 constexpr uint32_t DATA_INIT = 0x1234;
 constexpr uint32_t DATA_UPDATE = 0X4321;
+constexpr unsigned int ARR_SIZE = 5;
 
 struct TEST_WAIT {
   int compareOp;
@@ -203,31 +204,49 @@ void testWrite() {
 
   uint64_t* host_ptr64 = (uint64_t *) malloc(sizeof(uint64_t));
   uint32_t* host_ptr32 = (uint32_t *) malloc(sizeof(uint32_t));
+  uint64_t h_mem64[ARR_SIZE];
+  uint32_t h_mem32[ARR_SIZE];
+  uint64_t* h_mem64ptr = h_mem64;
+  uint32_t* h_mem32ptr = h_mem32;
   std::cout << " hipStreamWriteValue: testing ... \n";
 
   HIPCHECK(hipExtMallocWithFlags((void **)&signalPtr, 8, hipMallocSignalMemory));
 
   void* device_ptr64;
   void* device_ptr32;
-
+  uint64_t* d_mem64ptr;
+  uint32_t* d_mem32ptr;
   *host_ptr64 = 0x0;
   *host_ptr32 = 0x0;
   *signalPtr = 0x0;
 
   hipHostRegister(host_ptr64, sizeof(uint64_t), 0);
   hipHostRegister(host_ptr32, sizeof(uint32_t), 0);
+  hipHostRegister(h_mem64ptr, sizeof(uint64_t) * ARR_SIZE, 0);
+  hipHostRegister(h_mem32ptr, sizeof(uint32_t) * ARR_SIZE, 0);
 
   // Test writting registered host pointer
   HIPCHECK(hipStreamWriteValue64(stream, host_ptr64, value64, writeFlag));
   HIPCHECK(hipStreamWriteValue32(stream, host_ptr32, value32, writeFlag));
+  //test writting to an array
+  for (int indx = 0; indx < ARR_SIZE; indx++) {
+    HIPCHECK(hipStreamWriteValue64(stream, h_mem64ptr + indx, ARR_SIZE - indx, writeFlag));
+    HIPCHECK(hipStreamWriteValue32(stream, h_mem32ptr + indx, ARR_SIZE - indx, writeFlag));
+  }
+
   hipStreamSynchronize(stream);
 
   HIPASSERT(*host_ptr64 == value64);
   HIPASSERT(*host_ptr32 == value32);
-
+  for (int indx = 0; indx < ARR_SIZE ; indx++) {
+    HIPASSERT(*(h_mem64ptr + indx) == (ARR_SIZE - indx));
+    HIPASSERT(*(h_mem32ptr + indx) == (ARR_SIZE - indx));
+  }
   // Test writting device pointer
   hipHostGetDevicePointer((void**)&device_ptr64, host_ptr64, 0);
   hipHostGetDevicePointer((void**)&device_ptr32, host_ptr32, 0);
+  hipHostGetDevicePointer((void**)&d_mem64ptr, h_mem64ptr, 0);
+  hipHostGetDevicePointer((void**)&d_mem32ptr, h_mem32ptr, 0);
 
   // Reset values
   *host_ptr64 = 0x0;
@@ -235,10 +254,19 @@ void testWrite() {
 
   HIPCHECK(hipStreamWriteValue64(stream, device_ptr64, value64, writeFlag));
   HIPCHECK(hipStreamWriteValue32(stream, device_ptr32, value32, writeFlag));
+  for (int indx = 0; indx < ARR_SIZE; indx++) {
+    HIPCHECK(hipStreamWriteValue64(stream, d_mem64ptr + indx, indx, writeFlag));
+    HIPCHECK(hipStreamWriteValue32(stream, d_mem32ptr + indx, indx, writeFlag));
+  }
+
   hipStreamSynchronize(stream);
 
   HIPASSERT(*host_ptr64 == value64);
   HIPASSERT(*host_ptr32 == value32);
+  for (int indx = 0; indx < ARR_SIZE ; indx++) {
+    HIPASSERT(*(h_mem64ptr + indx) == indx);
+    HIPASSERT(*(h_mem32ptr + indx) == indx);
+  }
 
   // Test Writing to Signal Memory
   HIPCHECK(hipStreamWriteValue64(stream, signalPtr, value64, writeFlag));