SWDEV-294470 - [dtest] Catch2 unit tests for memset related tests. (#2345)

Change-Id: Ib227e75cb0bef9273bc787e47fa5b713086fac46

catch/unit/memory/CMakeLists.txt

@@ -33,6 +33,10 @@ set(TEST_SRC
     hipMemoryAllocateCoherent.cc
     hipMallocManaged_MultiScenario.cc
     hipManagedKeyword.cc
+    hipMemsetInvalidPtr.cc
+    hipMemset.cc
+    hipMemsetAsyncMultiThread.cc
+    hipMemsetAsyncAndKernel.cc
 )
 else()
 set(TEST_SRC
@@ -66,6 +70,10 @@ set(TEST_SRC
     hipMemoryAllocateCoherent.cc
     hipMallocManaged_MultiScenario.cc
     hipManagedKeyword.cc
+    hipMemsetInvalidPtr.cc
+    hipMemset.cc
+    hipMemsetAsyncMultiThread.cc
+    hipMemsetAsyncAndKernel.cc
 )
 endif()
 # Create shared lib of all tests

catch/unit/memory/hipMemset.cc

@@ -0,0 +1,281 @@
+/*
+ * Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+
+/**
+Testcase Scenarios :
+ 1) Test hipMemset small size buffers with unique memset values.
+ 2) Test hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32 apis with unique
+    number of elements and memset values.
+ 3) Test hipMemsetAsync, hipMemsetD8Async, hipMemsetD16Async, hipMemsetD32Async
+    apis with unique number of elements and memset values.
+ 4) Test two memset async operations at the same time.
+*/
+
+
+#include <hip_test_common.hh>
+
+
+// Table with unique number of elements and memset values.
+// (N, memsetval, memsetD32val, memsetD16val, memsetD8val)
+typedef std::tuple<size_t, char, int, int16_t, char> tupletype;
+static constexpr std::initializer_list<tupletype> tableItems {
+               std::make_tuple((4*1024*1024), 0x42, 0xDEADBEEF, 0xDEAD, 0xDE),
+               std::make_tuple((10)         , 0x42, 0x101     , 0x10,   0x1),
+               std::make_tuple((10013)      , 0x5a, 0xDEADBEEF, 0xDEAD, 0xDE),
+               std::make_tuple((256*1024*1024), 0xa6, 0xCAFEBABE, 0xCAFE, 0xCA)
+               };
+
+enum MemsetType {
+  hipMemsetTypeDefault,
+  hipMemsetTypeD8,
+  hipMemsetTypeD16,
+  hipMemsetTypeD32
+};
+
+template<typename T>
+static bool testhipMemset(T *A_h, T *A_d, T memsetval, enum MemsetType type,
+                  size_t numElements) {
+  size_t Nbytes = numElements * sizeof(T);
+  bool testResult = true;
+  constexpr auto MAX_OFFSET = 3;  // To memset on unaligned ptr.
+
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  A_h = reinterpret_cast<T*> (malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+
+  for (int offset = MAX_OFFSET; offset >= 0; offset --) {
+    if (type == hipMemsetTypeDefault) {
+      HIP_CHECK(hipMemset(A_d + offset, memsetval, numElements - offset));
+
+    } else if (type == hipMemsetTypeD8) {
+      HIP_CHECK(hipMemsetD8((hipDeviceptr_t)(A_d + offset), memsetval,
+                                                    numElements - offset));
+
+    } else if (type == hipMemsetTypeD16) {
+      HIP_CHECK(hipMemsetD16((hipDeviceptr_t)(A_d + offset), memsetval,
+                                                    numElements - offset));
+
+    } else if (type == hipMemsetTypeD32) {
+      HIP_CHECK(hipMemsetD32((hipDeviceptr_t)(A_d + offset), memsetval,
+                                                    numElements - offset));
+    }
+
+    HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
+    for (size_t i = offset; i < numElements; i++) {
+      if (A_h[i] != memsetval) {
+        testResult = false;
+        CAPTURE(i, A_h[i], memsetval);
+        break;
+      }
+    }
+  }
+
+  HIP_CHECK(hipFree(A_d));
+  free(A_h);
+  return testResult;
+}
+
+
+template<typename T>
+static bool testhipMemsetAsync(T *A_h, T *A_d, T memsetval,
+                                 enum MemsetType type, size_t numElements) {
+  size_t Nbytes = numElements * sizeof(T);
+  bool testResult = true;
+  constexpr auto MAX_OFFSET = 3;  // To memset on unaligned ptr.
+  hipStream_t stream;
+
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  A_h = reinterpret_cast<T*> (malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+
+  for (int offset = MAX_OFFSET; offset >= 0; offset --) {
+    if (type == hipMemsetTypeDefault) {
+      HIP_CHECK(hipMemsetAsync(A_d + offset, memsetval, numElements - offset,
+                                                                      stream));
+
+    } else if (type == hipMemsetTypeD8) {
+      HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)(A_d + offset), memsetval,
+                                                numElements - offset, stream));
+
+    } else if (type == hipMemsetTypeD16) {
+      HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)(A_d + offset), memsetval,
+                                                numElements - offset, stream));
+
+    } else if (type == hipMemsetTypeD32) {
+      HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)(A_d + offset), memsetval,
+                                                numElements - offset, stream));
+    }
+
+    HIP_CHECK(hipStreamSynchronize(stream));
+    HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
+    for (size_t i = offset; i < numElements; i++) {
+      if (A_h[i] != memsetval) {
+        testResult = false;
+        CAPTURE(i, A_h[i], memsetval);
+        break;
+      }
+    }
+  }
+
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipStreamDestroy(stream));
+  free(A_h);
+  return testResult;
+}
+
+
+/**
+ * Test hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32 apis with unique
+ * number of elements and memset values.
+ */
+TEST_CASE("Unit_hipMemset_SetMemoryWithOffset") {
+  char memsetval;
+  int memsetD32val;
+  int16_t memsetD16val;
+  char memsetD8val;
+  size_t N;
+  bool ret;
+
+  std::tie(N, memsetval, memsetD32val, memsetD16val, memsetD8val) =
+                 GENERATE(table<size_t, char, int, int16_t, char>(tableItems));
+
+
+  SECTION("Memset with hipMemsetTypeDefault") {
+    char *cA_d{nullptr}, *cA_h{nullptr};
+    ret = testhipMemset(cA_h, cA_d, memsetval, hipMemsetTypeDefault, N);
+    REQUIRE(ret == true);
+  }
+
+  SECTION("Memset with hipMemsetTypeD32") {
+    int32_t *iA_d{nullptr}, *iA_h{nullptr};
+    ret = testhipMemset(iA_h, iA_d, memsetD32val, hipMemsetTypeD32, N);
+    REQUIRE(ret == true);
+  }
+
+  SECTION("Memset with hipMemsetTypeD16") {
+    int16_t *siA_d{nullptr}, *siA_h{nullptr};
+    ret = testhipMemset(siA_h, siA_d, memsetD16val, hipMemsetTypeD16, N);
+    REQUIRE(ret == true);
+  }
+
+  SECTION("Memset with hipMemsetTypeD8") {
+    char *cA_d{nullptr}, *cA_h{nullptr};
+    ret = testhipMemset(cA_h, cA_d, memsetD8val, hipMemsetTypeD8, N);
+    REQUIRE(ret == true);
+  }
+}
+
+
+/**
+ * Test hipMemsetAsync, hipMemsetD8Async, hipMemsetD16Async, hipMemsetD32Async
+ * apis with unique number of elements and memset values.
+ */
+TEST_CASE("Unit_hipMemsetAsync_SetMemoryWithOffset") {
+  char memsetval;
+  int memsetD32val;
+  int16_t memsetD16val;
+  char memsetD8val;
+  size_t N;
+  bool ret;
+
+  std::tie(N, memsetval, memsetD32val, memsetD16val, memsetD8val) =
+                 GENERATE(table<size_t, char, int, int16_t, char>(tableItems));
+
+
+  SECTION("Memset with hipMemsetTypeDefault") {
+    char *cA_d{nullptr}, *cA_h{nullptr};
+    ret = testhipMemsetAsync(cA_h, cA_d, memsetval, hipMemsetTypeDefault, N);
+    REQUIRE(ret == true);
+  }
+
+  SECTION("Memset with hipMemsetTypeD32") {
+    int32_t *iA_d{nullptr}, *iA_h{nullptr};
+    ret = testhipMemsetAsync(iA_h, iA_d, memsetD32val, hipMemsetTypeD32, N);
+    REQUIRE(ret == true);
+  }
+
+  SECTION("Memset with hipMemsetTypeD16") {
+    int16_t *siA_d{nullptr}, *siA_h{nullptr};
+    ret = testhipMemsetAsync(siA_h, siA_d, memsetD16val, hipMemsetTypeD16, N);
+    REQUIRE(ret == true);
+  }
+
+  SECTION("Memset with hipMemsetTypeD8") {
+    char *cA_d{nullptr}, *cA_h{nullptr};
+    ret = testhipMemsetAsync(cA_h, cA_d, memsetD8val, hipMemsetTypeD8, N);
+    REQUIRE(ret == true);
+  }
+}
+
+/**
+ * Test hipMemset small size buffers with unique memset values.
+ */
+TEST_CASE("Unit_hipMemset_SmallBufferSizes") {
+  char *A_d, *A_h;
+  constexpr int memsetval = 0x24;
+
+  auto numElements = GENERATE(range(1, 4));
+  int numBytes = numElements * sizeof(char);
+
+  HIP_CHECK(hipMalloc(&A_d, numBytes));
+  A_h = reinterpret_cast<char*> (malloc(numBytes));
+
+  HIP_CHECK(hipMemset(A_d, memsetval, numBytes));
+  HIP_CHECK(hipMemcpy(A_h, A_d, numBytes, hipMemcpyDeviceToHost));
+
+  for (int i = 0; i < numBytes; i++) {
+    if (A_h[i] != memsetval) {
+      INFO("Mismatch at index:" << i << " computed:" << A_h[i]
+                                          << " memsetval:" << memsetval);
+      REQUIRE(false);
+    }
+  }
+
+  HIP_CHECK(hipFree(A_d));
+  free(A_h);
+}
+
+
+/**
+ * Test two memset async operations at the same time.
+ */
+TEST_CASE("Unit_hipMemset_2AsyncOperations") {
+  std::vector<float> v;
+  v.resize(2048);
+  float* p2, *p3;
+  hipMalloc(reinterpret_cast<void**>(&p2), 4096 + 4096*2);
+  p3 = p2+2048;
+  hipStream_t s;
+  hipStreamCreate(&s);
+  hipMemsetAsync(p2, 0, 32*32*4, s);
+  hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s);
+  hipStreamSynchronize(s);
+  for (int i = 0; i < 256; ++i) {
+    hipMemsetAsync(p2, 0, 32*32*4, s);
+    hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s);
+  }
+  hipStreamSynchronize(s);
+  hipDeviceSynchronize();
+  hipMemcpy(&v[0], p2, 1024, hipMemcpyDeviceToHost);
+  hipMemcpy(&v[1024], p3, 1024, hipMemcpyDeviceToHost);
+
+  REQUIRE(v[0] == 0);
+  REQUIRE(v[1024] == 1.75f);
+}

