diff --git a/projects/hip/tests/catch/hipTestMain/config/config_amd_linux.json b/projects/hip/tests/catch/hipTestMain/config/config_amd_linux.json
index c11a1b823e..fc2c1331ed 100644
--- a/projects/hip/tests/catch/hipTestMain/config/config_amd_linux.json
+++ b/projects/hip/tests/catch/hipTestMain/config/config_amd_linux.json
@@ -1,6 +1,8 @@
 {
     "DisabledTests":
 	[
+	   "# Following test is related to ticket EXSWCPHIPT-41",
+	   "Unit_hipStreamGetPriority_happy",
 	   "Unit_hipStreamPerThread_DeviceReset_1"
 	]
 
diff --git a/projects/hip/tests/catch/hipTestMain/config/config_amd_windows.json b/projects/hip/tests/catch/hipTestMain/config/config_amd_windows.json
index 0fc51bb435..988f09d269 100644
--- a/projects/hip/tests/catch/hipTestMain/config/config_amd_windows.json
+++ b/projects/hip/tests/catch/hipTestMain/config/config_amd_windows.json
@@ -73,7 +73,8 @@
 	   "Unit_hipGraphAddEventRecordNode_Functional_WithFlags",
 	   "Unit_hipGraphAddEventRecordNode_MultipleRun",
 	   "Unit_hipMalloc3D_Negative",
-	   "Unit_hipPointerGetAttribute_MappedMem"
+	   "Unit_hipPointerGetAttribute_MappedMem",
+	   "# Following test is related to ticket EXSWCPHIPT-41",
+	   "Unit_hipStreamGetPriority_happy"
 	]
-
 }
diff --git a/projects/hip/tests/catch/include/hip_test_common.hh b/projects/hip/tests/catch/include/hip_test_common.hh
index f7dc215478..496f202b8c 100644
--- a/projects/hip/tests/catch/include/hip_test_common.hh
+++ b/projects/hip/tests/catch/include/hip_test_common.hh
@@ -27,6 +27,7 @@ THE SOFTWARE.
 
 #define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__);
 
+// Not thread-safe
 #define HIP_CHECK(error)                                                                           \
   {                                                                                                \
     hipError_t localError = error;                                                                 \
@@ -37,6 +38,20 @@ THE SOFTWARE.
     }                                                                                              \
   }
 
+// Check that an expression, errorExpr, evaluates to the expected error_t, expectedError.
+#define HIP_CHECK_ERROR(errorExpr, expectedError)                                                  \
+  {                                                                                                \
+    hipError_t localError = errorExpr;                                                             \
+    INFO("Matching Errors: "                                                                       \
+         << " Expected Error: " << hipGetErrorString(expectedError)                                \
+         << " Expected Code: " << expectedError << '\n'                                            \
+         << "                  Actual Error:   " << hipGetErrorString(localError)                  \
+         << " Actual Code:   " << localError << "\nStr: " << #errorExpr                            \
+         << "\nIn File: " << __FILE__ << " At line: " << __LINE__);                                \
+    REQUIRE(localError == expectedError);                                                          \
+  }
+
+// Not thread-safe
 #define HIPRTC_CHECK(error)                                                                        \
   {                                                                                                \
     auto localError = error;                                                                       \
@@ -46,25 +61,26 @@ THE SOFTWARE.
       REQUIRE(false);                                                                              \
     }                                                                                              \
   }
+
 // Although its assert, it will be evaluated at runtime
 #define HIP_ASSERT(x)                                                                              \
   { REQUIRE((x)); }
 
 #ifdef __cplusplus
-  #include <iostream>
-  #include <iomanip>
-  #include <chrono>
+#include <iostream>
+#include <iomanip>
+#include <chrono>
 #endif
 
 #define HIPCHECK(error)                                                                            \
-    {                                                                                              \
-        hipError_t localError = error;                                                             \
-        if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) {      \
-            printf("error: '%s'(%d) from %s at %s:%d\n", hipGetErrorString(localError),            \
-                   localError, #error, __FILE__, __LINE__);                                        \
-            abort();                                                                               \
-        }                                                                                          \
-    }
+  {                                                                                                \
+    hipError_t localError = error;                                                                 \
+    if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) {          \
+      printf("error: '%s'(%d) from %s at %s:%d\n", hipGetErrorString(localError), localError,      \
+             #error, __FILE__, __LINE__);                                                          \
+      abort();                                                                                     \
+    }                                                                                              \
+  }
 
 #define HIPASSERT(condition)                                                                       \
     if (!(condition)) {                                                                            \
@@ -104,8 +120,8 @@ static inline int getDeviceCount() {
 
 // Returns the current system time in microseconds
 static inline long long get_time() {
-  return std::chrono::high_resolution_clock::now().time_since_epoch()
-      /std::chrono::microseconds(1);
+  return std::chrono::high_resolution_clock::now().time_since_epoch() /
+      std::chrono::microseconds(1);
 }
 
 static inline double elapsed_time(long long startTimeUs, long long stopTimeUs) {
@@ -147,7 +163,16 @@ inline bool isImageSupported() {
   return imageSupport != 0;
 }
 
+/**
+ * Causes the test to stop and be skipped at runtime.
+ * reason: Message describing the reason the test has been skipped.
+ */
+static inline void HIP_SKIP_TEST(char const* const reason) noexcept {
+  // ctest is setup to parse for "HIP_SKIP_THIS_TEST", at which point it will skip the test.
+  std::cout << "Skipping test. Reason: " << reason << '\n' << "HIP_SKIP_THIS_TEST" << std::endl;
 }
+}
+
 
 // This must be called in the beginning of image test app's main() to indicate whether image
 // is supported.
diff --git a/projects/hip/tests/catch/stress/memory/hipMalloc.cc b/projects/hip/tests/catch/stress/memory/hipMalloc.cc
new file mode 100644
index 0000000000..a3057a6e4c
--- /dev/null
+++ b/projects/hip/tests/catch/stress/memory/hipMalloc.cc
@@ -0,0 +1,50 @@
+/*
+   Copyright (c) 2022 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, INNCLUDING 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 ANNY 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.
+ */
+
+#include <hip_test_common.hh>
+
+#include <memory>
+
+// Stress allocation tests
+// Try to allocate as much memory as possible
+// But since max allocation can fail, we need to be happy with atleast 1/4th of memory
+TEST_CASE("Stress_hipMalloc_HighSizeAlloc") {
+  size_t devMemTotal{0}, devMemFree{0};
+  HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemTotal));
+  REQUIRE(devMemFree > 0);
+  REQUIRE(devMemTotal > 0);
+
+  char* d_ptr{nullptr};
+  size_t counter{0};
+
+  INFO("Free Mem Available: " << devMemFree << " bytes out of " << devMemTotal << " bytes!");
+  while (hipMalloc(&d_ptr, devMemFree) != hipSuccess && devMemFree > 1) {
+    counter++;
+    devMemFree >>= 1;  // reduce the memory to be allocated by half
+    INFO("Attempt to allocate " << devMemFree << " bytes out of " << devMemTotal
+                                << " bytes failed!");
+    REQUIRE(counter <= 2);  // Make sure that we are atleast able to allocate 1/4th of max memory
+  }
+
+  HIP_CHECK(hipMemset(d_ptr, 1, devMemFree));
+  auto ptr = std::unique_ptr<unsigned char[]>{new unsigned char[devMemFree]};
+  HIP_CHECK(hipMemcpy(ptr.get(), d_ptr, devMemFree, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipFree(d_ptr));
+  REQUIRE(std::all_of(ptr.get(), ptr.get() + devMemFree, [](unsigned char n) { return n == 1; }));
+}
diff --git a/projects/hip/tests/catch/stress/memory/hipMallocManagedStress.cc b/projects/hip/tests/catch/stress/memory/hipMallocManagedStress.cc
index 7e992cd152..e582056685 100644
--- a/projects/hip/tests/catch/stress/memory/hipMallocManagedStress.cc
+++ b/projects/hip/tests/catch/stress/memory/hipMallocManagedStress.cc
@@ -150,7 +150,7 @@ static int HmmAttrPrint() {
 //  The following test case allocation, host access, device access of HMM
 //   memory from size 1 to 10KB
 
-TEST_CASE("Unit_hipMallocManaged_MultiSize") {
+TEST_CASE("Stress_hipMallocManaged_MultiSize") {
   IfTestPassed = true;
   int managed = HmmAttrPrint();
   if (managed == 1) {
@@ -196,7 +196,7 @@ TEST_CASE("Unit_hipMallocManaged_MultiSize") {
 // The following test case tests the behavior of kernel with a HMM memory and
 // hipMalloc memory
 
-TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") {
+TEST_CASE("Stress_hipMallocManaged_KrnlWth2MemTypes") {
   IfTestPassed = true;
   int *Hmm = NULL, *Dptr = NULL, InitVal = 123;
   size_t NumElms = (1024 * 1024);
@@ -241,7 +241,7 @@ TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") {
 
 // The following test case tests when the same Hmm memory is used for
 // launching multiple different kernels will results in any issue
-TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") {
+TEST_CASE("Stress_hipMallocManaged_MultiKrnlHmmAccess") {
   int managed = HmmAttrPrint();
   if (managed) {
     int InitVal = 123, NumElms = (1024 * 1024);
@@ -253,7 +253,7 @@ TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") {
 }
 
 // Testing the allocation of/scenarios around max possible memory
-TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") {
+TEST_CASE("Stress_hipMallocManaged_ExtremeSizes") {
   int managed = HmmAttrPrint();
   if (managed == 1) {
     bool IfTestPassed = true;
diff --git a/projects/hip/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc b/projects/hip/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc
index d34d88c846..d1c39b600a 100644
--- a/projects/hip/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc
+++ b/projects/hip/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc
@@ -26,117 +26,131 @@ THE SOFTWARE.
 */
 
 #include <hip_test_common.hh>
-#define NUM 1024
-#define SIZE 1024 * 4
+
+constexpr size_t NUM = 1024;
+constexpr size_t SIZE = 1024 * 4;
 
 __device__ int globalIn[NUM];
 __device__ int globalOut[NUM];
 
 __global__ void Assign(int* Out) {
-    int tid = threadIdx.x + blockIdx.x * blockDim.x;
-    Out[tid] = globalIn[tid];
-    globalOut[tid] = globalIn[tid];
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
+  Out[tid] = globalIn[tid];
+  globalOut[tid] = globalIn[tid];
 }
 
-__device__ __constant__ int globalConstArr[NUM];
-__device__ __constant__ static float statConstVar = 1.0f;
+__device__ __constant__ int globalConst[NUM];
 
-__global__ void checkGlobalConstAddress(int* addr, bool* out) {
-    *out = (globalConstArr == addr);
-}
+__global__ void checkAddress(int* addr, bool* out) { *out = (globalConst == addr); }
 
-__global__ void checkStaticConstVarAddress(float* addr, bool* out) {
-    *out = (&statConstVar == addr);
-}
-
-/**
- Calling hipMemcpyTo/FromSymbolAsync() using user declared stream obj and hipStreamPerThread.
- */
 TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
-    int *A, *Am, *B, *Ad, *C, *Cm;
-    A = new int[NUM];
-    B = new int[NUM];
-    C = new int[NUM];
-    for (int i = 0; i < NUM; i++) {
-        A[i] = -1 * i;
-        B[i] = 0;
-        C[i] = 0;
-    }
+  int *A{nullptr}, *Am{nullptr}, *B{nullptr}, *Ad{nullptr}, *C{nullptr}, *Cm{nullptr};
+  A = new int[NUM];
+  B = new int[NUM];
+  C = new int[NUM];
 
-    HIP_CHECK(hipMalloc(&Ad, SIZE));
-    HIP_CHECK(hipHostMalloc(&Am, SIZE));
-    HIP_CHECK(hipHostMalloc(&Cm, SIZE));
-    for (int i = 0; i < NUM; i++) {
-        Am[i] = -1 * i;
-        Cm[i] = 0;
-    }
+  HIP_CHECK(hipMalloc((void**)&Ad, SIZE));
+  HIP_CHECK(hipHostMalloc((void**)&Am, SIZE));
+  HIP_CHECK(hipHostMalloc((void**)&Cm, SIZE));
 
-    hipStream_t stream;
+  for (size_t i = 0; i < NUM; i++) {
+    A[i] = -1 * static_cast<int>(i);
+    B[i] = 0;
+    C[i] = 0;
+    Am[i] = -1 * static_cast<int>(i);
+    Cm[i] = 0;
+  }
+
+
+  SECTION("Calling hipMemcpyTo/FromSymbol using stream") {
+    hipStream_t stream{};
     HIP_CHECK(hipStreamCreate(&stream));
-    HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0,
-              hipMemcpyHostToDevice, stream));
+    HIP_CHECK(
+        hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, hipMemcpyHostToDevice, stream));
     HIP_CHECK(hipStreamSynchronize(stream));
     hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
     HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
-    HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0,
-              hipMemcpyDeviceToHost, stream));
+    HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost,
+                                       stream));
     HIP_CHECK(hipStreamSynchronize(stream));
-    for (int i = 0; i < NUM; i++) {
-        REQUIRE(Am[i] == B[i]);
-        REQUIRE(Am[i] == Cm[i]);
-    }
-
-    for (int i = 0; i < NUM; i++) {
-        A[i] = -2 * i;
-        B[i] = 0;
-    }
-
-    HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0,
-                                hipMemcpyHostToDevice));
-    hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
-    HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
-    HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0,
-                                  hipMemcpyDeviceToHost));
-    for (int i = 0; i < NUM; i++) {
-        REQUIRE(A[i] == B[i]);
-        REQUIRE(A[i] == C[i]);
-    }
-
-    for (int i = 0; i < NUM; i++) {
-        A[i] = -3 * i;
-        B[i] = 0;
-    }
-    SECTION("Calling hipMemcpyTo/FromSymbol using user declared stream obj") {
-      HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0,
-                                        hipMemcpyHostToDevice, stream));
-      HIP_CHECK(hipStreamSynchronize(stream));
-      hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
-      HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
-      HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0,
-                                         hipMemcpyDeviceToHost, stream));
-      HIP_CHECK(hipStreamSynchronize(stream));
-    }
-    SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") {
-      HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0,
-                                   hipMemcpyHostToDevice, hipStreamPerThread));
-      HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
-      hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
-      HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
-      HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0,
-                                   hipMemcpyDeviceToHost, hipStreamPerThread));
-      HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
-    }
-    for (int i = 0; i < NUM; i++) {
-        REQUIRE(A[i] == B[i]);
-        REQUIRE(A[i] == C[i]);
-    }
     HIP_CHECK(hipStreamDestroy(stream));
