From 21cf5b0ae46d85183780956c2c72bd85dd2c2909 Mon Sep 17 00:00:00 2001
From: "Wen-Heng (Jack) Chung" <whchung@gmail.com>
Date: Mon, 4 Mar 2019 17:11:54 +0000
Subject: [PATCH] Add direct test for hipMemsetD32 and hipMemsetD32Async

[ROCm/clr commit: 365d08535ba9a9f5ce1a7152ed48c30409784bfe]
---
 .../tests/src/runtimeApi/memory/hipMemset.cpp | 61 ++++++++++++++++++-
 projects/clr/hipamd/tests/src/test_common.cpp |  7 +++
 projects/clr/hipamd/tests/src/test_common.h   |  1 +
 3 files changed, 66 insertions(+), 3 deletions(-)

diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp
index bcb0adbeef..8ac5beaf55 100644
--- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp
+++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp
@@ -26,11 +26,11 @@ THE SOFTWARE.
  * BUILD: %t %s ../../test_common.cpp
  * RUN: %t
  * //Small copy
- * RUN: %t -N 10    --memsetval 0x42
+ * RUN: %t -N 10    --memsetval 0x42 --memsetD32val 0x101
  * // Oddball size
- * RUN: %t -N 10013 --memsetval 0x5a
+ * RUN: %t -N 10013 --memsetval 0x5a --memsetD32val 0xDEADBEEF
  * // Big copy
- * RUN: %t -N 256M  --memsetval 0xa6
+ * RUN: %t -N 256M  --memsetval 0xa6 --memsetD32val 0xCAFEBABE
  * HIT_END
  */
 
@@ -62,6 +62,30 @@ bool testhipMemset(int memsetval,int p_gpuDevice)
     return testResult;
 }
 
+bool testhipMemsetD32(int memsetD32val,int p_gpuDevice)
+{
+    size_t Nbytes = N*sizeof(int);
+    printf ("testhipMemsetD32 N=%zu  memsetD32val=%8x device=%d\n", N, memsetD32val, p_gpuDevice);
+    int *A_d;
+    int *A_h;
+    bool testResult = true;
+
+    HIPCHECK ( hipMalloc(&A_d, Nbytes) );
+    A_h = (int*)malloc(Nbytes);
+    HIPCHECK ( hipMemsetD32(A_d, memsetD32val, N) );
+    HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
+
+    for (int i=0; i<N; i++) {
+        if (A_h[i] != memsetD32val) {
+            testResult = false;            printf("mismatch at index:%d computed:%08x, memsetD32val:%08x\n", i, A_h[i], memsetD32val);
+            break;
+        }
+    }
+    HIPCHECK(hipFree(A_d));
+    free(A_h);
+    return testResult;
+}
+
 bool testhipMemsetAsync(int memsetval,int p_gpuDevice)
 {
     size_t Nbytes = N*sizeof(int);
@@ -91,6 +115,35 @@ bool testhipMemsetAsync(int memsetval,int p_gpuDevice)
     return testResult;
 }
 
+bool testhipMemsetD32Async(int memsetD32val,int p_gpuDevice)
+{
+    size_t Nbytes = N*sizeof(int);
+    printf ("testhipMemsetD32Async N=%zu  memsetval=%8x device=%d\n", N, memsetD32val, p_gpuDevice);
+    int *A_d;
+    int *A_h;
+    bool testResult = true;
+
+    HIPCHECK ( hipMalloc((void**)&A_d, Nbytes) );
+    A_h = (int*)malloc(Nbytes);
+    hipStream_t stream;
+    HIPCHECK(hipStreamCreate(&stream));
+    HIPCHECK ( hipMemsetD32Async(A_d, memsetD32val, N, stream ));
+    HIPCHECK ( hipStreamSynchronize(stream));
+    HIPCHECK ( hipMemcpy(A_h, (void*)A_d, Nbytes, hipMemcpyDeviceToHost));
+
+    for (int i=0; i<N; i++) {
+        if (A_h[i] != memsetD32val) {
+            testResult = false;
+            printf("mismatch at index:%d computed:%02x, memsetD32val:%02x\n", i, A_h[i], memsetD32val);
+            break;
+        }
+    }
+    HIPCHECK(hipFree((void*)A_d));
+    HIPCHECK(hipStreamDestroy(stream));
+    free(A_h);
+    return testResult;
+}
+
 int main(int argc, char *argv[])
 {
     HipTest::parseStandardArguments(argc, argv, true);
@@ -98,6 +151,8 @@ int main(int argc, char *argv[])
     HIPCHECK(hipSetDevice(p_gpuDevice));
     testResult &= testhipMemset(memsetval, p_gpuDevice);
     testResult &= testhipMemsetAsync(memsetval, p_gpuDevice);
+    testResult &= testhipMemsetD32(memsetD32val, p_gpuDevice);
+    testResult &= testhipMemsetD32Async(memsetD32val, p_gpuDevice);
     if (testResult) passed();
     failed("Output Mismatch\n"); 
 }
diff --git a/projects/clr/hipamd/tests/src/test_common.cpp b/projects/clr/hipamd/tests/src/test_common.cpp
index 8acc03fabf..4b39446b35 100644
--- a/projects/clr/hipamd/tests/src/test_common.cpp
+++ b/projects/clr/hipamd/tests/src/test_common.cpp
@@ -24,6 +24,7 @@ THE SOFTWARE.
 // standard global variables that can be set on command line
 size_t N = 4 * 1024 * 1024;
 char memsetval = 0x42;
+int memsetD32val = 0xDEADBEEF;
 int iterations = 1;
 unsigned blocksPerCU = 6;  // to hide latency
 unsigned threadsPerBlock = 256;
@@ -99,6 +100,12 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg) {
                 failed("Bad memsetval argument");
             }
             memsetval = ex;
+        } else if (!strcmp(arg, "--memsetD32val")) {
+            int ex;
+            if (++i >= argc || !HipTest::parseInt(argv[i], &ex)) {
+                failed("Bad memsetD32val argument");
+            }
+            memsetD32val = ex;
         } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) {
             if (++i >= argc || !HipTest::parseInt(argv[i], &iterations)) {
                 failed("Bad iterations argument");
diff --git a/projects/clr/hipamd/tests/src/test_common.h b/projects/clr/hipamd/tests/src/test_common.h
index bacbf35a22..8820381826 100644
--- a/projects/clr/hipamd/tests/src/test_common.h
+++ b/projects/clr/hipamd/tests/src/test_common.h
@@ -98,6 +98,7 @@ THE SOFTWARE.
 // standard command-line variables:
 extern size_t N;
 extern char memsetval;
+extern int memsetD32val;
 extern int iterations;
 extern unsigned blocksPerCU;
 extern unsigned threadsPerBlock;