catch/unit/memory/hipMemsetAsyncAndKernel.cc

@@ -0,0 +1,193 @@
+/*
+ * Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+*/
+
+/*
+ * Test for checking order of execution of device kernel and
+ * hipMemsetAsync apis on all gpus
+ */
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+#define ITER 6
+#define N 1024 * 1024
+
+constexpr auto blocksPerCU = 6;  // to hide latency
+constexpr auto threadsPerBlock = 256;
+static unsigned blocks = 0;
+
+
+template <typename T>
+class MemSetKernelTest {
+ public:
+  T *A_h, *B_d, *B_h, *C_d;
+  T memSetVal;
+  size_t Nbytes;
+  bool testResult = true;
+  int validateCount = 0;
+  hipStream_t stream;
+
+  void memAllocate(T memSetValue) {
+    memSetVal = memSetValue;
+    Nbytes = N * sizeof(T);
+
+    A_h = reinterpret_cast<T*>(malloc(Nbytes));
+    HIP_ASSERT(A_h != nullptr);
+    HIP_CHECK(hipMalloc(&B_d , Nbytes));
+    B_h = reinterpret_cast<T*>(malloc(Nbytes));
+    HIP_ASSERT(B_h != nullptr);
+    HIP_CHECK(hipMalloc(&C_d , Nbytes));
+
+    for (int i = 0 ; i < N ; i++) {
+      B_h[i] = i;
+    }
+    HIP_CHECK(hipMemcpy(B_d , B_h , Nbytes , hipMemcpyHostToDevice));
+    HIP_CHECK(hipStreamCreate(&stream));
+  }
+
+  void memDeallocate() {
+    HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d));
+    free(B_h); free(A_h);
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+
+  void validateExecutionOrder() {
+    for (int p = 0 ; p < N ; p++) {
+      if (A_h[p] == memSetVal) {
+        validateCount+= 1;
+      }
+    }
+  }
+
+  bool resultAfterAllIterations() {
+    testResult = (validateCount == (ITER * N)) ? true : false;
+    memDeallocate();
+    return testResult;
+  }
+};
+
+static bool testhipMemsetAsyncWithKernel() {
+  MemSetKernelTest<char> obj;
+  constexpr char memsetval = 0x42;
+
+  obj.memAllocate(memsetval);
+  for (int k = 0 ; k < ITER ; k++) {
+    hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                    dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
+    HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream));
+    HIP_CHECK(hipStreamSynchronize(obj.stream));
+    HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
+
+    obj.validateExecutionOrder();
+  }
+  return obj.resultAfterAllIterations();
+}
+
+static bool testhipMemsetD32AsyncWithKernel() {
+  MemSetKernelTest <int32_t> obj;
+  constexpr int memsetD32val = 0xDEADBEEF;
+
+  obj.memAllocate(memsetD32val);
+  for (int k = 0 ; k < ITER ; k++) {
+    hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                    dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
+    HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)obj.C_d , obj.memSetVal,
+                                                              N, obj.stream));
+    HIP_CHECK(hipStreamSynchronize(obj.stream));
+    HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
+
+    obj.validateExecutionOrder();
+  }
+  return obj.resultAfterAllIterations();
+}
+
+static bool testhipMemsetD16AsyncWithKernel() {
+  MemSetKernelTest <int16_t> obj;
+  constexpr int16_t memsetD16val = 0xDEAD;
+
+  obj.memAllocate(memsetD16val);
+  for (int k = 0 ; k < ITER ; k++) {
+    hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                    dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
+    HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)obj.C_d , obj.memSetVal,
+                                                              N, obj.stream));
+    HIP_CHECK(hipStreamSynchronize(obj.stream));
+    HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
+
+    obj.validateExecutionOrder();
+  }
+  return obj.resultAfterAllIterations();
+}
+
+static bool testhipMemsetD8AsyncWithKernel() {
+  MemSetKernelTest <char> obj;
+  constexpr char memsetD8val = 0xDE;
+
+  obj.memAllocate(memsetD8val);
+  for (int k = 0; k < ITER; k++) {
+    hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                    dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
+    HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)obj.C_d, obj.memSetVal,
+                                                              N, obj.stream));
+    HIP_CHECK(hipStreamSynchronize(obj.stream));
+    HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
+
+    obj.validateExecutionOrder();
+  }
+  return obj.resultAfterAllIterations();
+}
+
+
+/*
+ * Test for checking order of execution of device kernel and
+ * hipMemsetAsync apis on all gpus
+ */
+TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") {
+  int numDevices = 0;
+  bool ret;
+
+  blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  REQUIRE(numDevices > 0);
+
+  auto devNum = GENERATE_COPY(range(0, numDevices));
+  HIP_CHECK(hipSetDevice(devNum));
+
+  SECTION("hipMemsetAsync With Kernel") {
+    ret = testhipMemsetAsyncWithKernel();
+    REQUIRE(ret == true);
+  }
+
+  SECTION("hipMemsetD32Async With Kernel") {
+    ret = testhipMemsetD32AsyncWithKernel();
+    REQUIRE(ret == true);
+  }
+
+  SECTION("hipMemsetD16Async With Kernel") {
+    ret = testhipMemsetD16AsyncWithKernel();
+    REQUIRE(ret == true);
+  }
+
+  SECTION("hipMemsetD8Async With Kernel") {
+    ret = testhipMemsetD8AsyncWithKernel();
+    REQUIRE(ret == true);
+  }
+}