-    HIP_CHECK(hipHostFree(Am));
-    HIP_CHECK(hipHostFree(Cm));
-    HIP_CHECK(hipFree(Ad));
-    delete[] A;
-    delete[] B;
-    delete[] C;
+    for (size_t i = 0; i < NUM; i++) {
+      REQUIRE(Am[i] == B[i]);
+      REQUIRE(Am[i] == Cm[i]);
+    }
+  }
+
+  SECTION("Calling hipMemcpyTo/FromSymbol - validate value in host memory") {
+    HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice));
+    hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
+    HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
+    HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost));
+
+    for (size_t i = 0; i < NUM; i++) {
+      REQUIRE(A[i] == B[i]);
+      REQUIRE(A[i] == C[i]);
+    }
+  }
+
+  SECTION("Calling hipMemcpyTo/FromSymbol using user declared stream obj") {
+    hipStream_t stream{};
+    HIP_CHECK(hipStreamCreate(&stream));
+    HIP_CHECK(
+        hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+    hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
+    HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
+    HIP_CHECK(
+        hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+    HIP_CHECK(hipStreamDestroy(stream));
+
+    for (size_t i = 0; i < NUM; i++) {
+      REQUIRE(A[i] == B[i]);
+      REQUIRE(A[i] == C[i]);
+    }
+  }
+
+  SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") {
+    HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice,
+                                     hipStreamPerThread));
+    HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+    hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
+    HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
+    HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost,
+                                       hipStreamPerThread));
+    HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+
+    for (size_t i = 0; i < NUM; i++) {
+      REQUIRE(A[i] == B[i]);
+      REQUIRE(A[i] == C[i]);
+    }
+  }
+
+  // Check for address on GPU and CPU side and compare it
+  // If address mismatch report error
+  // Validate size of symbol as well, compare it with output of hipGetSymbolSize
+  SECTION("Validate address on GPU") {
+    bool* checkOkD{nullptr};
+    bool checkOk = false;
+    size_t symbolSize = 0;
+    int* symbolAddress{nullptr};
+    HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst)));
+    HIP_CHECK(hipGetSymbolAddress((void**)&symbolAddress, HIP_SYMBOL(globalConst)));
+    HIP_CHECK(hipMalloc((void**)&checkOkD, sizeof(bool)));
+    hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD);
+    HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
+    HIP_CHECK(hipFree(checkOkD));
+    HIP_ASSERT(checkOk);
+    HIP_ASSERT((symbolSize == SIZE));
+  }
+
+  HIP_CHECK(hipHostFree(Am));
+  HIP_CHECK(hipHostFree(Cm));
+  HIP_CHECK(hipFree(Ad));
+  delete[] A;
+  delete[] B;
+  delete[] C;
 }
 
 /**
@@ -144,39 +158,64 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
  2) Validate get symbol address/size for static const variable.
  */
 TEST_CASE("Unit_hipGetSymbolAddressAndSize_Validation") {
-    bool *checkOkD;
-    bool checkOk = false;
-    size_t symbolSize{};
-    int *symbolArrAddress{};
-    float *symbolVarAddress{};
+  bool* checkOkD{nullptr};
+  bool checkOk = false;
+  size_t symbolSize{};
+  int* symbolArrAddress{};
+  float* symbolVarAddress{};
 
-    SECTION("Validate symbol size/address of global const array") {
-      HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConstArr)));
-      HIP_CHECK(hipGetSymbolAddress(
-                              reinterpret_cast<void **>(&symbolArrAddress),
-                                              HIP_SYMBOL(globalConstArr)));
-      HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
-      hipLaunchKernelGGL(checkGlobalConstAddress, dim3(1, 1, 1), dim3(1, 1, 1),
-                                             0, 0, symbolArrAddress, checkOkD);
-      HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool),
-                                                       hipMemcpyDeviceToHost));
-      HIP_CHECK(hipFree(checkOkD));
-      HIP_ASSERT(checkOk);
-      HIP_ASSERT(symbolSize == SIZE);
-    }
+  SECTION("Validate symbol size/address of global const array") {
+    HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConstArr)));
+    HIP_CHECK(hipGetSymbolAddress(reinterpret_cast<void**>(&symbolArrAddress),
+                                  HIP_SYMBOL(globalConstArr)));
+    HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
+    hipLaunchKernelGGL(checkGlobalConstAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
+                       symbolArrAddress, checkOkD);
+    HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
+    HIP_CHECK(hipFree(checkOkD));
+    HIP_ASSERT(checkOk);
+    HIP_ASSERT(symbolSize == SIZE);
+  }
 
-    SECTION("Validate symbol size/address of static const variable") {
-      HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(statConstVar)));
-      HIP_CHECK(hipGetSymbolAddress(
-                               reinterpret_cast<void **>(&symbolVarAddress),
-                                                 HIP_SYMBOL(statConstVar)));
-      HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
-      hipLaunchKernelGGL(checkStaticConstVarAddress, dim3(1, 1, 1),
-                              dim3(1, 1, 1), 0, 0, symbolVarAddress, checkOkD);
-      HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool),
-                                                       hipMemcpyDeviceToHost));
-      HIP_CHECK(hipFree(checkOkD));
-      HIP_ASSERT(checkOk);
-      HIP_ASSERT(symbolSize == sizeof(float));
-    }
+  SECTION("Validate symbol size/address of static const variable") {
+    HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(statConstVar)));
+    HIP_CHECK(
+        hipGetSymbolAddress(reinterpret_cast<void**>(&symbolVarAddress), HIP_SYMBOL(statConstVar)));
+    HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
+    hipLaunchKernelGGL(checkStaticConstVarAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
+                       symbolVarAddress, checkOkD);
+    HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
+    HIP_CHECK(hipFree(checkOkD));
+    HIP_ASSERT(checkOk);
+    HIP_ASSERT(symbolSize == sizeof(float));
+  }
+}
+
+TEST_CASE("Unit_hipGetSymbolAddress_Negative") {
+  SECTION("Invalid symbol") {
+    int notADeviceSymbol{0};
+    int* addr{nullptr};
+    HIP_CHECK_ERROR(
+        hipGetSymbolAddress(reinterpret_cast<void**>(&addr), HIP_SYMBOL(notADeviceSymbol)),
+        hipErrorInvalidSymbol);
+  }
+
+  SECTION("Nullptr symbol") {
+    int* addr{nullptr};
+    HIP_CHECK_ERROR(hipGetSymbolAddress(reinterpret_cast<void**>(&addr), nullptr),
+                    hipErrorInvalidSymbol);
+  }
+}
+
+TEST_CASE("Unit_hipGetSymbolSize_Negative") {
+  SECTION("Invalid symbol") {
+    int notADeviceSymbol{0};
+    size_t dsize{0};
+    HIP_CHECK_ERROR(hipGetSymbolSize(&dsize, HIP_SYMBOL(notADeviceSymbol)), hipErrorInvalidSymbol);
+  }
+
+  SECTION("Nullptr symbol") {
+    size_t size{0};
+    HIP_CHECK_ERROR(hipGetSymbolSize(&size, nullptr), hipErrorInvalidSymbol);
+  }
 }