catch/unit/memory/hipMemsetAsyncMultiThread.cc

@@ -0,0 +1,243 @@
+/*
+ * Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+
+/*
+ * Test that validates functionality of hipmemsetAsync apis over multi threads
+ */
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+
+
+#define NUM_THREADS 20
+#define ITER 10
+#define N (4*1024*1024)
+
+
+template <typename T>
+class MemSetAsyncMthreadTest {
+ public:
+  T *A_h, *A_d, *B_h;
+  T memSetVal;
+  size_t Nbytes;
+  bool testResult = true;
+  int validateCount = 0;
+  hipStream_t stream;
+
+  void memAllocate(T memSetValue) {
+    memSetVal = memSetValue;
+    Nbytes = N * sizeof(T);
+
+    A_h = reinterpret_cast<T*>(malloc(Nbytes));
+    HIP_ASSERT(A_h != nullptr);
+
+    HIP_CHECK(hipMalloc(&A_d, Nbytes));
+    B_h = reinterpret_cast<T*>(malloc(Nbytes));
+    HIP_ASSERT(B_h != nullptr);
+
+    HIP_CHECK(hipStreamCreate(&stream));
+  }
+
+  void threadCompleteStatus() {
+    for (int k = 0 ; k < N ; k++) {
+      if ((A_h[k] == memSetVal) && (B_h[k] == memSetVal)) {
+        validateCount+= 1;
+      }
+    }
+  }
+
+  bool resultAfterAllIterations() {
+    memDeallocate();
+    testResult = (validateCount == (ITER * N)) ? true: false;
+    return testResult;
+  }
+
+  void memDeallocate() {
+    HIP_CHECK(hipFree(A_d));
+    free(A_h);
+    free(B_h);
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+};
+
+template <typename T>
+void queueJobsForhipMemsetAsync(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
+                                hipStream_t stream) {
+  HIPCHECK(hipMemsetAsync(A_d, memSetVal, N, stream));
+  HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
+}
+
+template <typename T>
+void queueJobsForhipMemsetD32Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
+                                   hipStream_t stream) {
+  HIPCHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memSetVal, N, stream));
+  HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
+}
+
+template <typename T>
+void queueJobsForhipMemsetD16Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
+                                   hipStream_t stream) {
+  HIPCHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memSetVal, N, stream));
+  HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
+}
+
+template <typename T>
+void queueJobsForhipMemsetD8Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
+                                  hipStream_t stream) {
+  HIPCHECK(hipMemsetD8Async((hipDeviceptr_t)A_d, memSetVal, N, stream));
+  HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
+}
+
+/* Queue hipMemsetAsync jobs on multiple threads and verify they all
+ * finished on all threads successfully
+ */
+bool testhipMemsetAsyncWithMultiThread() {
+  MemSetAsyncMthreadTest <char> obj;
+  constexpr char memsetval = 0x42;
+  obj.memAllocate(memsetval);
+  std::thread t[NUM_THREADS];
+
+  for (int i = 0 ; i < ITER ; i++) {
+    for (int k = 0 ; k < NUM_THREADS ; k++) {
+      if (k%2) {
+        t[k] = std::thread(queueJobsForhipMemsetAsync<char>, obj.A_d, obj.A_h,
+                           obj.memSetVal, obj.Nbytes, obj.stream);
+      } else {
+        t[k] = std::thread(queueJobsForhipMemsetAsync<char>, obj.A_d, obj.B_h,
+                             obj.memSetVal, obj.Nbytes, obj.stream);
+      }
+    }
+
+    for (int j = 0 ; j < NUM_THREADS ; j++) {
+      t[j].join();
+    }
+
+    HIP_CHECK(hipStreamSynchronize(obj.stream));
+    obj.threadCompleteStatus();
+  }
+  return obj.resultAfterAllIterations();
+}
+
+bool testhipMemsetD32AsyncWithMultiThread() {
+  MemSetAsyncMthreadTest <int32_t> obj;
+  constexpr int memsetD32val = 0xDEADBEEF;
+  obj.memAllocate(memsetD32val);
+  std::thread t[NUM_THREADS];
+
+  for (int i = 0 ; i < ITER ; i++) {
+    for (int k = 0 ; k < NUM_THREADS ; k++) {
+      if (k%2) {
+        t[k] = std::thread(queueJobsForhipMemsetD32Async<int32_t>, obj.A_d,
+                           obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream);
+      } else {
+        t[k] = std::thread(queueJobsForhipMemsetD32Async<int32_t>, obj.A_d,
+                           obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream);
+      }
+    }
+
+    for (int j = 0 ; j < NUM_THREADS ; j++) {
+      t[j].join();
+    }
+
+    HIP_CHECK(hipStreamSynchronize(obj.stream));
+    obj.threadCompleteStatus();
+  }
+  return obj.resultAfterAllIterations();
+}
+
+bool testhipMemsetD16AsyncWithMultiThread() {
+  MemSetAsyncMthreadTest <int16_t> obj;
+  constexpr int16_t memsetD16val = 0xDEAD;
+  obj.memAllocate(memsetD16val);
+  std::thread t[NUM_THREADS];
+
+  for (int i = 0 ; i < ITER ; i++) {
+    for (int k = 0 ; k < NUM_THREADS ; k++) {
+      if (k%2) {
+        t[k] = std::thread(queueJobsForhipMemsetD16Async<int16_t>, obj.A_d,
+                           obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream);
+      } else {
+        t[k] = std::thread(queueJobsForhipMemsetD16Async<int16_t>, obj.A_d,
+                           obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream);
+      }
+    }
+
+    for (int j = 0 ; j < NUM_THREADS ; j++) {
+      t[j].join();
+    }
+
+    HIP_CHECK(hipStreamSynchronize(obj.stream));
+    obj.threadCompleteStatus();
+  }
+  return obj.resultAfterAllIterations();
+}
+
+bool testhipMemsetD8AsyncWithMultiThread() {
+  MemSetAsyncMthreadTest <char> obj;
+  constexpr char memsetD8val = 0xDE;
+  obj.memAllocate(memsetD8val);
+  std::thread t[NUM_THREADS];
+
+  for (int i = 0 ; i < ITER ; i++) {
+    for (int k = 0 ; k < NUM_THREADS ; k++) {
+      if (k%2) {
+        t[k] = std::thread(queueJobsForhipMemsetD8Async<char>, obj.A_d,
+                           obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream);
+      } else {
+        t[k] = std::thread(queueJobsForhipMemsetD8Async<char>, obj.A_d,
+                           obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream);
+      }
+    }
+    for (int j = 0 ; j < NUM_THREADS ; j++) {
+      t[j].join();
+    }
+
+    HIP_CHECK(hipStreamSynchronize(obj.stream));
+    obj.threadCompleteStatus();
+  }
+  return obj.resultAfterAllIterations();
+}
+
+
+/*
+ * Test that validates functionality of hipmemsetAsync apis over multi threads
+ */
+TEST_CASE("Unit_hipMemsetAsync_QueueJobsMultithreaded") {
+  bool ret;
+
+  SECTION("hipMemsetAsync With MultiThread") {
+    ret = testhipMemsetAsyncWithMultiThread();
+    REQUIRE(ret == true);
+  }
+
+  SECTION("hipMemsetD32Async With MultiThread") {
+    ret = testhipMemsetD32AsyncWithMultiThread();
+    REQUIRE(ret == true);
+  }
+
+  SECTION("hipMemsetD16Async With MultiThread") {
+    ret = testhipMemsetD16AsyncWithMultiThread();
+    REQUIRE(ret == true);
+  }
+
+  SECTION("hipMemsetD8Async With MultiThread") {
+    ret = testhipMemsetD8AsyncWithMultiThread();
+    REQUIRE(ret == true);
+  }
+}