diff --git a/projects/hip/tests/catch/unit/memory/CMakeLists.txt b/projects/hip/tests/catch/unit/memory/CMakeLists.txt
index d4e821463c..3ce80dd780 100644
--- a/projects/hip/tests/catch/unit/memory/CMakeLists.txt
+++ b/projects/hip/tests/catch/unit/memory/CMakeLists.txt
@@ -42,14 +42,13 @@ set(TEST_SRC
     hipPointerGetAttributes.cc
     hipHostGetFlags.cc
     hipMallocManaged_MultiScenario.cc
-    hipMemsetInvalidPtr.cc
+    hipMemsetNegative.cc
     hipMemset.cc
     hipMemsetAsyncMultiThread.cc
     hipMemset3D.cc
     hipMemset2D.cc
     hipHostMallocTests.cc
     hipMemset3DFunctional.cc
-    hipMemset3DNegative.cc
     hipMemset3DRegressMultiThread.cc
     hipMallocManagedFlagsTst.cc
     hipMemPrefetchAsyncExtTsts.cc
@@ -108,14 +107,13 @@ set(TEST_SRC
     hipHostRegister.cc
     hipHostGetFlags.cc
     hipMallocManaged_MultiScenario.cc
-    hipMemsetInvalidPtr.cc
+    hipMemsetNegative.cc
     hipMemset.cc
     hipMemsetAsyncMultiThread.cc
     hipMemset3D.cc
     hipMemset2D.cc
     hipHostMallocTests.cc
     hipMemset3DFunctional.cc
-    hipMemset3DNegative.cc
     hipMemset3DRegressMultiThread.cc
     hipMallocManagedFlagsTst.cc
     hipMemPrefetchAsyncExtTsts.cc
diff --git a/projects/hip/tests/catch/unit/memory/hipMallocArray.cc b/projects/hip/tests/catch/unit/memory/hipMallocArray.cc
index 3e521bfcb5..a73566b40d 100644
--- a/projects/hip/tests/catch/unit/memory/hipMallocArray.cc
+++ b/projects/hip/tests/catch/unit/memory/hipMallocArray.cc
@@ -25,7 +25,6 @@ hipMallocArray API test scenarios
 4. Multithreaded scenario
 */
 
-
 #include <hip_test_common.hh>
 
 static constexpr auto NUM_W{4};
@@ -53,20 +52,16 @@ static void MallocArray_DiffSizes(int gpu) {
   std::vector<size_t> array_size;
   array_size.push_back(NUM_W);
   array_size.push_back(BIGNUM_W);
-  for (auto &size : array_size) {
+  for (auto& size : array_size) {
     hipArray* A_d[ARRAY_LOOP];
     size_t tot, avail, ptot, pavail;
     hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
     HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
     for (int i = 0; i < ARRAY_LOOP; i++) {
       if (size == NUM_W) {
-        HIP_CHECK(hipMallocArray(&A_d[i], &desc,
-              NUM_W, NUM_H,
-              hipArrayDefault));
+        HIP_CHECK(hipMallocArray(&A_d[i], &desc, NUM_W, NUM_H, hipArrayDefault));
       } else {
-        HIP_CHECK(hipMallocArray(&A_d[i], &desc,
-              BIGNUM_W, BIGNUM_H,
-              hipArrayDefault));
+        HIP_CHECK(hipMallocArray(&A_d[i], &desc, BIGNUM_W, BIGNUM_H, hipArrayDefault));
       }
     }
     for (int i = 0; i < ARRAY_LOOP; i++) {
@@ -79,10 +74,6 @@ static void MallocArray_DiffSizes(int gpu) {
   }
 }
 
-static void MallocArrayThreadFunc(int gpu) {
-  MallocArray_DiffSizes(gpu);
-}
-
 /*
  * This testcase verifies the negative scenarios of
  * hipMallocArray API
@@ -92,56 +83,31 @@ TEST_CASE("Unit_hipMallocArray_Negative") {
   hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
 #if HT_NVIDIA
   SECTION("NullPointer to Array") {
-    REQUIRE(hipMallocArray(nullptr, &desc,
-            NUM_W, NUM_H, hipArrayDefault) != hipSuccess);
+    REQUIRE(hipMallocArray(nullptr, &desc, NUM_W, NUM_H, hipArrayDefault) != hipSuccess);
   }
 
   SECTION("NullPointer to Channel Descriptor") {
-    REQUIRE(hipMallocArray(&A_d, nullptr,
-            NUM_W, NUM_H, hipArrayDefault) != hipSuccess);
+    REQUIRE(hipMallocArray(&A_d, nullptr, NUM_W, NUM_H, hipArrayDefault) != hipSuccess);
   }
 #endif
   SECTION("Width 0 in hipMallocArray") {
-    REQUIRE(hipMallocArray(&A_d, &desc,
-            0, NUM_H, hipArrayDefault) != hipSuccess);
+    REQUIRE(hipMallocArray(&A_d, &desc, 0, NUM_H, hipArrayDefault) != hipSuccess);
   }
 
   SECTION("Height 0 in hipMallocArray") {
-    REQUIRE(hipMallocArray(&A_d, &desc,
-            NUM_W, 0, hipArrayDefault) == hipSuccess);
+    REQUIRE(hipMallocArray(&A_d, &desc, NUM_W, 0, hipArrayDefault) == hipSuccess);
   }
 
-  SECTION("Invalid Flag") {
-    REQUIRE(hipMallocArray(&A_d, &desc,
-            NUM_W, NUM_H, 100) != hipSuccess);
-  }
+  SECTION("Invalid Flag") { REQUIRE(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, 100) != hipSuccess); }
 
   SECTION("Max int values") {
-    REQUIRE(hipMallocArray(&A_d, &desc,
-            std::numeric_limits<int>::max(),
-            std::numeric_limits<int>::max(),
-            hipArrayDefault) != hipSuccess);
+    REQUIRE(hipMallocArray(&A_d, &desc, std::numeric_limits<int>::max(),
+                           std::numeric_limits<int>::max(), hipArrayDefault) != hipSuccess);
   }
 }
-/*
- * This testcase verifies the basic scenario of
- * hipMallocArray API for different datatypes
- * of size 10
- */
-TEMPLATE_TEST_CASE("Unit_hipMallocArray_Basic",
-                   "", int, unsigned int, float) {
-  hipArray* A_d;
-  hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
-  REQUIRE(hipMallocArray(&A_d, &desc,
-                         NUM_W, NUM_H,
-                         hipArrayDefault) == hipSuccess);
-  HIP_CHECK(hipFreeArray(A_d));
-}
 
 
-TEST_CASE("Unit_hipMallocArray_DiffSizes") {
-  MallocArray_DiffSizes(0);
-}
+TEST_CASE("Unit_hipMallocArray_DiffSizes") { MallocArray_DiffSizes(0); }
 
 
 /*
@@ -156,10 +122,10 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") {
   size_t tot, avail, ptot, pavail;
   HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
   for (int i = 0; i < devCnt; i++) {
-    threadlist.push_back(std::thread(MallocArrayThreadFunc, i));
+    threadlist.push_back(std::thread(MallocArray_DiffSizes, i));
   }
 
-  for (auto &t : threadlist) {
+  for (auto& t : threadlist) {
     t.join();
   }
   HIP_CHECK(hipMemGetInfo(&avail, &tot));
@@ -170,3 +136,266 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") {
   }
 }
 
+
+constexpr size_t BlockSize = 16;
+
+template <class T, size_t N> struct type_and_size {
+  using type = T;
+  static constexpr size_t size = N;
+};
+
+// scalars are interpreted as a vector of 1 length.
+// template <size_t N> using int_constant = std::integral_constant<size_t, N>;
+template <typename T> struct vector_info;
+template <> struct vector_info<int> : type_and_size<int, 1> {};
+template <> struct vector_info<float> : type_and_size<float, 1> {};
+template <> struct vector_info<short> : type_and_size<short, 1> {};
+template <> struct vector_info<char> : type_and_size<char, 1> {};
+template <> struct vector_info<unsigned int> : type_and_size<unsigned int, 1> {};
+template <> struct vector_info<unsigned short> : type_and_size<unsigned short, 1> {};
+template <> struct vector_info<unsigned char> : type_and_size<unsigned char, 1> {};
+
+template <> struct vector_info<int2> : type_and_size<int, 2> {};
+template <> struct vector_info<float2> : type_and_size<float, 2> {};
+template <> struct vector_info<short2> : type_and_size<short, 2> {};
+template <> struct vector_info<char2> : type_and_size<char, 2> {};
+template <> struct vector_info<uint2> : type_and_size<unsigned int, 2> {};
+template <> struct vector_info<ushort2> : type_and_size<unsigned short, 2> {};
+template <> struct vector_info<uchar2> : type_and_size<unsigned char, 2> {};
+
+template <> struct vector_info<int4> : type_and_size<int, 4> {};
+template <> struct vector_info<float4> : type_and_size<float, 4> {};
+template <> struct vector_info<short4> : type_and_size<short, 4> {};
+template <> struct vector_info<char4> : type_and_size<char, 4> {};
+template <> struct vector_info<uint4> : type_and_size<unsigned int, 4> {};
+template <> struct vector_info<ushort4> : type_and_size<unsigned short, 4> {};
+template <> struct vector_info<uchar4> : type_and_size<unsigned char, 4> {};
+
+// Kernels ///////////////////////////////////////
+
+// read from a texture using normalized coordinates
+constexpr size_t ChannelToRead = 1;
+template <typename T>
+__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height,
+                                bool textureGather) {
+  // Calculate normalized texture coordinates
+  const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
+  const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
+  const float u = x / (float)width;
+
+  // Read from texture and write to global memory
+  if (height == 0) {
+    output[x] = tex1D<T>(texObj, u);
+  } else {
+    const float v = y / (float)height;
+    output[y * width + x] =
+        textureGather ? tex2Dgather<T>(texObj, u, v, ChannelToRead) : tex2D<T>(texObj, u, v);
+  }
+}
+
+template <typename T> __device__ void addOne(T* a) {
+  using scalar_type = typename vector_info<T>::type;
+  auto as = reinterpret_cast<scalar_type*>(a);
+  for (size_t i = 0; i < vector_info<T>::size; ++i) {
+    as[i] = as[i] + static_cast<scalar_type>(1);
+  }
+}
+
+// read from a surface and write to another
+template <typename T> __global__ void incSurface(hipSurfaceObject_t surf, size_t height) {
+  // Calculate surface coordinates
+  unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
+  unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
+  if (height == 0) {
+    T data;
+    surf1Dread(&data, surf, x * sizeof(T));
+    addOne(&data);  // change the value to show that write works
+    surf1Dwrite(data, surf, x * sizeof(T));
+  } else {
+    T data;
+    surf2Dread(&data, surf, x * sizeof(T), y);
+    addOne(&data);  // change the value to show that write works
+    surf2Dwrite(data, surf, x * sizeof(T), y);
+  }
+}
+
+// Helpers ///////////////////////////////////////
+
+template <typename T> size_t getAllocSize(const size_t width, const size_t height) noexcept {
+  return sizeof(T) * width * (height ? height : 1);
+}
+
+template <typename T> void checkDataIsAscending(const std::vector<T>& hostData) {
+  bool allMatch = true;
+  size_t i = 0;
+  for (; i < hostData.size(); ++i) {
+    allMatch = allMatch && hostData[i] == static_cast<T>(i);
+    if (!allMatch) break;
+  }
+  INFO("hostData[" << i << "] == " << static_cast<T>(hostData[i]));
+  REQUIRE(allMatch);
+}
+
+// Tests /////////////////////////////////////////
+
+// Test the default array by generating a texture from it then reading from that texture.
+// Textures are read-only so write to the array then copy from the texture into normal device memory
+template <typename T>
+void testArrayAsTexture(hipArray_t arrayPtr, const size_t width, const size_t height) {
+  using scalar_type = typename vector_info<T>::type;
+  constexpr auto vec_size = vector_info<T>::size;
+
+  const auto h = height ? height : 1;
+  const size_t pitch = width * sizeof(T);  // no padding
+  const auto size = pitch * h;
+
+  // create an array to initialize the hip array, then later use it to hold the result
+  std::vector<scalar_type> hostData(width * h * vec_size);
+
+  // Setup backing array
+  // assign ascending values to the data array to show indexing is working.
+  std::iota(std::begin(hostData), std::end(hostData), 0);
+  HIP_CHECK(
+      hipMemcpy2DToArray(arrayPtr, 0, 0, hostData.data(), pitch, pitch, h, hipMemcpyHostToDevice));
+
+
+  // create texture
+  hipTextureObject_t textObj{};
+  hipResourceDesc resDesc{};
+  memset(&resDesc, 0, sizeof(hipResourceDesc));
+  // enum to store how to resDesc.res union is being used
+  resDesc.resType = hipResourceTypeArray;
+  resDesc.res.array.array = arrayPtr;
+
+  hipTextureDesc textDesc{};
+  memset(&textDesc, 0, sizeof(hipTextureDesc));
+  textDesc.filterMode =
+      hipFilterModePoint;  // use the actual values in the texture, not normalized data
+  textDesc.readMode = hipReadModeElementType;  // don't convert the data to floats
+  textDesc.normalizedCoords = 1;               // use normalized coordinates (0.0-1.0)
+
+  HIP_CHECK(hipCreateTextureObject(&textObj, &resDesc, &textDesc, nullptr));
+
+
+  // run kernel
+  T* device_data{};
+  HIP_CHECK(hipMalloc(&device_data, size));
+  readFromTexture<<<dim3(width / BlockSize, height ? height / BlockSize : 1, 1),
+                    dim3(BlockSize, height ? BlockSize : 1, 1)>>>(device_data, textObj, width,
+                                                                  height, false);
+  HIP_CHECK(hipGetLastError());  // check for errors when running the kernel
+
+  // copy data back and then test it
+  std::fill(std::begin(hostData), std::end(hostData), 0);
+  HIP_CHECK(hipMemcpy(hostData.data(), device_data, size, hipMemcpyDeviceToHost));
+
+  checkDataIsAscending(hostData);
+
+  // clean up
+  HIP_CHECK(hipDestroyTextureObject(textObj));
+  HIP_CHECK(hipFree(device_data));
+}
+
+// Test the an array created with the SurfaceLoadStore flag by generating a surface and reading from
+// it and writing to it.
+template <typename T>
+void testArrayAsSurface(hipArray_t arrayPtr, const size_t width, const size_t height) {
+  using scalar_type = typename vector_info<T>::type;
+  constexpr auto vec_size = vector_info<T>::size;
+
+  const auto h = height ? height : 1;
+  const size_t pitch = width * sizeof(T);  // no padding
+  const auto size = pitch * h;
+
+  std::vector<scalar_type> hostData(width * h * vec_size);
+
+  // Setup backing array
+  // assign ascending values to the data array to show indexing is working.
+  std::iota(std::begin(hostData), std::end(hostData), 0);
+  HIP_CHECK(
+      hipMemcpy2DToArray(arrayPtr, 0, 0, hostData.data(), pitch, pitch, h, hipMemcpyHostToDevice));
+
+
+  // create surface
+  hipSurfaceObject_t surfObj{};
+  hipResourceDesc resDesc;
+  memset(&resDesc, 0, sizeof(hipResourceDesc));
+  resDesc.resType = hipResourceTypeArray;
+
+  resDesc.res.array.array = arrayPtr;
+  HIP_CHECK(hipCreateSurfaceObject(&surfObj, &resDesc));
+
+
+  // run kernel
+  T* device_data{};
+  HIP_CHECK(hipMalloc(&device_data, size));
+  // This will increment the values of the surface, so this is undone later
+  incSurface<T><<<dim3(width / BlockSize, height ? height / BlockSize : 1, 1),
+                  dim3(BlockSize, height ? BlockSize : 1, 1)>>>(surfObj, height);
+  HIP_CHECK(hipGetLastError());  // check for errors when running the kernel
+
+
+  // copy data back and then test it
+  std::fill(std::begin(hostData), std::end(hostData), 0);
+  HIP_CHECK(hipMemcpy2DFromArray(hostData.data(), pitch, arrayPtr, 0, 0, pitch, h,
+                                 hipMemcpyDeviceToHost));
+
+
+  // undo the increment
+  std::for_each(std::begin(hostData), std::end(hostData),
+                [](scalar_type& x) { x -= static_cast<scalar_type>(1); });
+  checkDataIsAscending(hostData);
+
+  // clean up
+  HIP_CHECK(hipDestroySurfaceObject(surfObj));
+  HIP_CHECK(hipFree(device_data));
+}
+
+size_t getFreeMem() {
+  size_t free = 0, total = 0;
+  HIP_CHECK(hipMemGetInfo(&free, &total));
+  return free;
+}
+
+// The happy path of a default array and a SurfaceLoadStore array should work
+// Selection of types chosen to reduce compile times
+TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, short2, char, uchar2,
+                   char4, float, float2, float4) {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-62");
+#endif
+
+  hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
+
+  size_t init_free = getFreeMem();
+
+  // pointer to the array in device memory
+  hipArray_t arrayPtr{};
+  size_t width = 1024;
+  size_t height = GENERATE(0, 1024);
+
+  SECTION("hipArrayDefault") {
+    INFO("flag is hipArrayDefault");
+    INFO("height: " << height);
+
+    HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArrayDefault));
+    testArrayAsTexture<TestType>(arrayPtr, width, height);
+  }
+#if HT_NVIDIA  // surfaces and texture gather not supported on AMD
+  SECTION("hipArraySurfaceLoadStore") {
+    INFO("flag is hipArraySurfaceLoadStore");
+    INFO("height: " << height);
+
+    HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArraySurfaceLoadStore));
+    testArrayAsSurface<TestType>(arrayPtr, width, height);
+  }
+#endif
+
+  size_t final_free = getFreeMem();
+
+  const size_t alloc_size = getAllocSize<TestType>(width, height);
+  // alloc will be chunked, so this is not exact
+  REQUIRE(init_free - final_free >= alloc_size);
+
+  HIP_CHECK(hipFreeArray(arrayPtr));
+}
diff --git a/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc b/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc
index 98bc2ab014..592ce86897 100644
--- a/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc
+++ b/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2022 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
@@ -51,14 +51,13 @@ Testcase Scenarios :
 #include <hip_test_checkers.hh>
 #include <hip_test_kernels.hh>
 
-
-#include <vector>
-#include <limits>
 #include <atomic>
+#include <limits>
+#include <vector>
 
 
 /* Buffer size for bigger chunks in alloc/free cycles */
-static constexpr auto BuffSizeBC = 5*1024*1024;
+static constexpr auto BuffSizeBC = 5 * 1024 * 1024;
 
 /* Buffer size for smaller chunks in alloc/free cycles */
 static constexpr auto BuffSizeSC = 16;
@@ -68,19 +67,18 @@ static constexpr auto BuffSizeSC = 16;
 static constexpr auto NumDiv = 100;
 
 /* Max alloc/free iterations for smaller chunks */
-static constexpr auto MaxAllocFree_SmallChunks = (5000000/NumDiv);
+static constexpr auto MaxAllocFree_SmallChunks = (5000000 / NumDiv);
 
 /* Max alloc/free iterations for bigger chunks */
 static constexpr auto MaxAllocFree_BigChunks = 10000;
 
 /* Max alloc and pool iterations */
-static constexpr auto MaxAllocPoolIter = (2000000/NumDiv);
+static constexpr auto MaxAllocPoolIter = (2000000 / NumDiv);
 
 /* Test status shared across threads */
 static std::atomic<bool> g_thTestPassed{true};
 
 
-
 /**
  * Validates data consistency on supplied gpu
  */
@@ -103,9 +101,8 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
   HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
   HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
 
-  hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
-                     0, 0, static_cast<const int*>(A_d),
-                     static_cast<const int*>(B_d), C_d, N);
+  hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
+                     static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
 
   HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
 
@@ -121,9 +118,8 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
 
   if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
     // In concurrent calls on one GPU, we cannot verify leaking in this way
-    UNSCOPED_INFO(
-        "validateMemoryOnGPU : Memory allocation mismatch observed."
-        << "Possible memory leak.");
+    UNSCOPED_INFO("validateMemoryOnGPU : Memory allocation mismatch observed."
+                  << "Possible memory leak.");
     TestPassed = false;
   }
 
@@ -138,7 +134,7 @@ static bool regressAllocInLoop(int gpu) {
   bool TestPassed = true;
   size_t tot, avail, ptot, pavail, numBytes;
   int i = 0;
-  int *ptr;
+  int* ptr;
 
   HIP_CHECK(hipSetDevice(gpu));
   numBytes = BuffSizeBC;
@@ -150,11 +146,12 @@ static bool regressAllocInLoop(int gpu) {
     HIP_CHECK(hipMemGetInfo(&avail, &tot));
     HIP_CHECK(hipFree(ptr));
 
-    if (pavail-avail < numBytes) {  // We expect pavail-avail >= numBytes
-      UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " <<
-        numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" <<
-        pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" <<
-        tot << ", pavail-avail=" << pavail-avail);
+    if (pavail - avail < numBytes) {  // We expect pavail-avail >= numBytes
+      UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " << numBytes
+                                      << " not matching with hipMemGetInfo - FAIL."
+                                      << "pavail=" << pavail << ", ptot=" << ptot
+                                      << ", avail=" << avail << ", tot=" << tot
+                                      << ", pavail-avail=" << pavail - avail);
       TestPassed = false;
       break;
     }
@@ -173,8 +170,8 @@ static bool regressAllocInLoop(int gpu) {
   HIP_CHECK(hipMemGetInfo(&avail, &tot));
 
   if ((pavail != avail) || (ptot != tot)) {
-    UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." <<
-      "Possible memory leak.");
+    UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed."
+                  << "Possible memory leak.");
     TestPassed &= false;
   }
 
@@ -203,9 +200,8 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
   HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
   HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
 
-  hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
-                     0, 0, static_cast<const int*>(A_d),
-                     static_cast<const int*>(B_d), C_d, N);
+  hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
+                     static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
 
   HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
 
@@ -238,7 +234,7 @@ static bool regressAllocInLoopMthread(int gpu) {
   bool TestPassed = true;
   size_t tot, avail, ptot, pavail, numBytes;
   int i = 0;
-  int *ptr;
+  int* ptr;
 
   HIPCHECK(hipSetDevice(gpu));
   numBytes = BuffSizeBC;
@@ -250,11 +246,12 @@ static bool regressAllocInLoopMthread(int gpu) {
     HIPCHECK(hipMemGetInfo(&avail, &tot));
     HIPCHECK(hipFree(ptr));
 
-    if (pavail-avail < numBytes) {  // We expect pavail-avail >= numBytes
-      UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " <<
-        numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" <<
-        pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" <<
-        tot << ", pavail-avail=" << pavail-avail);
+    if (pavail - avail < numBytes) {  // We expect pavail-avail >= numBytes
+      UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " << numBytes
+                                      << " not matching with hipMemGetInfo - FAIL."
+                                      << "pavail=" << pavail << ", ptot=" << ptot
+                                      << ", avail=" << avail << ", tot=" << tot
+                                      << ", pavail-avail=" << pavail - avail);
       TestPassed = false;
       break;
     }
@@ -273,8 +270,8 @@ static bool regressAllocInLoopMthread(int gpu) {
   HIPCHECK(hipMemGetInfo(&avail, &tot));
 
   if ((pavail != avail) || (ptot != tot)) {
-    UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." <<
-      "Possible memory leak.");
+    UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed."
+                  << "Possible memory leak.");
     TestPassed &= false;
   }
 
@@ -285,18 +282,15 @@ static bool regressAllocInLoopMthread(int gpu) {
  * Thread func to regress alloc and check data consistency
  */
 static void threadFunc(int gpu) {
-  g_thTestPassed = regressAllocInLoopMthread(gpu)
-                          && validateMemoryOnGpuMThread(gpu);
+  g_thTestPassed = regressAllocInLoopMthread(gpu) && validateMemoryOnGpuMThread(gpu);
 
-  UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" <<
-                                         g_thTestPassed.load());
+  UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" << g_thTestPassed.load());
 }
 
 
 /* Performs Argument Validation of api */
 TEST_CASE("Unit_hipMalloc_ArgumentValidation") {
-  int *ptr;
-  hipError_t ret;
+  int* ptr{nullptr};
 
   SECTION("hipMalloc() when size(0)") {
     HIP_CHECK(hipMalloc(&ptr, 0));
@@ -304,21 +298,17 @@ TEST_CASE("Unit_hipMalloc_ArgumentValidation") {
     REQUIRE(ptr == nullptr);
   }
 
-  SECTION("hipFree() when freeing nullptr ") {
-    ptr = nullptr;
-    // api should return success and shudnt crash
+  SECTION("hipFree() when freeing nullptr") {
     HIP_CHECK(hipFree(ptr));
   }
 
   SECTION("hipMalloc() with invalid argument") {
-    constexpr auto sizeBytes = 100;
-    ret = hipMalloc(nullptr, sizeBytes);
-    REQUIRE(ret != hipSuccess);
+    HIP_CHECK_ERROR(hipMalloc(nullptr, 100), hipErrorInvalidValue);
   }
 
   SECTION("hipMalloc() with max size_t") {
-    ret = hipMalloc(&ptr, std::numeric_limits<std::size_t>::max());
-    REQUIRE(ret != hipSuccess);
+    HIP_CHECK_ERROR(hipMalloc(&ptr, std::numeric_limits<std::size_t>::max()),
+                    hipErrorMemoryAllocation);
   }
 }
 
@@ -344,12 +334,12 @@ TEST_CASE("Unit_hipMalloc_LoopRegressionAllocFreeCycles") {
  * of time.
  */
 TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
-  size_t avail, tot, pavail, ptot;
-  bool ret;
-  hipError_t err;
-  std::vector<int *> ptrlist;
+  size_t avail{0}, tot{0}, pavail{0}, ptot{0};
+  bool ret{false};
+  hipError_t err{};
+  std::vector<int*> ptrlist{};
   constexpr auto BuffSize = 10;
-  int devCnt, *ptr;
+  int devCnt{0}, *ptr{nullptr};
 
   // Get GPU count
   HIP_CHECK(hipGetDeviceCount(&devCnt));
@@ -358,14 +348,13 @@ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
   HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
 
   // Allocate small chunks of memory million times
-  for (int i = 0; i < MaxAllocPoolIter ; i++) {
+  for (int i = 0; i < MaxAllocPoolIter; i++) {
     if ((err = hipMalloc(&ptr, BuffSize)) != hipSuccess) {
       HIP_CHECK(hipMemGetInfo(&avail, &tot));
 
-      INFO("Loop regression pool allocation failure. " <<
-      "Total gpu memory " << tot/(1024.0*1024.0) <<", Free memory " <<
-      avail/(1024.0*1024.0) << " iter " << i << " error "
-      << hipGetErrorString(err));
+      INFO("Loop regression pool allocation failure. "
+           << "Total gpu memory " << tot / (1024.0 * 1024.0) << ", Free memory "
+           << avail / (1024.0 * 1024.0) << " iter " << i << " error " << hipGetErrorString(err));
 
       REQUIRE(false);
     }
@@ -375,7 +364,7 @@ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
   }
 
   // Free ptrs at later point of time
-  for ( auto &t : ptrlist ) {
+  for (auto& t : ptrlist) {
     HIP_CHECK(hipFree(t));
   }
 
@@ -404,7 +393,7 @@ TEST_CASE("Unit_hipMalloc_Multithreaded_MultiGPU") {
     threadlist.push_back(std::thread(threadFunc, i));
   }
 
-  for (auto &t : threadlist) {
+  for (auto& t : threadlist) {
     t.join();
   }
 
diff --git a/projects/hip/tests/catch/unit/memory/hipMemset3DNegative.cc b/projects/hip/tests/catch/unit/memory/hipMemset3DNegative.cc
deleted file mode 100644
index 798c90e1ac..0000000000
--- a/projects/hip/tests/catch/unit/memory/hipMemset3DNegative.cc
+++ /dev/null
@@ -1,236 +0,0 @@
-/*
-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 hipMemset3D() with uninitialized devPitchedPtr.
- 2) Test hipMemset3DAsync() with uninitialized devPitchedPtr.
-
- 3) Reset devPitchedPtr to zero and check return value for hipMemset3D().
- 4) Reset devPitchedPtr to zero and check return value for hipMemset3DAsync().
-
- 5) Test hipMemset3D() with extent.width as max size_t and keeping height,
- depth as valid values.
- 6) Test hipMemset3DAsync() with extent.width as max size_t and keeping height,
- depth as valid values.
- 7) Test hipMemset3D() with extent.height as max size_t and keeping width,
- depth as valid values.
- 8) Test hipMemset3DAsync() with extent.height as max size_t and keeping width,
- depth as valid values.
- 9) Test hipMemset3D() with extent.depth as max size_t and keeping height,
- width as valid values.
-10) Test hipMemset3DAsync() with extent.depth as max size_t and keeping height,
- width as valid values.
-
-11) Device Ptr out bound and extent(0) passed for hipMemset3D().
-12) Device Ptr out bound and extent(0) passed for hipMemset3DAsync().
-
-13) Device Ptr out bound and valid extent passed for hipMemset3D().
-14) Device Ptr out bound and valid extent passed for hipMemset3DAsync().
-*/
-
-#include <hip_test_common.hh>
-
-TEST_CASE("Unit_hipMemset3D_Negative") {
-  hipError_t ret;
-  hipPitchedPtr devPitchedPtr;
-  constexpr int memsetval = 1;
-  constexpr size_t numH = 256, numW = 256;
-  constexpr size_t depth = 10;
-  constexpr size_t width = numW * sizeof(char);
-  hipExtent extent = make_hipExtent(width, numH, depth);
-
-  HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
-
-  SECTION("Using uninitialized devpitched ptr") {
-    hipPitchedPtr devPitchedUnPtr;
-
-    ret = hipMemset3D(devPitchedUnPtr, memsetval, extent);
-    REQUIRE(ret != hipSuccess);
-  }
-
-  SECTION("Reset devPitchedPtr to zero") {
-    hipPitchedPtr rdevPitchedPtr{};
-
-    ret = hipMemset3D(rdevPitchedPtr, memsetval, extent);
-    REQUIRE(ret != hipSuccess);
-  }
-
-  SECTION("Pass extent fields as max size_t") {
-    hipExtent extMW = make_hipExtent(std::numeric_limits<std::size_t>::max(),
-                                     numH,
-                                     depth);
-    hipExtent extMH = make_hipExtent(width,
-                                     std::numeric_limits<std::size_t>::max(),
-                                     depth);
-    hipExtent extMD = make_hipExtent(width,
-                                     numH,
-                                     std::numeric_limits<std::size_t>::max());
-
-    ret = hipMemset3D(devPitchedPtr, memsetval, extMW);
-    REQUIRE(ret != hipSuccess);
-
-    ret = hipMemset3D(devPitchedPtr, memsetval, extMH);
-    REQUIRE(ret != hipSuccess);
-
-    if ((TestContext::get()).isAmd()) {
-      ret = hipMemset3D(devPitchedPtr, memsetval, extMD);
-      REQUIRE(ret != hipSuccess);
-    } else {
-      WARN("Test is skipped for max depth."
-           << "Cuda doesn't check the maximum depth of extent field");
-    }
-  }
-
-  SECTION("Device Ptr out bound and extent(0) passed for memset") {
-    size_t pitch = devPitchedPtr.pitch;
-    size_t slicePitch = pitch * extent.height;
-    constexpr auto advanceOffset = 10;
-
-    // Point devptr to end of allocated memory
-    char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
-                       + depth * slicePitch;
-
-    // Advance devptr further to go out of boundary
-    devPtrMod = devPtrMod + advanceOffset;
-    hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
-                                                    numW * sizeof(char), numH);
-    hipExtent extent0{};
-    ret = hipMemset3D(modDevPitchedPtr, memsetval, extent0);
-
-    // api expected to check extent0 and return success before going for ptr.
-    REQUIRE(ret == hipSuccess);
-  }
-
-  SECTION("Device Ptr out bound and valid extent passed for memset") {
-    size_t pitch = devPitchedPtr.pitch;
-    size_t slicePitch = pitch * extent.height;
-    constexpr auto advanceOffset = 10;
-
-    // Point devptr to end of allocated memory
-    char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
-                       + depth * slicePitch;
-
-    // Advance devptr further to go out of boundary
-    devPtrMod = devPtrMod + advanceOffset;
-    hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
-                                                    numW * sizeof(char), numH);
-    ret = hipMemset3D(modDevPitchedPtr, memsetval, extent);
-
-    REQUIRE(ret != hipSuccess);
-  }
-
-  HIP_CHECK(hipFree(devPitchedPtr.ptr));
-}
-
-TEST_CASE("Unit_hipMemset3DAsync_Negative") {
-  hipError_t ret;
-  hipPitchedPtr devPitchedPtr;
-  hipStream_t stream;
-  constexpr int memsetval = 1;
-  constexpr size_t numH = 256;
-  constexpr size_t numW = 256;
-  constexpr size_t depth = 10;
-  constexpr size_t width = numW * sizeof(char);
-  hipExtent extent = make_hipExtent(width, numH, depth);
-
-  HIP_CHECK(hipStreamCreate(&stream));
-  HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
-
-  SECTION("Using uninitialized devpitched ptr") {
-    hipPitchedPtr devPitchedUnPtr;
-
-    ret = hipMemset3DAsync(devPitchedUnPtr, memsetval, extent, stream);
-    REQUIRE(ret != hipSuccess);
-  }
-
-  SECTION("Reset devPitchedPtr to zero") {
-    hipPitchedPtr rdevPitchedPtr{};
-
-    ret = hipMemset3DAsync(rdevPitchedPtr, memsetval, extent, stream);
-    REQUIRE(ret != hipSuccess);
-  }
-
-  SECTION("Pass extent fields as max size_t") {
-    hipExtent extMW = make_hipExtent(std::numeric_limits<std::size_t>::max(),
-                                     numH,
-                                     depth);
-    hipExtent extMH = make_hipExtent(width,
-                                     std::numeric_limits<std::size_t>::max(),
-                                     depth);
-    hipExtent extMD = make_hipExtent(width,
-                                     numH,
-                                     std::numeric_limits<std::size_t>::max());
-
-    ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMW, stream);
-    REQUIRE(ret != hipSuccess);
-
-    ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMH, stream);
-    REQUIRE(ret != hipSuccess);
-
-    if ((TestContext::get()).isAmd()) {
-      ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMD, stream);
-      REQUIRE(ret != hipSuccess);
-    } else {
-      WARN("Test is skipped for max depth."
-           << "Cuda doesn't check the maximum depth of extent field");
-    }
-  }
-
-  SECTION("Device Ptr out bound and extent(0) passed for memset") {
-    size_t pitch = devPitchedPtr.pitch;
-    size_t slicePitch = pitch * extent.height;
-    constexpr auto advanceOffset = 10;
-
-    // Point devptr to end of allocated memory
-    char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
-                       + depth * slicePitch;
-
-    // Advance devptr further to go out of boundary
-    devPtrMod = devPtrMod + advanceOffset;
-    hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
-                                                    numW * sizeof(char), numH);
-    hipExtent extent0{};
-    ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent0, stream);
-
-    // api expected to check extent0 and return success before going for ptr.
-    REQUIRE(ret == hipSuccess);
-  }
-
-  SECTION("Device Ptr out bound and valid extent passed for memset") {
-    size_t pitch = devPitchedPtr.pitch;
-    size_t slicePitch = pitch * extent.height;
-    constexpr auto advanceOffset = 10;
-
-    // Point devptr to end of allocated memory
-    char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
-                       + depth * slicePitch;
-
-    // Advance devptr further to go out of boundary
-    devPtrMod = devPtrMod + advanceOffset;
-    hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
-                                                    numW * sizeof(char), numH);
-    ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent, stream);
-
-    REQUIRE(ret != hipSuccess);
-  }
-
-  HIP_CHECK(hipStreamDestroy(stream));
-  HIP_CHECK(hipFree(devPitchedPtr.ptr));
-}
diff --git a/projects/hip/tests/catch/unit/memory/hipMemsetInvalidPtr.cc b/projects/hip/tests/catch/unit/memory/hipMemsetInvalidPtr.cc
deleted file mode 100644
index 8c87c66ea8..0000000000
--- a/projects/hip/tests/catch/unit/memory/hipMemsetInvalidPtr.cc
+++ /dev/null
@@ -1,127 +0,0 @@
-/*
- * 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);
-  }
-}
diff --git a/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc b/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc
new file mode 100644
index 0000000000..a06638904e
--- /dev/null
+++ b/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc
@@ -0,0 +1,229 @@
+/*
+ * Copyright (c) 2022 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.
+ */
+
+#include <hip_test_common.hh>
+#include <limits>
+#include "hip/driver_types.h"
+
+static constexpr size_t memsetVal{0x42};
+static constexpr hipExtent validExtent{184, 57, 16};
+static constexpr size_t height{validExtent.height};
+static constexpr size_t width{validExtent.width};
+static constexpr hipStream_t nullStream{nullptr};
+
+inline void testHipMemsetApis(void* dst, int value, size_t sizeBytes,
+                              hipError_t expectedReturn = hipErrorInvalidValue) {
+  HIP_CHECK_ERROR(hipMemset(dst, value, sizeBytes), expectedReturn);
+  HIP_CHECK_ERROR(hipMemsetAsync(dst, value, sizeBytes, nullStream), expectedReturn);
+
+  hipDeviceptr_t devicePtrDst = reinterpret_cast<hipDeviceptr_t>(dst);
+  HIP_CHECK_ERROR(hipMemsetD32(devicePtrDst, value, sizeBytes), expectedReturn);
+  HIP_CHECK_ERROR(hipMemsetD32Async(devicePtrDst, value, sizeBytes, nullStream), expectedReturn);
+  HIP_CHECK_ERROR(hipMemsetD16(devicePtrDst, value, sizeBytes), expectedReturn);
+  HIP_CHECK_ERROR(hipMemsetD16Async(devicePtrDst, value, sizeBytes, nullStream), expectedReturn);
+  HIP_CHECK_ERROR(hipMemsetD8(devicePtrDst, value, sizeBytes), expectedReturn);
+  HIP_CHECK_ERROR(hipMemsetD8Async(devicePtrDst, value, sizeBytes, nullStream), expectedReturn);
+}
+
+inline void testHipMemset2DApis(void* dst, size_t pitch, int value, size_t width, size_t height,
+                                hipError_t expectedReturn = hipErrorInvalidValue) {
+  HIP_CHECK_ERROR(hipMemset2D(dst, pitch, value, width, height), expectedReturn);
+  HIP_CHECK_ERROR(hipMemset2DAsync(dst, pitch, value, width, height, nullStream), expectedReturn);
+}
+
+
+inline void testHipMemset3DApis(hipPitchedPtr& pitchedDevPtr, int value, const hipExtent& extent,
+                                hipError_t expectedReturn = hipErrorInvalidValue) {
+  HIP_CHECK_ERROR(hipMemset3D(pitchedDevPtr, value, extent), expectedReturn);
+  HIP_CHECK_ERROR(hipMemset3DAsync(pitchedDevPtr, value, extent, nullStream), expectedReturn);
+}
+
+TEST_CASE("Unit_hipMemset_Negative_InvalidPtr") {
+  void* dst;
+
+  SECTION("Uninitialized Dst") {}
+  SECTION("Nullptr as dst") { dst = nullptr; }
+
+  std::unique_ptr<char[]> hostPtr;
+  SECTION("Host Pointer as Dst") {
+    hostPtr.reset(new char[width]);
+    dst = hostPtr.get();
+  }
+
+  testHipMemsetApis(dst, memsetVal, width);
+}
+
+
+TEST_CASE("Unit_hipMemset_Negative_OutOfBoundsSize") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-20");
+#endif
+
+#if !HT_AMD
+  void* dst;
+  constexpr size_t outOfBoundsSize{width + 1};
+  HIP_CHECK(hipMalloc(&dst, width));
+
+  testHipMemsetApis(dst, memsetVal, outOfBoundsSize);
+  HIP_CHECK(hipFree(dst));
+#endif
+}
+
+TEST_CASE("Unit_hipMemset_Negative_OutOfBoundsPtr") {
+  void* dst;
+  HIP_CHECK(hipMalloc(&dst, width));
+  void* outOfBoundsPtr{reinterpret_cast<char*>(dst) + width + 1};
+  testHipMemsetApis(outOfBoundsPtr, memsetVal, width);
+  HIP_CHECK(hipFree(dst));
+}
+
+TEST_CASE("Unit_hipMemset2D_Negative_InvalidPtr") {
+  void* dst;
+  SECTION("Uninitialized Dst") {}
+  SECTION("Nullptr as Dst") { dst = nullptr; }
+
+  std::unique_ptr<char[]> hostPtr;
+  SECTION("Host Pointer as Dst") {
+    hostPtr.reset(new char[height * width]);
+    dst = hostPtr.get();
+  }
+
+  void* A_d;
+  size_t pitch_A;
+  HIP_CHECK(hipMallocPitch(&A_d, &pitch_A, width, height));
+  testHipMemset2DApis(dst, pitch_A, memsetVal, width, height);
+  hipFree(A_d);
+}
+
+TEST_CASE("Unit_hipMemset2D_Negative_InvalidSizes") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-52");
+#endif
+
+  void* dst;
+  size_t realPitch;
+  HIP_CHECK(hipMallocPitch(&dst, &realPitch, width, height));
+
+  SECTION("Invalid Pitch") {
+    size_t invalidPitch = 1;
+    testHipMemset2DApis(dst, invalidPitch, memsetVal, width, height);
+  }
+
+  SECTION("Invalid Width") {
+    size_t invalidWidth = realPitch + 1;
+    testHipMemset2DApis(dst, realPitch, memsetVal, invalidWidth, height);
+  }
+
+#if !HT_AMD /* EXSWCPHIPT-52 */
+  SECTION("Invalid height") {
+    size_t invalidHeight = height + 1;
+    testHipMemset2DApis(dst, realPitch, memsetVal, width, invalidHeight);
+  }
+#endif
+  HIP_CHECK(hipFree(dst));
+}
+
+TEST_CASE("Unit_hipMemset2D_Negative_OutOfBoundsPtr") {
+  void* dst;
+  size_t realPitch;
+
+  HIP_CHECK(hipMallocPitch(&dst, &realPitch, width, height));
+  void* outOfBoundsPtr{reinterpret_cast<char*>(dst) + realPitch * height + 1};
+  testHipMemset2DApis(outOfBoundsPtr, realPitch, memsetVal, width, height);
+  HIP_CHECK(hipFree(dst));
+}
+
+
+TEST_CASE("Unit_hipMemset3D_Negative_InvalidPtr") {
+  hipPitchedPtr pitchedDevPtr;
+
+  SECTION("Uninitialized PitchedDevPtr") {}
+  SECTION("Zero Initialized PitchedDevPtr") { pitchedDevPtr = {}; }
+
+  testHipMemset3DApis(pitchedDevPtr, memsetVal, validExtent);
+}
+
+TEST_CASE("Unit_hipMemset3D_Negative_ModifiedPtr") {
+  hipPitchedPtr pitchedDevPtr;
+
+  HIP_CHECK(hipMalloc3D(&pitchedDevPtr, validExtent));
+  void* allocatedMemory{pitchedDevPtr.ptr};
+
+  SECTION("Nullptr Dst") { pitchedDevPtr.ptr = nullptr; }
+
+  std::unique_ptr<char[]> hostPtr;
+  SECTION("Host Pointer as Dst") {
+    hostPtr.reset(new char[validExtent.width * validExtent.height * validExtent.depth]);
+    pitchedDevPtr.ptr = hostPtr.get();
+  }
+
+  SECTION("Invalid Pitch") { pitchedDevPtr.pitch = 1; }
+
+  CAPTURE(pitchedDevPtr.ptr, pitchedDevPtr.pitch, pitchedDevPtr.xsize, pitchedDevPtr.ysize);
+  testHipMemset3DApis(pitchedDevPtr, memsetVal, validExtent);
+  HIP_CHECK(hipFree(allocatedMemory));
+}
+
+TEST_CASE("Unit_hipMemset3D_Negative_InvalidSizes") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-52");
+#endif
+
+  hipPitchedPtr pitchedDevPtr;
+  HIP_CHECK(hipMalloc3D(&pitchedDevPtr, validExtent));
+  hipExtent invalidExtent{validExtent};
+
+  SECTION("Max Width") { invalidExtent.width = std::numeric_limits<std::size_t>::max(); }
+
+  SECTION("Max Height") { invalidExtent.height = std::numeric_limits<std::size_t>::max(); }
+
+#if !HT_NVIDIA /* This case hangs on Nvidia */
+  SECTION("Max Depth") { invalidExtent.depth = std::numeric_limits<std::size_t>::max(); }
+#endif
+
+  SECTION("Invalid Width") { invalidExtent.width = pitchedDevPtr.pitch + 1; }
+
+#if !HT_AMD /* EXSWCPHIPT-52 */
+  SECTION("Invalid height") { invalidExtent.height += 1; }
+
+  SECTION("Invalid depth") { invalidExtent.depth += 1; }
+#endif
+
+  CAPTURE(invalidExtent.width, invalidExtent.height, invalidExtent.depth);
+  testHipMemset3DApis(pitchedDevPtr, memsetVal, invalidExtent);
+  HIP_CHECK(hipFree(pitchedDevPtr.ptr));
+}
+
+TEST_CASE("Unit_hipMemset3D_Negative_OutOfBounds") {
+  hipPitchedPtr pitchedDevPtr;
+
+  HIP_CHECK(hipMalloc3D(&pitchedDevPtr, validExtent));
+  hipPitchedPtr outOfBoundsPtr{pitchedDevPtr};
+  outOfBoundsPtr.ptr = reinterpret_cast<char*>(pitchedDevPtr.ptr) +
+      pitchedDevPtr.pitch * validExtent.height * validExtent.depth + 1;
+
+  SECTION("Extent Equal to 0") {
+    hipExtent zeroExtent{0, 0, 0};
+    testHipMemset3DApis(outOfBoundsPtr, memsetVal, zeroExtent, hipSuccess);
+  }
+
+  SECTION("Valid Extent") { testHipMemset3DApis(outOfBoundsPtr, memsetVal, validExtent); }
+
+  HIP_CHECK(hipFree(pitchedDevPtr.ptr));
+}
diff --git a/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc b/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc
index 4db534c223..1e02b696fc 100644
--- a/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc
+++ b/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc
@@ -336,3 +336,72 @@ TEST_CASE("Unit_hipPointerGetAttributes_MultiThread") {
     std::thread t4(clusterAllocs, 1000, 1, 4);
     if (serialize) t4.join();
 }