catch/unit/memory/hipMemsetInvalidPtr.cc

@@ -0,0 +1,127 @@
+/*
+ * Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+*/
+
+/**
+Testcase Scenarios :
+ 1) Test hipMemset apis with invalid pointer and invalid 2D pitch.
+ 2) Test hipMemsetAsync apis with invalid pointer and invalid 2D pitch.
+*/
+
+
+#include <hip_test_common.hh>
+
+#define N 50
+#define MEMSETVAL 0x42
+
+/**
+ * Testcase validates hipMemset apis behavior with
+ * invalid pointer and invalid 2D pitch value.
+ */
+TEST_CASE("Unit_hipMemset_InvalidPtrTests") {
+  hipError_t ret;
+  constexpr int Nbytes = N*sizeof(char);
+  char *A_d;
+
+  SECTION("hipMemset with null") {
+    ret = hipMemset(NULL, MEMSETVAL , Nbytes);
+    REQUIRE(ret != hipSuccess);
+  }
+
+  SECTION("hipMemset with hostptr") {
+    char *A_h;
+    A_h = reinterpret_cast<char *>(malloc(Nbytes));
+
+    ret = hipMemset(A_h, MEMSETVAL, Nbytes);
+    REQUIRE(ret != hipSuccess);
+
+    free(A_h);
+  }
+
+  SECTION("hipMemsetD32 with null") {
+    ret = hipMemsetD32(NULL, MEMSETVAL , Nbytes);
+    REQUIRE(ret != hipSuccess);
+  }
+
+  SECTION("hipMemsetD16 with null") {
+    ret = hipMemsetD16(NULL, MEMSETVAL , Nbytes);
+    REQUIRE(ret != hipSuccess);
+  }
+
+  SECTION("hipMemsetD8 with null") {
+    ret = hipMemsetD8(NULL, MEMSETVAL , Nbytes);
+    REQUIRE(ret != hipSuccess);
+  }
+
+  SECTION("hipMemset2D with null") {
+    constexpr size_t NUM_H = 256, NUM_W = 256;
+    size_t pitch_A;
+    size_t width = NUM_W * sizeof(char);
+
+    HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
+                                                             width , NUM_H));
+    ret = hipMemset2D(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H);
+    REQUIRE(ret != hipSuccess);
+
+    hipFree(A_d);
+  }
+}
+
+
+/**
+ * Testcase validates hipMemsetAsync apis behavior with
+ * invalid pointer and invalid 2D pitch value.
+ */
+TEST_CASE("Unit_hipMemsetAsync_InvalidPtrTests") {
+  hipError_t ret;
+  constexpr int Nbytes = N*sizeof(char);
+  char *A_d;
+
+  SECTION("hipMemsetAsync with null") {
+    ret = hipMemsetAsync(NULL, MEMSETVAL, Nbytes , 0);
+    REQUIRE(ret != hipSuccess);
+  }
+
+  SECTION("hipMemsetD32Async with null") {
+    ret = hipMemsetD32Async(NULL, MEMSETVAL , Nbytes, 0);
+    REQUIRE(ret != hipSuccess);
+  }
+
+  SECTION("hipMemsetD16Async with null") {
+    ret = hipMemsetD16Async(NULL, MEMSETVAL , Nbytes, 0);
+    REQUIRE(ret != hipSuccess);
+  }
+
+  SECTION("hipMemsetD8Async with null") {
+    ret = hipMemsetD8Async(NULL, MEMSETVAL , Nbytes, 0);
+    REQUIRE(ret != hipSuccess);
+  }
+
+  SECTION("hipMemset2DAsync with null") {
+    constexpr size_t NUM_H = 256, NUM_W = 256;
+    size_t pitch_A;
+    size_t width = NUM_W * sizeof(char);
+
+    HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
+                                                             width , NUM_H));
+    ret = hipMemset2DAsync(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H, 0);
+    REQUIRE(ret != hipSuccess);
+
+    hipFree(A_d);
+  }
+}