+
+TEST_CASE("Unit_hipPointerGetAttributes_Negative") {
+  SECTION("Invalid Attributes Pointer") {
+    int* dPtr{nullptr};
+    HIP_CHECK(hipMalloc(&dPtr, sizeof(int)));
+    HIP_CHECK_ERROR(hipPointerGetAttributes(nullptr, dPtr), hipErrorInvalidValue);
+    HIP_CHECK(hipFree(dPtr));
+  }
+
+  SECTION("Invalid Device Pointer") {
+    hipPointerAttribute_t attributes{};
+    HIP_CHECK_ERROR(hipPointerGetAttributes(&attributes, nullptr), hipErrorInvalidValue);
+  }
+}
+
+// Run this test for all devices for DeviceMemory, HostMemory, ManagedMemory and MappedMemory
+TEST_CASE("Unit_hipPointerGetAttributes_GpuIter") {
+  int deviceCount{0};
+  HIP_CHECK(hipGetDeviceCount(&deviceCount));
+  REQUIRE(deviceCount != 0);
+
+  // Memory Types
+  enum MemoryTypes { DeviceMemory = 0, HostMemory, MappedMemory };
+  auto MemoryType =
+      GENERATE(MemoryTypes::DeviceMemory, MemoryTypes::HostMemory, MemoryTypes::MappedMemory);
+
+  INFO("Memory Type: " << MemoryType);
+
+  int* ptr{nullptr};
+  for (int i = 0; i < deviceCount; i++) {
+    HIP_CHECK(hipSetDevice(i));
+
+    ptr = nullptr;
+    if (MemoryType == MemoryTypes::DeviceMemory) {
+      HIP_CHECK(hipMalloc(&ptr, sizeof(int)));
+    } else if (MemoryType == MemoryTypes::HostMemory) {
+      HIP_CHECK(hipHostMalloc(&ptr, sizeof(int)));
+    } else if (MemoryType == MemoryTypes::MappedMemory) {
+      HIP_CHECK(hipHostMalloc(&ptr, sizeof(int), hipHostMallocMapped));
+    }
+
+    REQUIRE(ptr != nullptr);
+    hipPointerAttribute_t attributes{};
+    HIP_CHECK(hipPointerGetAttributes(&attributes, ptr));
+
+    REQUIRE(attributes.device == i);  // Device Check
+
+    // Memory address and type check
+    if (MemoryType == MemoryTypes::DeviceMemory) {
+      REQUIRE(attributes.memoryType == hipMemoryTypeDevice);
+      REQUIRE(attributes.devicePointer == ptr);
+      REQUIRE(attributes.hostPointer == nullptr);
+    } else if (MemoryType == MemoryTypes::HostMemory) {
+      REQUIRE(attributes.memoryType == hipMemoryTypeHost);
+      REQUIRE(attributes.hostPointer == ptr);
+    } else if (MemoryType == MemoryTypes::MappedMemory) {
+      int* devicePtr{nullptr};
+      HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&devicePtr), ptr, 0));
+      REQUIRE(attributes.hostPointer == ptr);
+      REQUIRE(attributes.hostPointer == devicePtr);
+    }
+
+    if (MemoryType == MemoryTypes::DeviceMemory) {
+      HIP_CHECK(hipFree(ptr));
+    } else if (MemoryType == MemoryTypes::MappedMemory || MemoryType == MemoryTypes::HostMemory) {
+      HIP_CHECK(hipHostFree(ptr));
+    }
+  }
+}
diff --git a/projects/hip/tests/catch/unit/stream/hipStreamGetPriority.cc b/projects/hip/tests/catch/unit/stream/hipStreamGetPriority.cc
index 72839db20d..ac68a9396a 100644
--- a/projects/hip/tests/catch/unit/stream/hipStreamGetPriority.cc
+++ b/projects/hip/tests/catch/unit/stream/hipStreamGetPriority.cc
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2022 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
@@ -16,59 +16,77 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-/**
+
+/*
 Testcase Scenarios :
 1) Negative tests for hipStreamGetPriority api.
 2) Create stream and check default priority of stream is within range.
 3) Create stream with high or low priority and check priority is set as expected.
-4) Create stream with higher priority or lower priority for the priority range returned.
+4) Create stream with higher priority or lower priority for the priority range returned, the stream
+priority should be clamped to the priority range.
 5) Create stream with CUMask and check priority is returned as expected.
 */
 
 #include <hip_test_common.hh>
 
 /**
- * Negative tests for hipStreamGetPriority api.
+ * Check the error returned when an invalid pointer to a priority is used.
  */
-TEST_CASE("Unit_hipStreamGetPriority_Negative") {
-  hipStream_t stream = 0;
-  REQUIRE(hipStreamGetPriority(stream, nullptr) == hipErrorInvalidValue);
+TEST_CASE("Unit_hipStreamGetPriority_InvalidPriorityPointer") {
+  hipStream_t stream{};
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK_ERROR(hipStreamGetPriority(stream, nullptr), hipErrorInvalidValue);
+  HIP_CHECK(hipStreamDestroy(stream));
 }
 
 /**
- * Create stream with higher priority for the priority range returned.
+ * Create stream and check priority.
  */
-TEST_CASE("Unit_hipStreamGetPriority_higher") {
+TEST_CASE("Unit_hipStreamGetPriority_happy") {
   int priority_low = 0;
   int priority_high = 0;
   int devID = GENERATE(range(0, HipTest::getDeviceCount()));
   HIP_CHECK(hipSetDevice(devID));
   HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
-  hipStream_t stream;
-  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
-                                                           priority_high-1));
+  hipStream_t stream{};
   int priority = 0;
-  HIP_CHECK(hipStreamGetPriority(stream, &priority));
-  REQUIRE(priority == priority_high);
-  HIP_CHECK(hipStreamDestroy(stream));
-}
-
-/**
- * Create stream with lower priority for the priority range returned.
- */
-TEST_CASE("Unit_hipStreamGetPriority_lower") {
-  int priority_low = 0;
-  int priority_high = 0;
-  int devID = GENERATE(range(0, HipTest::getDeviceCount()));
-  HIP_CHECK(hipSetDevice(devID));
-  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
-  hipStream_t stream;
-  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
-                                                           priority_low+1));
-  int priority = 0;
-  HIP_CHECK(hipStreamGetPriority(stream, &priority));
-  REQUIRE(priority_low == priority);
-  HIP_CHECK(hipStreamDestroy(stream));
+  SECTION("Null Stream") {
+    HIP_CHECK(hipStreamGetPriority(nullptr, &priority));
+    // valid priority
+    REQUIRE(priority_low >= priority);
+    REQUIRE(priority >= priority_high);
+  }
+  SECTION("Created Stream") {
+    SECTION("Default Priority") {
+      HIP_CHECK(hipStreamCreate(&stream));
+      HIP_CHECK(hipStreamGetPriority(stream, &priority));
+      // valid priority
+      // Lower the value higher the priority, higher the value lower the priority
+      REQUIRE(priority_low >= priority);
+      REQUIRE(priority >= priority_high);
+    }
+    SECTION("High Priority") {
+      HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
+      HIP_CHECK(hipStreamGetPriority(stream, &priority));
+      REQUIRE(priority == priority_high);
+    }
+    SECTION("Higher Priority") {
+      HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high - 1));
+      HIP_CHECK(hipStreamGetPriority(stream, &priority));
+      REQUIRE(priority == priority_high);
+    }
+    SECTION("Low Priority") {
+      HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
+      HIP_CHECK(hipStreamGetPriority(stream, &priority));
+      REQUIRE(priority_low == priority);
+    }
+    SECTION("Lower Priority") {
+      HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_low + 1));
+      HIP_CHECK(hipStreamGetPriority(stream, &priority));
+      REQUIRE(priority_low == priority);
+    }
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
 }
 
 #if HT_AMD
@@ -76,7 +94,7 @@ TEST_CASE("Unit_hipStreamGetPriority_lower") {
  * Create stream with CUMask and check priority is returned as expected.
  */
 TEST_CASE("Unit_hipStreamGetPriority_StreamsWithCUMask") {
-  hipStream_t stream;
+  hipStream_t stream{};
   int priority = 0;
   int priority_normal = 0;
   int priority_low = 0;