SWDEV-299127 - Merge 'develop' into 'amd-staging'

Change-Id: Ibe710eba031ddce646fdae4593615174ee717a29

.jenkins/Jenkinsfile

@@ -58,8 +58,8 @@ def hipBuildTest(String backendLabel) {
                         set -x
                         # Check if backend label contains string "amd" or backend host is a server with amd gpu
                         if [[ $backendLabel =~ amd ]]; then
-                            LLVM_PATH=/opt/rocm/llvm ctest -E 'cooperative_streams_least_capacity.tst|cooperative_streams_half_capacity.tst|cooperative_streams_full_capacity.tst|grid_group_data_sharing.tst|hipIpcMemAccessTest.tst|p2p_copy_coherency.tst'
                             sleep 120
+                            LLVM_PATH=/opt/rocm/llvm ctest -E 'cooperative_streams_least_capacity.tst|cooperative_streams_half_capacity.tst|cooperative_streams_full_capacity.tst|grid_group_data_sharing.tst|hipIpcMemAccessTest.tst|p2p_copy_coherency.tst'
                         else
                             make test
                         fi
@@ -100,9 +100,9 @@ def hipBuildTest(String backendLabel) {
                         set -x
                         # Check if backend label contains string "amd" or backend host is a server with amd gpu
                         if [[ $backendLabel =~ amd ]]; then
+                            sleep 120
                             export HT_CONFIG_FILE="$HIP_DIR/tests/catch/hipTestMain/config/config_amd_linux.json"
                             LLVM_PATH=/opt/rocm/llvm ctest -E 'Unit_hipGraphChildGraphNodeGetGraph_Functional|Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative|Unit_hipPtrGetAttribute_Simple|Unit_hipStreamPerThread_DeviceReset_2'
-                            sleep 120
                         else
                             make test
                         fi

docs/markdown/hip_build.md

@@ -141,7 +141,7 @@ After build and install HIP commands, catch tests can be built via the following
 ```
 cd "$HIP_DIR"
 mkdir -p build; cd build
-export HIP_PATH=$HIPAMD_DIR/build
+export HIP_PATH=$HIPAMD_DIR/build/install
 cmake  ../tests/catch/ -DHIP_PLATFORM=amd
 make -j$(nproc) build_tests
 ctest # run tests

docs/markdown/hip_deprecated_api_list.md

@@ -1,9 +1,11 @@
 # HIP Deprecated APIs
 
+
 ## HIP Context Management APIs
 
 CUDA supports cuCtx API, the Driver API that defines "Context" and "Devices" as separate entities. Contexts contain a single device, and a device can theoretically have multiple contexts. HIP initially added limited support for these API to facilitate easy porting from existing driver codes. These API are marked as deprecated now since there are better alternate interface (such as hipSetDevice or the stream API) to achieve the required functions.
-
+### hipCtxCreate
+### hipCtxDestroy
 ### hipCtxPopCurrent
 ### hipCtxPushCurrent
 ### hipCtxSetCurrent
@@ -19,6 +21,7 @@ CUDA supports cuCtx API, the Driver API that defines "Context" and "Devices" as
 ### hipCtxEnablePeerAccess
 ### hipCtxDisablePeerAccess
 
+
 ## HIP Memory Management APIs
 
 ### hipMallocHost
@@ -31,4 +34,49 @@ Should use "hipHostMalloc" instead.
 Should use "hipHostMalloc" instead.
 
 ### hipFreeHost
-Should use "hipHostFree" instead.
+Should use "hipHostFree" instead.
+
+### hipMemcpyToArray
+### hipMemcpyFromArray
+
+
+## HIP Profiler Control APIs
+
+### hipProfilerStart
+Should use roctracer/rocTX instead
+
+### hipProfilerStop
+Should use roctracer/rocTX instead
+
+
+## HIP Texture Management APIs
+
+###hipGetTextureReference
+###hipTexRefSetAddressMode
+###hipTexRefSetArray
+###hipTexRefSetFilterMode
+###hipTexRefSetFlags
+###hipTexRefSetFormat
+###hipBindTexture
+###hipBindTexture2D
+###hipBindTextureToArray
+###hipGetTextureAlignmentOffset
+###hipUnbindTexture
+###hipTexRefGetAddress
+###hipTexRefGetAddressMode
+###hipTexRefGetFilterMode
+###hipTexRefGetFlags
+###hipTexRefGetFormat
+###hipTexRefGetMaxAnisotropy
+###hipTexRefGetMipmapFilterMode
+###hipTexRefGetMipmapLevelBias
+###hipTexRefGetMipmapLevelClamp
+###hipTexRefGetMipMappedArray
+###hipTexRefSetAddress
+###hipTexRefSetAddress2D
+###hipTexRefSetMaxAnisotropy
+###hipTexRefSetBorderColor
+###hipTexRefSetMipmapFilterMode
+###hipTexRefSetMipmapLevelBias
+###hipTexRefSetMipmapLevelClamp
+###hipTexRefSetMipmappedArray

docs/markdown/hip_programming_guide.md

@@ -129,6 +129,9 @@ For more details on hipRTC APIs, refer to HIP-API.pdf in GitHub (https://github.
 
 The link here(https://github.com/ROCm-Developer-Tools/HIP/blob/main/tests/src/hiprtc/saxpy.cpp) shows an example how to program HIP application using runtime compilation mechanism, and detail hipRTC programming guide is also available in Github (https://github.com/ROCm-Developer-Tools/HIP/blob/main/docs/markdown/hip_rtc.md).
 
+## HIP Graph
+HIP graph is supported. For more details, refer to the HIP API Guide.
+
 ## Device-Side Malloc
 
 HIP-Clang now supports device-side malloc and free.

include/hip/hip_runtime_api.h

@@ -657,7 +657,7 @@ enum hipLimit_t {
 #define hipDeviceScheduleBlockingSync 0x4
 #define hipDeviceScheduleMask 0x7
 #define hipDeviceMapHost 0x8
-#define hipDeviceLmemResizeToMax 0x16
+#define hipDeviceLmemResizeToMax 0x10
 /** Default HIP array allocation flag.*/
 #define hipArrayDefault 0x00
 #define hipArrayLayered 0x01

samples/0_Intro/square/square.hipref.cpp

@@ -38,7 +38,7 @@ THE SOFTWARE.
  */
 template <typename T>
 __global__ void vector_square(T* C_d, const T* A_d, size_t N) {
-    size_t offset = (blockIdx.x * blockDim_x + threadIdx.x);
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
     size_t stride = blockDim.x * gridDim.x;
 
     for (size_t i = offset; i < N; i += stride) {

tests/catch/CMakeLists.txt

@@ -41,6 +41,7 @@ if(NOT DEFINED ROCM_PATH)
         endif()
     endif()
 endif()
+file(TO_CMAKE_PATH "${ROCM_PATH}" ROCM_PATH)
 message(STATUS "ROCM Path: ${ROCM_PATH}")
 
 
@@ -179,6 +180,8 @@ if(NOT DEFINED OFFLOAD_ARCH_STR
     else()
         message(STATUS "ROCm Agent Enumurator found no valid architectures")
     endif()
+elseif(DEFINED OFFLOAD_ARCH_STR)
+    string(REPLACE "--offload-arch=" "" HIP_GPU_ARCH_LIST ${OFFLOAD_ARCH_STR})
 endif()
 
 if(DEFINED OFFLOAD_ARCH_STR)

tests/catch/README.md

@@ -144,7 +144,7 @@ If there arises a condition where certain flag is disabled and due to which a te
   ```cpp
   TEST_CASE("TestOnlyOnXnack") {
     if(!XNACKEnabled) {
-      HIP_SKIP_TEST("Test only runs on system with XNACK enabled");
+      HipTest::HIP_SKIP_TEST("Test only runs on system with XNACK enabled");
       return;
     }
     // Rest of test functionality
@@ -162,11 +162,11 @@ These macros are to be called in multi process tests, inside a process which get
 There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn a process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf output.
 Sample Usage:
 ```cpp
-hip::SpawnProc proc(<relative path of exe with test folder>, <optional bool value, if output is to be recorded>);
+hip::SpawnProc proc(<name of exe>, <optional bool value, if output is to be recorded>);
 REQUIRE(0 == proc.run()); // Test of return value of the proc
 REQUIRE(exepctedOutput == proc.getOutput()); // Test on expected output of the process
 ```
-The process can be a standalone exe (see tests/catch/unit/printfExe for more information).
+The process must be a standalone exe inside the same folder as other tests.
 
 ## Enabling New Tests
 Initially, the new tests can be enabled via using ```-DHIP_CATCH_TEST=1```. After porting existing tests, this will be turned on by default.

tests/catch/hipTestMain/config/config_amd_linux.json

@@ -3,7 +3,8 @@
 	[
 	   "# Following test is related to ticket EXSWCPHIPT-41",
 	   "Unit_hipStreamGetPriority_happy",
-	   "Unit_hipStreamPerThread_DeviceReset_1"
+	   "Unit_hipStreamPerThread_DeviceReset_1",
+	   "Unit_hipGraphAddKernelNode_Negative"
 	]
 
 }

tests/catch/hipTestMain/config/config_amd_windows.json

@@ -29,6 +29,23 @@
 	   "Unit_hipDeviceGetUuid",
 	   "Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Functional",
 	   "Unit_hipGraphExecEventWaitNodeSetEvent_Negative",
-	   "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory"
+	   "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory",
+	   "Unit_hipGraphMemcpyNodeSetParams_Functional",
+	   "Unit_hipMalloc3D_ValidatePitch",
+	   "Unit_hipArrayCreate_happy",
+	   "Unit_hipGraphAddKernelNode_Negative",
+	   "Unit_hipHostRegister_Negative - int",
+	   "Unit_hipHostRegister_Negative - float",
+	   "Unit_hipHostRegister_Negative - double",
+	   "Unit_hipMemAllocPitch_ValidatePitch",
+	   "Unit_hipArrayCreate_happy - int",
+	   "Unit_hipArrayCreate_happy - int4",
+	   "Unit_hipArrayCreate_happy - short2",
+	   "Unit_hipArrayCreate_happy - char",
+	   "Unit_hipArrayCreate_happy - char4",
+	   "Unit_hipArrayCreate_happy - float",
+	   "Unit_hipArrayCreate_happy - float2",
+	   "Unit_hipArrayCreate_happy - float4",
+	   "Unit_hipMemVmm_Basic"
 	]
 }

tests/catch/include/hip_test_helper.hh

@@ -38,7 +38,7 @@ static inline int getGeviceCount() {
 }
 
 // Get Free Memory from the system
-static size_t getMemoryAmount() {
+static inline size_t getMemoryAmount() {
 #ifdef __linux__
   struct sysinfo info{};
   sysinfo(&info);

tests/catch/include/hip_test_process.hh

@@ -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
@@ -21,7 +21,9 @@ THE SOFTWARE.
 */
 
 #pragma once
+
 #include "hip_test_common.hh"
+#include "hip_test_filesystem.hh"
 
 #include <string>
 #include <array>
@@ -29,9 +31,17 @@ THE SOFTWARE.
 #include <random>
 #include <fstream>
 #include <streambuf>
-#include "hip_test_filesystem.hh"
 
 namespace hip {
+/*
+Class to spawn a process in isolation and test its standard output and return status
+Good for printf tests and environment variable tests
+
+How to use:
+Have the stand alone exe in the same folder
+Init a class using hip::SpawnProc proc("ExeName", yes_or_no_to_capture_output);
+proc.run("Optional command line args");
+*/
 class SpawnProc {
   std::string exeName;
   std::string resultStr;
@@ -53,18 +63,31 @@ class SpawnProc {
  public:
   SpawnProc(std::string exeName_, bool captureOutput_ = false)
       : exeName(exeName_), captureOutput(captureOutput_) {
-    auto dir = fs::path(TestContext::get().currentPath()).parent_path();
+    auto dir = fs::path(TestContext::get().currentPath());
     dir /= exeName;
     exeName = dir.string();
+
+    INFO("Testing that exe exists: " << exeName);
+    REQUIRE(fs::exists(exeName));
+
     if (captureOutput) {
       auto path = fs::temp_directory_path();
       path /= getRandomString();
       tmpFileName = path.string();
+      INFO("Testing that capture file does not exist already: " << tmpFileName);
+      REQUIRE(!fs::exists(tmpFileName));
     }
   }
 
-  int run() {
+  int run(std::string commandLineArgs = "") {
     std::string execCmd = exeName;
+
+    // Append command line args
+    if (commandLineArgs.size() > 0) {
+      execCmd += " ";  // Add space for command line args
+      execCmd += commandLineArgs;
+    }
+
     if (captureOutput) {
       execCmd += " > ";
       execCmd += tmpFileName;
@@ -77,7 +100,11 @@ class SpawnProc {
       resultStr =
           std::string((std::istreambuf_iterator<char>(t)), std::istreambuf_iterator<char>());
     }
+#if HT_LINUX
+    return WEXITSTATUS(res);
+#else
     return res;
+#endif
   }
 
   std::string getOutput() { return resultStr; }

tests/catch/packaging/hip-tests.txt

@@ -30,7 +30,7 @@ install(FILES @PROJECT_BINARY_DIR@/CTestTestfile.cmake
 # Packaging steps
 #############################
 set(CPACK_SET_DESTDIR TRUE)
-set(CPACK_INSTALL_PREFIX @ROCM_PATH@/test/hip/)
+set(CPACK_INSTALL_PREFIX "@ROCM_PATH@/test/hip/")
 set(PKG_NAME catch-@HIP_PLATFORM@)
 set(CPACK_PACKAGE_NAME ${PKG_NAME})
 set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [CATCH TESTS]")

tests/catch/unit/CMakeLists.txt

@@ -27,7 +27,6 @@ add_subdirectory(event)
 add_subdirectory(occupancy)
 add_subdirectory(device)
 add_subdirectory(printf)
-add_subdirectory(printfExe)
 add_subdirectory(texture)
 add_subdirectory(streamperthread)
 add_subdirectory(kernel)

tests/catch/unit/device/CMakeLists.txt

@@ -18,7 +18,13 @@ set(TEST_SRC
     hipDeviceGetUuid.cc
 )
 
+set_source_files_properties(hipGetDeviceCount.cc PROPERTIES COMPILE_FLAGS -std=c++17)
+
+add_executable(getDeviceCount EXCLUDE_FROM_ALL getDeviceCount_exe.cc)
+
 hip_add_exe_to_target(NAME DeviceTest
                       TEST_SRC ${TEST_SRC}
                       TEST_TARGET_NAME build_tests
                       COMPILE_OPTIONS -std=c++14)
+
+add_dependencies(DeviceTest getDeviceCount)

tests/catch/unit/device/getDeviceCount_exe.cc

@@ -0,0 +1,89 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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.
+*/
+
+#include <hip/hip_runtime.h>
+#include <iostream>
+#include <stdlib.h>
+
+bool UNSETENV(std::string var) {
+  int result = -1;
+  #ifdef __unix__
+    result = unsetenv(var.c_str());
+  #else
+    result = _putenv((var + '=').c_str());
+  #endif
+  return (result == 0) ? true: false;
+}
+
+bool SETENV(std::string var, std::string value, int overwrite) {
+  int result = -1;
+  #ifdef __unix__
+    result = setenv(var.c_str(), value.c_str(), overwrite);
+  #else
+    result = _putenv((var + '=' + value).c_str());
+  #endif
+  return (result == 0) ? true: false;
+}
+
+// Expects 1 command line arg, which is the Device Visible String
+int main(int argc, char** argv) {
+  if (argc != 2) {
+    std::cerr << "Invalid number of args passed.\n"
+              << "argc : " << argc << std::endl;
+    for (int i = 0; i < argc; i++) {
+      std::cerr << "  argv[" << i << "] : " << argv[0] << std::endl;
+    }
+    std::cerr << "The program expects device visibility string i.e. 0,1,2" << std::endl;
+    return -1;
+  }
+
+  // disable visible_devices env from shell
+#ifdef __HIP_PLATFORM_NVCC__
+  UNSETENV("CUDA_VISIBLE_DEVICES");
+  SETENV("CUDA_VISIBLE_DEVICES", argv[1], 1);
+  auto init_res = hipInit(0);
+  if (hipSuccess != init_res) {
+    std::cerr << "CUDA INIT API returned : " << hipGetErrorString(init_res) << std::endl;
+    return -1;
+  }
+#else
+  UNSETENV("ROCR_VISIBLE_DEVICES");
+  UNSETENV("HIP_VISIBLE_DEVICES");
+  SETENV("ROCR_VISIBLE_DEVICES", argv[1], 1);
+  SETENV("HIP_VISIBLE_DEVICES", argv[1], 1);
+#endif
+
+  int count = 0;
+  auto res = hipGetDeviceCount(&count);
+  if (hipSuccess != res) {
+    std::cerr << "HIP API returned : " << hipGetErrorString(res) << std::endl;
+    return -1;
+  }
+
+#ifdef __HIP_PLATFORM_NVCC__
+  UNSETENV("CUDA_VISIBLE_DEVICES");
+#else
+  UNSETENV("ROCR_VISIBLE_DEVICES");
+  UNSETENV("HIP_VISIBLE_DEVICES");
+#endif
+  return count;
+}

tests/catch/unit/device/hipGetDeviceCount.cc

@@ -23,6 +23,7 @@ THE SOFTWARE.
  */
 
 #include <hip_test_common.hh>
+#include <hip_test_process.hh>
 
 /**
  * hipGetDeviceCount tests
@@ -32,3 +33,25 @@ TEST_CASE("Unit_hipGetDeviceCount_NegTst") {
   // Scenario1
   REQUIRE_FALSE(hipGetDeviceCount(nullptr) == hipSuccess);
 }
+
+TEST_CASE("Unit_hipGetDeviceCount_HideDevices") {
+  int deviceCount = HipTest::getDeviceCount();
+  if (deviceCount < 2) {
+    HipTest::HIP_SKIP_TEST("This test requires more than 2 GPUs. Skipping.");
+    return;
+  }
+
+  for (int i = deviceCount; i >= 1; i--) {
+    std::string visibleStr;
+    for (int j = 0; j < i; j++) {  // Generate a string which has first i devices
+      visibleStr += std::to_string(j);
+      if (j != (i - 1)) {
+        visibleStr += ",";
+      }
+    }
+
+    hip::SpawnProc proc("getDeviceCount", true);
+    INFO("Output from process : " << proc.getOutput());
+    REQUIRE(proc.run(visibleStr) == i);
+  }
+}

tests/catch/unit/event/Unit_hipEventRecord.cc

@@ -19,72 +19,102 @@ 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.
 */
+
 // Test hipEventRecord serialization behavior.
-// Through manual inspection of the reported timestamps, can determine if recording a NULL event
-// forces synchronization : set
-#include <hip_test_checkers.hh>
-#include <kernels.hh>
-#include <hip_test_context.hh>
+
 #include <hip_test_common.hh>
 
+#include <kernels.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_context.hh>
+
 TEST_CASE("Unit_hipEventRecord") {
-    size_t N = 4 * 1024 * 1024;
-    unsigned threadsPerBlock = 256;
-    int iterations = 1;
+  constexpr size_t N = 1024;
+  constexpr int iterations = 1;
 
-    unsigned blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
-    if (blocks > 1024) blocks = 1024;
-    if (blocks == 0) blocks = 1;
+  constexpr int blocks = 1024;
 
-    printf("N=%zu (A+B+C= %6.1f MB total) blocks=%u threadsPerBlock=%u iterations=%d\n", N,
-           ((double)3 * N * sizeof(float)) / 1024 / 1024, blocks, threadsPerBlock, iterations);
-    printf("iterations=%d\n", iterations);
+  constexpr size_t Nbytes = N * sizeof(float);
 
-    size_t Nbytes = N * sizeof(float);
+  float *A_h, *B_h, *C_h;
+  float *A_d, *B_d, *C_d;
+  HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N);
 
-    float *A_h, *B_h, *C_h;
-    float *A_d, *B_d, *C_d;
-    HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N);
+  enum TestType {
+    WithFlags_Default = hipEventDefault,
+    WithFlags_Blocking = hipEventBlockingSync,
+    WithFlags_DisableTiming = hipEventDisableTiming,
+#if HT_AMD
+    WithFlags_ReleaseToDevice = hipEventReleaseToDevice,
+    WithFlags_ReleaseToSystem = hipEventReleaseToSystem,
+#endif
+    WithoutFlags
+  };
 
-    hipEvent_t start, stop;
+#if HT_AMD
+  auto flags = GENERATE(WithFlags_Default, WithFlags_Blocking, WithFlags_DisableTiming,
+                        WithFlags_ReleaseToDevice, WithFlags_ReleaseToSystem, WithoutFlags);
+#endif
 
-    // NULL stream check:
+#if HT_NVIDIA
+  auto flags =
+      GENERATE(WithFlags_Default, WithFlags_Blocking, WithFlags_DisableTiming, WithoutFlags);
+#endif
+
+
+  hipEvent_t start{}, stop{};
+
+  if (flags == WithoutFlags) {
     HIP_CHECK(hipEventCreate(&start));
     HIP_CHECK(hipEventCreate(&stop));
+  } else {
+    HIP_CHECK(hipEventCreateWithFlags(&start, flags));
+    HIP_CHECK(hipEventCreateWithFlags(&stop, flags));
+  }
 
-    HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
-    HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
 
-    for (int i = 0; i < iterations; i++) {
-        //--- START TIMED REGION
-        long long hostStart = HipTest::get_time();
-        // Record the start event
-        HIP_CHECK(hipEventRecord(start, NULL));
+  for (int i = 0; i < iterations; i++) {
+    //--- START TIMED REGION
+    long long hostStart = HipTest::get_time();
+    // Record the start event
+    HIP_CHECK(hipEventRecord(start, NULL));
 
-        HipTest::launchKernel<float>(HipTest::vectorADD<float>, blocks, threadsPerBlock, 0, 0,
-static_cast<const float*>(A_d), static_cast<const float*>(B_d), C_d, N);
+    HipTest::launchKernel<float>(HipTest::vectorADD<float>, blocks, 1, 0, 0,
+                                 static_cast<const float*>(A_d), static_cast<const float*>(B_d),
+                                 C_d, N);
 
-        HIP_CHECK(hipEventRecord(stop, NULL));
-        HIP_CHECK(hipEventSynchronize(stop));
-        long long hostStop = HipTest::get_time();
-        //--- STOP TIMED REGION
+    HIP_CHECK(hipEventRecord(stop, NULL));
+    HIP_CHECK(hipEventSynchronize(stop));
+    long long hostStop = HipTest::get_time();
+    //--- STOP TIMED REGION
 
-        float eventMs = 1.0f;
-        HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop));
-        float hostMs = HipTest::elapsed_time(hostStart, hostStop);
+    float hostMs = HipTest::elapsed_time(hostStart, hostStop);
 
-        printf("host_time (chrono)                =%6.3fms\n", hostMs);
-        printf("kernel_time (hipEventElapsedTime) =%6.3fms\n", eventMs);
-        printf("\n");
+    INFO("host_time (chrono)                = " << hostMs);
 
-        // Make sure timer is timing something...
-        REQUIRE(eventMs > 0.0f);
+    // Make sure timer is timing something...
+    if (flags != WithFlags_DisableTiming) {
+      float eventMs = 1.0f;
+      HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop));
+      INFO("kernel_time (hipEventElapsedTime) = " << eventMs);
+      REQUIRE(eventMs > 0.0f);
     }
+  }
 
-    HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
 
-    HIP_CHECK(hipEventDestroy(start));
-    HIP_CHECK(hipEventDestroy(stop));
+  HIP_CHECK(hipEventDestroy(start));
+  HIP_CHECK(hipEventDestroy(stop));
 
-    HipTest::checkVectorADD(A_h, B_h, C_h, N, true);
+  HipTest::checkVectorADD(A_h, B_h, C_h, N, true);
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+  TestContext::get().cleanContext();
 }
+
+TEST_CASE("Unit_hipEventRecord_Negative") {
+  SECTION("Nullptr event") {
+    HIP_CHECK_ERROR(hipEventRecord(nullptr, nullptr), hipErrorInvalidResourceHandle);
+  }
+}

tests/catch/unit/graph/CMakeLists.txt

@@ -64,6 +64,10 @@ set(TEST_SRC
   hipStreamEndCapture.cc
   hipGraphMemcpyNodeSetParamsFromSymbol.cc
   hipGraphExecEventWaitNodeSetEvent.cc
+  hipGraphAddMemsetNode.cc
+  hipGraphAddKernelNode.cc
+  hipGraphMemcpyNodeGetParams.cc
+  hipGraphMemcpyNodeSetParams.cc
 )
 
 hip_add_exe_to_target(NAME GraphsTest

tests/catch/unit/graph/hipGraphAddKernelNode.cc

@@ -0,0 +1,101 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_kernels.hh>
+
+/* Test verifies hipGraphAddKernelNode API Negative scenarios.
+ */
+
+TEST_CASE("Unit_hipGraphAddKernelNode_Negative") {
+  constexpr int N = 1024;
+  size_t NElem{N};
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+  int *A_d, *B_d, *C_d;
+  hipGraph_t graph;
+  hipError_t ret;
+  hipGraphNode_t kNode;
+  hipKernelNodeParams kNodeParams{};
+  std::vector<hipGraphNode_t> dependencies;
+
+  HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N));
+  HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N));
+  HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kNodeParams.gridDim = dim3(blocks);
+  kNodeParams.blockDim = dim3(threadsPerBlock);
+  kNodeParams.sharedMemBytes = 0;
+  kNodeParams.kernelParams = reinterpret_cast<void **>(kernelArgs);
+  kNodeParams.extra = nullptr;
+
+  SECTION("Pass pGraphNode as nullptr") {
+    ret = hipGraphAddKernelNode(nullptr, graph, nullptr, 0, &kNodeParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass Graph as nullptr") {
+    ret = hipGraphAddKernelNode(&kNode, nullptr, nullptr, 0, &kNodeParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass invalid numDependencies") {
+    ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 11, &kNodeParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass invalid numDependencies and valid list for dependencies") {
+    HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
+    dependencies.push_back(kNode);
+    ret = hipGraphAddKernelNode(&kNode, graph,
+              dependencies.data(), dependencies.size()+1, &kNodeParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass NodeParams as nullptr") {
+    ret = hipGraphAddKernelNode(&kNode, graph,
+              dependencies.data(), dependencies.size(), nullptr);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass NodeParams func datamember as nullptr") {
+    kNodeParams.func = nullptr;
+    ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams);
+    REQUIRE(hipSuccess != ret);
+  }
+  SECTION("Pass kernelParams datamember as nullptr") {
+    kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+    kNodeParams.kernelParams = nullptr;
+    ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+#if HT_AMD
+// On Cuda setup this test case getting failed
+  SECTION("Try adding kernel node after destroy the already created graph") {
+    kNodeParams.kernelParams = reinterpret_cast<void **>(kernelArgs);
+    HIP_CHECK(hipGraphDestroy(graph));
+    ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+#endif
+
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipFree(B_d));
+  HIP_CHECK(hipFree(C_d));
+}
+

tests/catch/unit/graph/hipGraphAddMemsetNode.cc

@@ -0,0 +1,123 @@
+/*
+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.
+*/
+
+/**
+Negative Testcase Scenarios for api hipGraphAddMemsetNode :
+1) Pass pGraphNode as nullptr and check if api returns error.
+2) Pass pGraphNode as un-initialize object and check.
+3) Pass Graph as nullptr and check if api returns error.
+4) Pass Graph as empty object(skipping graph creation), api should return error code.
+5) Pass pDependencies as nullptr, api should return success.
+6) Pass numDependencies is max(size_t) and pDependencies is not valid ptr, api expected to return error code.
+7) Pass pDependencies is nullptr, but numDependencies is non-zero, api expected to return error.
+8) Pass pMemsetParams as nullptr and check if api returns error code.
+9) Pass pMemsetParams as un-initialize object and check if api returns error code.
+10) Pass hipMemsetParams::dst as nullptr should return error code.
+11) Pass hipMemsetParams::element size other than 1, 2, or 4 and check api should return error code.
+12) Pass hipMemsetParams::height as zero and check api should return error code.
+*/
+
+#include <hip_test_common.hh>
+
+/**
+ * Negative Test for API hipGraphAddMemsetNode
+ */
+
+TEST_CASE("Unit_hipGraphAddMemsetNode_Negative") {
+  hipError_t ret;
+  hipGraph_t graph;
+  hipGraphNode_t memsetNode;
+  char *devData;
+
+  HIP_CHECK(hipMalloc(&devData, 1024));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  hipMemsetParams memsetParams{};
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(devData);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = 1024;
+  memsetParams.height = 1;
+
+  SECTION("Pass pGraphNode as nullptr") {
+    ret = hipGraphAddMemsetNode(nullptr, graph, nullptr, 0, &memsetParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass pGraphNode as un-initialize object") {
+    hipGraphNode_t memsetNode_1;
+    ret = hipGraphAddMemsetNode(&memsetNode_1, graph,
+                                nullptr, 0, &memsetParams);
+    REQUIRE(hipSuccess == ret);
+  }
+  SECTION("Pass graph as nullptr") {
+    ret = hipGraphAddMemsetNode(&memsetNode, nullptr,
+                                nullptr, 0, &memsetParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass Graph as empty object") {
+    hipGraph_t graph_1{};
+    ret = hipGraphAddMemsetNode(&memsetNode, graph_1,
+                                nullptr, 0, &memsetParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass pDependencies as nullptr") {
+    ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams);
+    REQUIRE(hipSuccess == ret);
+  }
+  SECTION("Pass numDependencies is max and pDependencies is not valid ptr") {
+    ret = hipGraphAddMemsetNode(&memsetNode, graph,
+                                nullptr, INT_MAX, &memsetParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass pDependencies as nullptr, but numDependencies is non-zero") {
+    ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 9, &memsetParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass pMemsetParams as nullptr") {
+    ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, nullptr);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass pMemsetParams as un-initialize object") {
+    hipMemsetParams memsetParams1;
+    ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+                                &memsetParams1);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass hipMemsetParams::dst as nullptr") {
+    memsetParams.dst = nullptr;
+    ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass hipMemsetParams::element size other than 1, 2, or 4") {
+    memsetParams.dst = reinterpret_cast<void*>(devData);
+    memsetParams.elementSize = 9;
+    ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass hipMemsetParams::height as zero") {
+    memsetParams.elementSize = sizeof(char);
+    memsetParams.height = 0;
+    ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  HIP_CHECK(hipFree(devData));
+  HIP_CHECK(hipGraphDestroy(graph));
+}

tests/catch/unit/graph/hipGraphMemcpyNodeGetParams.cc

@@ -0,0 +1,232 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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 :
+Negative -
+1) Pass node as nullptr and verify api returns error code.
+2) Pass un-initialize node and verify api returns error code.
+3) Pass pNodeParams as nullptr and verify api returns error code.
+Functional -
+1) Create a graph, add Memcpy node to graph with desired node params.
+   Verify api fetches the node params mentioned while adding Memcpy node.
+2) Set Memcpy node params with hipGraphMemcpyNodeSetParams,
+   now get the params and verify both are same.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+
+#define SIZE 10
+#define UPDATESIZE 8
+
+/* Test verifies hipGraphMemcpyNodeGetParams API Negative scenarios.
+ */
+TEST_CASE("Unit_hipGraphMemcpyNodeGetParams_Negative") {
+  constexpr int width{SIZE}, height{SIZE}, depth{SIZE};
+  hipArray *devArray;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparms;
+  int* hData;
+  uint32_t size = width * height * depth * sizeof(int);
+  hData = reinterpret_cast<int*>(malloc(size));
+  REQUIRE(hData != nullptr);
+  memset(hData, 0, size);
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        hData[i*width*height + j*width + k] = i*width*height + j*width + k;
+      }
+    }
+  }
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
+                                              0, 0, 0, formatKind);
+  HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width,
+                             height, depth), hipArrayDefault));
+  memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.extent = make_hipExtent(width , height, depth);
+  myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int),
+                                      width, height);
+  myparms.dstArray = devArray;
+  myparms.kind = hipMemcpyHostToDevice;
+
+  hipGraph_t graph;
+  hipError_t ret;
+  hipGraphNode_t memcpyNode;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms));
+
+  SECTION("Pass node as nullptr") {
+    ret = hipGraphMemcpyNodeGetParams(nullptr, &myparms);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass un-initilize node") {
+    hipGraphNode_t memcpyNode_uninit{};
+    ret = hipGraphMemcpyNodeGetParams(memcpyNode_uninit, &myparms);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass GetNodeParams as nullptr") {
+    ret = hipGraphMemcpyNodeGetParams(memcpyNode, nullptr);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  HIP_CHECK(hipFreeArray(devArray));
+  free(hData);
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+
+/* Test verifies hipGraphMemcpyNodeGetParams API Functional scenarios.
+ */
+
+static bool compareHipPos(hipPos hPos1, hipPos hPos2) {
+  if ((hPos1.x == hPos2.x) && (hPos1.y == hPos2.y) && (hPos1.z == hPos2.z))
+    return true;
+  else
+    return false;
+}
+static bool compareHipExtent(hipExtent hExt1, hipExtent hExt2) {
+  if ((hExt1.width == hExt2.width) && (hExt1.height == hExt2.height) &&
+      (hExt1.depth == hExt2.depth))
+    return true;
+  else
+    return false;
+}
+static bool compareHipPitchedPtr(hipPitchedPtr hpPtr1, hipPitchedPtr hpPtr2) {
+  if ((reinterpret_cast<int *>(hpPtr1.ptr) ==
+       reinterpret_cast<int *>(hpPtr2.ptr))
+       && (hpPtr1.pitch == hpPtr2.pitch)
+       #if HT_AMD
+       && (hpPtr1.xsize == hpPtr2.xsize)
+       /* xsize check below is disabled on nvidia as xsize value
+        * is not being updated properly due to issue with CUDA api */
+       #endif
+       && (hpPtr1.ysize == hpPtr2.ysize))
+    return true;
+  else
+    return false;
+}
+
+static bool memcpyNodeCompare(hipMemcpy3DParms *mNode1,
+                              hipMemcpy3DParms *mNode2) {
+  if (mNode1->srcArray != mNode2->srcArray)
+    return false;
+  if (!compareHipPos(mNode1->srcPos, mNode2->srcPos))
+    return false;
+  if (!compareHipPitchedPtr(mNode1->srcPtr, mNode2->srcPtr))
+    return false;
+  if (mNode1->dstArray != mNode2->dstArray)
+    return false;
+  if (!compareHipPos(mNode1->dstPos, mNode2->dstPos))
+    return false;
+  if (!compareHipPitchedPtr(mNode1->dstPtr, mNode2->dstPtr))
+    return false;
+  if (!compareHipExtent(mNode1->extent, mNode2->extent))
+    return false;
+  if (mNode1->kind != mNode2->kind)
+    return false;
+  return true;
+}
+
+TEST_CASE("Unit_hipGraphMemcpyNodeGetParams_Functional") {
+  constexpr int width{SIZE}, height{SIZE}, depth{SIZE};
+  hipArray *devArray;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparms;
+  int* hData;
+  uint32_t size = width * height * depth * sizeof(int);
+  hData = reinterpret_cast<int*>(malloc(size));
+  REQUIRE(hData != nullptr);
+  memset(hData, 0, size);
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        hData[i*width*height + j*width + k] = i*width*height + j*width + k;
+      }
+    }
+  }
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
+                                              0, 0, 0, formatKind);
+  HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width,
+                             height, depth), hipArrayDefault));
+  memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.extent = make_hipExtent(width , height, depth);
+  myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int),
+                                      width, height);
+  myparms.dstArray = devArray;
+  myparms.kind = hipMemcpyHostToDevice;
+
+  hipGraph_t graph;
+  hipGraphNode_t memcpyNode;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms));
+
+  SECTION("Get Memcpy Param and verify.") {
+    hipMemcpy3DParms m3DGetParams;
+    REQUIRE(hipSuccess == hipGraphMemcpyNodeGetParams(memcpyNode,
+                                                      &m3DGetParams));
+    // Validating the result
+    REQUIRE(true == memcpyNodeCompare(&myparms, &m3DGetParams));
+  }
+
+  SECTION("Set memcpy params and Get param and verify.") {
+    hipMemcpy3DParms myparms1, m3DGetParams1;
+    constexpr int width1{UPDATESIZE}, height1{UPDATESIZE}, depth1{UPDATESIZE};
+    hipArray *devArray1;
+    hipChannelFormatKind formatKind1 = hipChannelFormatKindSigned;
+    int* hData1;
+    uint32_t size1 = width1 * height1 * depth1 * sizeof(int);
+    hData1 = reinterpret_cast<int*>(malloc(size1));
+    REQUIRE(hData1 != nullptr);
+    memset(hData1, 0, size1);
+    for (int i = 0; i < depth1; i++) {
+      for (int j = 0; j < height1; j++) {
+        for (int k = 0; k < width1; k++) {
+          hData1[i*width1*height1 + j*width1 + k] = i*width1*height1 +
+                                                    j*width1 + k;
+        }
+      }
+    }
+    hipChannelFormatDesc channelDesc1 = hipCreateChannelDesc(sizeof(int)*8,
+                                              0, 0, 0, formatKind1);
+    HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc1,
+              make_hipExtent(width1, height1, depth1), hipArrayDefault));
+    memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms));
+    myparms1.srcPos = make_hipPos(0, 0, 0);
+    myparms1.dstPos = make_hipPos(0, 0, 0);
+    myparms1.extent = make_hipExtent(width1 , height1, depth1);
+    myparms1.srcPtr = make_hipPitchedPtr(hData1, width1 * sizeof(int),
+                                         width1, height1);
+    myparms1.dstArray = devArray1;
+    myparms1.kind = hipMemcpyHostToDevice;
+
+    REQUIRE(hipSuccess == hipGraphMemcpyNodeSetParams(memcpyNode, &myparms1));
+    REQUIRE(hipSuccess == hipGraphMemcpyNodeGetParams(memcpyNode,
+                                                      &m3DGetParams1));
+    REQUIRE(true == memcpyNodeCompare(&myparms1, &m3DGetParams1));
+
+    HIP_CHECK(hipFreeArray(devArray1));
+    free(hData1);
+  }
+  HIP_CHECK(hipFreeArray(devArray));
+  free(hData);
+  HIP_CHECK(hipGraphDestroy(graph));
+}

tests/catch/unit/graph/hipGraphMemcpyNodeSetParams.cc

@@ -0,0 +1,215 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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 :
+Negative -
+1) Pass node as nullptr and verify api returns error code.
+2) Pass un-initialize node and verify api returns error code.
+3) Pass pNodeParams as nullptr and verify api returns error code.
+Functional -
+1) Add Memcpy node to graph, update the Memcpy node params with set and
+   launch the graph and check updated params are taking effect.
+2) Add Memcpy node to graph, launch graph, then update the Memcpy node params
+   with set and launch the graph and check updated params are taking effect.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+
+#define SIZE 10
+
+/* Test verifies hipGraphMemcpyNodeSetParams API Negative scenarios.
+ */
+TEST_CASE("Unit_hipGraphMemcpyNodeSetParams_Negative") {
+  constexpr int width{SIZE}, height{SIZE}, depth{SIZE};
+  hipArray *devArray;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparms;
+  int* hData;
+  uint32_t size = width * height * depth * sizeof(int);
+  hData = reinterpret_cast<int*>(malloc(size));
+  REQUIRE(hData != nullptr);
+  memset(hData, 0, size);
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        hData[i*width*height + j*width + k] = i*width*height + j*width + k;
+      }
+    }
+  }
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
+                                              0, 0, 0, formatKind);
+  HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width,
+                             height, depth), hipArrayDefault));
+  memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.extent = make_hipExtent(width , height, depth);
+  myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int),
+                                      width, height);
+  myparms.dstArray = devArray;
+  myparms.kind = hipMemcpyHostToDevice;
+
+  hipGraph_t graph;
+  hipError_t ret;
+  hipGraphNode_t memcpyNode;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms));
+
+  SECTION("Pass node as nullptr") {
+    ret = hipGraphMemcpyNodeSetParams(nullptr, &myparms);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass un-initialize node") {
+    hipGraphNode_t memcpyNode_uninit{};
+    ret = hipGraphMemcpyNodeSetParams(memcpyNode_uninit, &myparms);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass SetNodeParams as nullptr") {
+    ret = hipGraphMemcpyNodeSetParams(memcpyNode, nullptr);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  HIP_CHECK(hipFreeArray(devArray));
+  free(hData);
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+
+/* Test verifies hipGraphMemcpyNodeSetParams API Functional scenarios.
+ */
+TEST_CASE("Unit_hipGraphMemcpyNodeSetParams_Functional") {
+  constexpr int width{SIZE}, height{SIZE}, depth{SIZE};
+  hipArray *devArray;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparms, myparms1;
+  uint32_t size = width * height * depth * sizeof(int);
+
+  int *hData = reinterpret_cast<int*>(malloc(size));
+  REQUIRE(hData != nullptr);
+  memset(hData, 0, size);
+  int *hDataTemp = reinterpret_cast<int*>(malloc(size));
+  REQUIRE(hDataTemp != nullptr);
+  memset(hDataTemp, 0, size);
+  int *hOutputData = reinterpret_cast<int *>(malloc(size));
+  REQUIRE(hOutputData != nullptr);
+  memset(hOutputData, 0,  size);
+  int *hOutputData1 = reinterpret_cast<int *>(malloc(size));
+  REQUIRE(hOutputData1 != nullptr);
+  memset(hOutputData1, 0,  size);
+
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        hData[i*width*height + j*width + k] = i*width*height + j*width + k;
+      }
+    }
+  }
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
+                                              0, 0, 0, formatKind);
+  HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width,
+                             height, depth), hipArrayDefault));
+  memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
+
+  // Host to Device
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.extent = make_hipExtent(width , height, depth);
+  myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int),
+                                      width, height);
+  myparms.dstArray = devArray;
+  myparms.kind = hipMemcpyHostToDevice;
+
+  hipGraph_t graph;
+  hipGraphNode_t memcpyNode;
+  std::vector<hipGraphNode_t> dependencies;
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms));
+  dependencies.push_back(memcpyNode);
+
+  // Device to host
+  memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms));
+  myparms1.srcPos = make_hipPos(0, 0, 0);
+  myparms1.dstPos = make_hipPos(0, 0, 0);
+  myparms1.dstPtr = make_hipPitchedPtr(hDataTemp, width * sizeof(int),
+                                      width, height);
+  myparms1.srcArray = devArray;
+  myparms1.extent = make_hipExtent(width, height, depth);
+  myparms1.kind = hipMemcpyDeviceToHost;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                  dependencies.size(), &myparms1));
+
+  SECTION("Update the memcpyNode and check") {
+    // Device to host with updated host ptr hDataTemp -> hOutputData
+    memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms));
+    myparms1.srcPos = make_hipPos(0, 0, 0);
+    myparms1.dstPos = make_hipPos(0, 0, 0);
+    myparms1.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(int),
+                                        width, height);
+    myparms1.srcArray = devArray;
+    myparms1.extent = make_hipExtent(width, height, depth);
+    myparms1.kind = hipMemcpyDeviceToHost;
+
+    HIP_CHECK(hipGraphMemcpyNodeSetParams(memcpyNode, &myparms1));
+
+    // Instantiate and launch the graph
+    HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+    HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+    HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+    // Check result
+    HipTest::checkArray(hData, hOutputData, width, height, depth);
+  }
+
+  SECTION("Update the memcpyNode again and check") {
+    // Device to host with updated host ptr hOutputData -> hOutputData1
+    memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms));
+    myparms1.srcPos = make_hipPos(0, 0, 0);
+    myparms1.dstPos = make_hipPos(0, 0, 0);
+    myparms1.dstPtr = make_hipPitchedPtr(hOutputData1, width * sizeof(int),
+                                        width, height);
+    myparms1.srcArray = devArray;
+    myparms1.extent = make_hipExtent(width, height, depth);
+    myparms1.kind = hipMemcpyDeviceToHost;
+
+    HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                    dependencies.size(), &myparms1));
+    HIP_CHECK(hipGraphMemcpyNodeSetParams(memcpyNode, &myparms1));
+
+    // Instantiate and launch the graph
+    HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+    HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+    HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+    // Check result
+    HipTest::checkArray(hData, hOutputData1, width, height, depth);
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  HIP_CHECK(hipFreeArray(devArray));
+  free(hData);
+  free(hDataTemp);
+  free(hOutputData);
+  free(hOutputData1);
+}

tests/catch/unit/memory/CMakeLists.txt

@@ -1,4 +1,4 @@
-# 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
@@ -21,6 +21,7 @@
 # Common Tests - Test independent of all platforms
 if(HIP_PLATFORM MATCHES "amd")
 set(TEST_SRC
+    DriverContext.cc
     memset.cc
     malloc.cc
     hipMemcpy2DToArray.cc
@@ -38,6 +39,8 @@ set(TEST_SRC
     hipMemcpyAllApiNegative.cc
     hipMemcpy_MultiThread.cc
     hipHostRegister.cc
+    hipHostUnregister.cc
+    hipMallocPitch.cc
     hipMemPtrGetInfo.cc
     hipPointerGetAttributes.cc
     hipHostGetFlags.cc
@@ -73,11 +76,13 @@ set(TEST_SRC
     hipHostMalloc.cc
     hipMemcpy.cc
     hipMemcpyAsync.cc
+    hipMemsetFunctional.cc
     hipMallocPitch.cc
     hipMallocArray.cc
     hipMalloc3D.cc
     hipMalloc3DArray.cc
     hipArrayCreate.cc
+    hipArray3DCreate.cc
     hipDrvMemcpy3D.cc
     hipDrvMemcpy3DAsync.cc
     hipPointerGetAttribute.cc
@@ -85,9 +90,11 @@ set(TEST_SRC
     hipMallocMngdMultiThread.cc
     hipMemPrefetchAsync.cc
     hipArray.cc
+    hipMemVmm.cc
 )
 else()
 set(TEST_SRC
+    DriverContext.cc
     memset.cc
     malloc.cc
     hipMemcpy2DToArray.cc
@@ -105,6 +112,8 @@ set(TEST_SRC
     hipMemcpyAllApiNegative.cc
     hipMemcpy_MultiThread.cc
     hipHostRegister.cc
+    hipHostUnregister.cc
+    hipMallocPitch.cc
     hipHostGetFlags.cc
     hipHostGetDevicePointer.cc
     hipMallocManaged_MultiScenario.cc
@@ -136,11 +145,13 @@ set(TEST_SRC
     hipHostMalloc.cc
     hipMemcpy.cc
     hipMemcpyAsync.cc
+    hipMemsetFunctional.cc
     hipMallocPitch.cc
     hipMallocArray.cc
     hipMalloc3D.cc
     hipMalloc3DArray.cc
     hipArrayCreate.cc
+    hipArray3DCreate.cc
     hipDrvMemcpy3D.cc
     hipDrvMemcpy3DAsync.cc
     hipPointerGetAttribute.cc
@@ -159,4 +170,5 @@ endif()
 
 hip_add_exe_to_target(NAME MemoryTest
                       TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests)
+                      TEST_TARGET_NAME build_tests
+                      COMPILE_OPTIONS -std=c++14)

tests/catch/unit/memory/DriverContext.cc

@@ -0,0 +1,36 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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.
+*/
+
+#include "DriverContext.hh"
+#include <hip_test_common.hh>
+
+DriverContext::DriverContext() {
+  HIP_CHECK(hipInit(0));
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  HIP_CHECK(hipDevicePrimaryCtxRetain(&ctx, device));
+  HIP_CHECK(hipCtxPushCurrent(ctx));
+}
+
+DriverContext::~DriverContext() {
+  HIP_CHECK(hipCtxPopCurrent(&ctx));
+  HIP_CHECK(hipDevicePrimaryCtxRelease(device));
+}

tests/catch/unit/memory/DriverContext.hh

@@ -0,0 +1,39 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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.
+*/
+
+#pragma once
+
+#include <hip_test_context.hh>
+
+class DriverContext {
+ private:
+  hipCtx_t ctx;
+  hipDevice_t device;
+
+ public:
+  DriverContext();
+  ~DriverContext();
+
+  // Rule of three
+  DriverContext(const DriverContext& other) = delete;
+  DriverContext(DriverContext&& other) noexcept = delete;
+};

tests/catch/unit/memory/hipArray3DCreate.cc

@@ -0,0 +1,321 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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.
+*/
+
+#include <limits>
+#include "DriverContext.hh"
+#include "hipArrayCommon.hh"
+#include "hip_test_common.hh"
+
+namespace {
+void checkArrayIsExpected(const hiparray array, const HIP_ARRAY3D_DESCRIPTOR& expected_desc) {
+// hipArray3DGetDescriptor doesn't currently exist (EXSWCPHIPT-87)
+#if HT_AMD
+  std::ignore = array;
+  std::ignore = expected_desc;
+#else
+  CUDA_ARRAY3D_DESCRIPTOR queried_desc;
+  cuArray3DGetDescriptor(&queried_desc, array);
+
+  REQUIRE(queried_desc.Width == expected_desc.Width);
+  REQUIRE(queried_desc.Height == expected_desc.Height);
+  REQUIRE(queried_desc.Depth == expected_desc.Depth);
+  REQUIRE(queried_desc.Format == expected_desc.Format);
+  REQUIRE(queried_desc.NumChannels == expected_desc.NumChannels);
+  REQUIRE(queried_desc.Flags == expected_desc.Flags);
+#endif
+}
+
+void testInvalidDescription(HIP_ARRAY3D_DESCRIPTOR desc) {
+  hiparray array;
+  HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue);
+}
+}  // namespace
+
+TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, short4, float,
+                   float2, float4) {
+  using vec_info = vector_info<TestType>;
+  DriverContext ctx;
+
+  HIP_ARRAY3D_DESCRIPTOR desc{};
+  desc.Format = vec_info::format;
+  desc.NumChannels = vec_info::size;
+#if HT_AMD
+  desc.Flags = 0;
+#else
+  desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST);
+#endif
+
+  constexpr size_t size = 64;
+
+  std::vector<hipExtent> extents{
+      {size, 0, 0},       // 1D array
+      {size, size, 0},    // 2D array
+      {size, size, size}  // 3D array
+  };
+
+  for (auto& extent : extents) {
+    desc.Width = extent.width;
+    desc.Height = extent.height;
+    desc.Depth = extent.depth;
+
+    CAPTURE(desc.Width, desc.Height, desc.Depth);
+
+    hiparray array;
+    HIP_CHECK(hipArray3DCreate(&array, &desc));
+    checkArrayIsExpected(array, desc);
+    HIP_CHECK(hipArrayDestroy(array));
+  }
+}
+
+TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, ushort2,
+                   unsigned char, float, float4) {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-97");
+  return;
+#endif
+
+  using vec_info = vector_info<TestType>;
+  DriverContext ctx;
+
+  hiparray array;
+  HIP_ARRAY3D_DESCRIPTOR desc{};
+  desc.Format = vec_info::format;
+  desc.NumChannels = vec_info::size;
+#if HT_AMD
+  desc.Flags = 0;
+#else
+  desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST);
+  if (desc.Flags == CUDA_ARRAY3D_SURFACE_LDST) {
+    HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58");
+    return;
+  }
+#endif
+  CAPTURE(desc.Flags);
+
+  const Sizes sizes(desc.Flags);
+  CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D);
+
+  const size_t s = 64;
+  SECTION("Happy") {
+    // stored in a vector so some values can be ifdef'd out
+    std::vector<hipExtent> extentsToTest{
+        make_hipExtent(sizes.max1D, 0, 0),                              // 1D max
+        make_hipExtent(sizes.max2D[0], s, 0),                           // 2D max width
+        make_hipExtent(s, sizes.max2D[1], 0),                           // 2D max height
+        make_hipExtent(sizes.max2D[0], sizes.max2D[1], 0),              // 2D max
+        make_hipExtent(sizes.max3D[0], s, s),                           // 3D max width
+        make_hipExtent(s, sizes.max3D[1], s),                           // 3D max height
+        make_hipExtent(s, s, sizes.max3D[2]),                           // 3D max depth
+        make_hipExtent(s, sizes.max3D[1], sizes.max3D[2]),              // 3D max height and depth
+        make_hipExtent(sizes.max3D[0], s, sizes.max3D[2]),              // 3D max width and depth
+        make_hipExtent(sizes.max3D[0], sizes.max3D[1], s),              // 3D max width and height
+        make_hipExtent(sizes.max3D[0], sizes.max3D[1], sizes.max3D[2])  // 3D max
+    };
+    const auto extent =
+        GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest)));
+
+    desc.Width = extent.width;
+    desc.Height = extent.height;
+    desc.Depth = extent.depth;
+
+    CAPTURE(desc.Width, desc.Height, desc.Depth);
+
+    auto maxArrayCreateError = hipArray3DCreate(&array, &desc);
+    // this can try to alloc many GB of memory, so out of memory is acceptable
+    if (maxArrayCreateError == hipErrorOutOfMemory) return;
+    HIP_CHECK(maxArrayCreateError);
+    checkArrayIsExpected(array, desc);
+    HIP_CHECK(hipArrayDestroy(array));
+  }
+  SECTION("Negative") {
+    std::vector<hipExtent> extentsToTest {
+      make_hipExtent(sizes.max1D + 1, 0, 0),                          // 1D max
+          make_hipExtent(sizes.max2D[0] + 1, s, 0),                   // 2D max width
+          make_hipExtent(s, sizes.max2D[1] + 1, 0),                   // 2D max height
+          make_hipExtent(sizes.max2D[0] + 1, sizes.max2D[1] + 1, 0),  // 2D max
+          make_hipExtent(sizes.max3D[0] + 1, s, s),                   // 3D max width
+          make_hipExtent(s, sizes.max3D[1] + 1, s),                   // 3D max height
+#if !HT_NVIDIA                                       // leads to hipSuccess on NVIDIA
+          make_hipExtent(s, s, sizes.max3D[2] + 1),  // 3D max depth
+#endif
+          make_hipExtent(s, sizes.max3D[1] + 1, sizes.max3D[2] + 1),  // 3D max height and depth
+          make_hipExtent(sizes.max3D[0] + 1, s, sizes.max3D[2] + 1),  // 3D max width and depth
+          make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, s),  // 3D max width and height
+          make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, sizes.max3D[2] + 1)  // 3D max
+    };
+    const auto extent =
+        GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest)));
+
+    desc.Width = extent.width;
+    desc.Height = extent.height;
+    desc.Depth = extent.depth;
+
+    CAPTURE(desc.Width, desc.Height, desc.Depth);
+
+    HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue);
+  }
+}
+
+#if HT_NVIDIA
+constexpr std::array<unsigned int, 9> validFlags{
+    0,
+    hipArraySurfaceLoadStore,
+    hipArrayLayered,
+    hipArrayLayered | hipArraySurfaceLoadStore,
+    hipArrayCubemap,
+    hipArrayCubemap | hipArrayLayered,
+    hipArrayCubemap | hipArraySurfaceLoadStore,
+    hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore,
+    hipArrayTextureGather};
+#else
+constexpr std::array<unsigned int, 5> validFlags{
+    0, hipArrayCubemap, hipArrayCubemap | hipArrayLayered,
+    hipArrayCubemap | hipArraySurfaceLoadStore,
+    hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore};
+#endif
+
+constexpr HIP_ARRAY3D_DESCRIPTOR defaultDescriptor(unsigned int flags, size_t size) {
+  HIP_ARRAY3D_DESCRIPTOR desc{};
+  desc.Format = HIP_AD_FORMAT_FLOAT;
+  desc.NumChannels = 4;
+  desc.Flags = flags;
+  desc.Width = size;
+  desc.Height = size;
+  desc.Depth = size;
+
+#if HT_NVIDIA
+  if (flags == CUDA_ARRAY3D_TEXTURE_GATHER) {
+    desc.Depth = 0;
+  }
+#endif
+  return desc;
+}
+
+// Providing the array pointer as nullptr should return an error
+TEST_CASE("Unit_hipArray3DCreate_Negative_NullArrayPtr") {
+  auto desc = defaultDescriptor(0, 64);
+
+  DriverContext ctx;
+  HIP_CHECK_ERROR(hipArray3DCreate(nullptr, &desc), hipErrorInvalidValue);
+}
+
+// Providing the description pointer as nullptr should return an error
+TEST_CASE("Unit_hipArray3DCreate_Negative_NullDescPtr") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-130");
+  return;
+#endif
+
+  DriverContext ctx;
+  hiparray array;
+  HIP_CHECK_ERROR(hipArray3DCreate(&array, nullptr), hipErrorInvalidValue);
+}
+
+
+// Zero width arrays are not allowed
+TEST_CASE("Unit_hipArray3DCreate_Negative_ZeroWidth") {
+  DriverContext ctx;
+
+  unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+  auto desc = defaultDescriptor(flags, 6);
+  desc.Width = 0;
+  CAPTURE(desc.Flags);
+
+  testInvalidDescription(desc);
+}
+
+// Zero height arrays are only allowed for 1D arrays and layered arrays
+TEST_CASE("Unit_hipArray3DCreate_Negative_ZeroHeight") {
+  DriverContext ctx;
+
+  unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+  auto desc = defaultDescriptor(flags, 6);
+#if HT_NVIDIA
+  std::array<unsigned int, 2> exceptions{CUDA_ARRAY3D_LAYERED,
+                                         CUDA_ARRAY3D_LAYERED | CUDA_ARRAY3D_SURFACE_LDST};
+#else
+  std::array<unsigned int, 0> exceptions{};
+#endif
+  desc.Height = 0;
+
+  if (std::find(std::begin(exceptions), std::end(exceptions), desc.Flags) == std::end(exceptions)) {
+    // flag is not in list of exceptions
+    testInvalidDescription(desc);
+  }
+}
+
+// Arrays must be created with a valid data format
+TEST_CASE("Unit_hipArray3DCreate_Negative_InvalidFormat") {
+  DriverContext ctx;
+
+  unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+  auto desc = defaultDescriptor(flags, 6);
+
+  desc.Format = static_cast<hipArray_Format>(0xDEADBEEF);
+  REQUIRE(std::find(std::begin(driverFormats), std::end(driverFormats), desc.Format) ==
+          std::end(driverFormats));
+
+  testInvalidDescription(desc);
+}
+
+// An array must have either 1,2, or 4 channels
+TEST_CASE("Unit_hipArray3DCreate_Negative_NumChannels") {
+  DriverContext ctx;
+  unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+  auto desc = defaultDescriptor(flags, 6);
+  desc.NumChannels = GENERATE(0, 3, 5);
+
+  testInvalidDescription(desc);
+}
+
+// Using invalid flags should result in an error
+TEST_CASE("Unit_hipArray3DCreate_Negative_InvalidFlags") {
+  DriverContext ctx;
+
+  // FIXME: use the same flags for both tests when the values exist for hip
+#if HT_NVIDIA
+  unsigned int flags =
+      GENERATE(0xDEADBEEF, CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST,
+               CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_CUBEMAP,
+               CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST | CUDA_ARRAY3D_CUBEMAP);
+#else
+  unsigned int flags = 0xDEADBEEF;
+#endif
+
+  CAPTURE(flags);
+
+  auto desc = defaultDescriptor(flags, 6);
+
+
+  REQUIRE(std::find(std::begin(validFlags), std::end(validFlags), desc.Flags) ==
+          std::end(validFlags));
+
+  testInvalidDescription(desc);
+}
+
+
+// hipArray3DCreate should handle the max numeric value gracefully.
+TEST_CASE("Unit_hipArray3DCreate_Negative_NumericLimit") {
+  DriverContext ctx;
+
+  unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+  auto desc = defaultDescriptor(flags, std::numeric_limits<size_t>::max());
+
+  testInvalidDescription(desc);
+}

tests/catch/unit/memory/hipArrayCommon.hh

@@ -0,0 +1,211 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+
+constexpr size_t BlockSize = 16;
+
+template <class T, size_t N, hipArray_Format Format> struct type_and_size_and_format {
+  using type = T;
+  static constexpr size_t size = N;
+  static constexpr hipArray_Format format = Format;
+};
+
+// Create a map of type to scalar type, vector size and scalar type format enum.
+// This is useful for creating simpler function that depend on the vector size.
+template <typename T> struct vector_info;
+template <>
+struct vector_info<int> : type_and_size_and_format<int, 1, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float> : type_and_size_and_format<float, 1, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short> : type_and_size_and_format<short, 1, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char> : type_and_size_and_format<char, 1, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<unsigned int>
+    : type_and_size_and_format<unsigned int, 1, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<unsigned short>
+    : type_and_size_and_format<unsigned short, 1, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<unsigned char>
+    : type_and_size_and_format<unsigned char, 1, HIP_AD_FORMAT_UNSIGNED_INT8> {};
+
+template <>
+struct vector_info<int2> : type_and_size_and_format<int, 2, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float2> : type_and_size_and_format<float, 2, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short2> : type_and_size_and_format<short, 2, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char2> : type_and_size_and_format<char, 2, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<uint2>
+    : type_and_size_and_format<unsigned int, 2, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<ushort2>
+    : type_and_size_and_format<unsigned short, 2, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<uchar2>
+    : type_and_size_and_format<unsigned char, 2, HIP_AD_FORMAT_UNSIGNED_INT8> {};
+
+template <>
+struct vector_info<int4> : type_and_size_and_format<int, 4, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float4> : type_and_size_and_format<float, 4, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short4> : type_and_size_and_format<short, 4, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char4> : type_and_size_and_format<char, 4, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<uint4>
+    : type_and_size_and_format<unsigned int, 4, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<ushort4>
+    : type_and_size_and_format<unsigned short, 4, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<uchar4>
+    : type_and_size_and_format<unsigned char, 4, HIP_AD_FORMAT_UNSIGNED_INT8> {};
+
+// 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> 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);
+}
+
+inline size_t getFreeMem() {
+  size_t free = 0, total = 0;
+  HIP_CHECK(hipMemGetInfo(&free, &total));
+  return free;
+}
+
+struct Sizes {
+  int max1D;
+  std::array<int, 2> max2D;
+  std::array<int, 3> max3D;
+
+  Sizes(unsigned int flag) {
+    int device;
+    HIP_CHECK(hipGetDevice(&device));
+    static_assert(
+        hipArrayDefault == 0,
+        "hipArrayDefault is assumed to be equivalent to 0 for the following switch statment");
+#if HT_NVIDIA
+    static_assert(hipArraySurfaceLoadStore == CUDA_ARRAY3D_SURFACE_LDST,
+                  "hipArraySurface is assumed to be equivalent to CUDA_ARRAY3D_SURFACE_LDST for "
+                  "the following switch statment");
+#endif
+    switch (flag) {
+      case hipArrayDefault: {  // 0
+        hipDeviceProp_t prop;
+        HIP_CHECK(hipGetDeviceProperties(&prop, device));
+        max1D = prop.maxTexture1D;
+        max2D = {prop.maxTexture2D[0], prop.maxTexture2D[1]};
+        max3D = {prop.maxTexture3D[0], prop.maxTexture3D[1], prop.maxTexture3D[2]};
+        return;
+      }
+      case hipArraySurfaceLoadStore: {  // CUDA_ARRAY3D_SURFACE_LDST
+        int value;
+        HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface1D, device));
+        max1D = value;
+        HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface2D, device));
+        max2D = {value, value};
+        HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface3D, device));
+        max3D = {value, value, value};
+        return;
+      }
+      default: {
+        INFO("Array flag not supported");
+        REQUIRE(false);
+        return;
+      }
+    }
+  }
+};
+
+inline const char* channelFormatString(hipChannelFormatKind formatKind) noexcept {
+  switch (formatKind) {
+    case hipChannelFormatKindFloat:
+      return "float";
+    case hipChannelFormatKindSigned:
+      return "signed";
+    case hipChannelFormatKindUnsigned:
+      return "unsigned";
+    default:
+      return "error";
+  }
+}
+
+// All the possible formats for channel data in an array.
+static const std::vector<hipArray_Format> driverFormats{
+    HIP_AD_FORMAT_UNSIGNED_INT8, HIP_AD_FORMAT_UNSIGNED_INT16, HIP_AD_FORMAT_UNSIGNED_INT32,
+    HIP_AD_FORMAT_SIGNED_INT8,   HIP_AD_FORMAT_SIGNED_INT16,   HIP_AD_FORMAT_SIGNED_INT32,
+    HIP_AD_FORMAT_HALF,          HIP_AD_FORMAT_FLOAT};
+
+// Helpful for printing errors
+inline const char* formatToString(hipArray_Format f) {
+  switch (f) {
+    case HIP_AD_FORMAT_UNSIGNED_INT8:
+      return "Unsigned Int 8";
+    case HIP_AD_FORMAT_UNSIGNED_INT16:
+      return "Unsigned Int 16";
+    case HIP_AD_FORMAT_UNSIGNED_INT32:
+      return "Unsigned Int 32";
+    case HIP_AD_FORMAT_SIGNED_INT8:
+      return "Signed Int 8";
+    case HIP_AD_FORMAT_SIGNED_INT16:
+      return "Signed Int 16";
+    case HIP_AD_FORMAT_SIGNED_INT32:
+      return "Signed Int 32";
+    case HIP_AD_FORMAT_HALF:
+      return "Float 16";
+    case HIP_AD_FORMAT_FLOAT:
+      return "Float 32";
+    default:
+      return "not found";
+  }
+}

tests/catch/unit/memory/hipArrayCreate.cc

@@ -24,12 +24,16 @@ hipArrayCreate API test scenarios
 3. Multithreaded scenario
 */
 
+#include <array>
+#include <numeric>
 #include <hip_test_common.hh>
+#include "hipArrayCommon.hh"
+#include "DriverContext.hh"
 
-static constexpr auto NUM_W{4};
-static constexpr auto BIGNUM_W{100};
-static constexpr auto NUM_H{4};
-static constexpr auto BIGNUM_H{100};
+static constexpr size_t NUM_W{4};
+static constexpr size_t BIGNUM_W{100};
+static constexpr size_t NUM_H{4};
+static constexpr size_t BIGNUM_H{100};
 static constexpr auto ARRAY_LOOP{100};
 
 /*
@@ -47,77 +51,35 @@ static constexpr auto ARRAY_LOOP{100};
  */
 
 static void ArrayCreate_DiffSizes(int gpu) {
-  HIP_CHECK(hipSetDevice(gpu));
-  std::vector<size_t> array_size;
-  array_size.push_back(NUM_W);
-  array_size.push_back(BIGNUM_W);
-  for (auto &size : array_size) {
-    HIP_ARRAY array[ARRAY_LOOP];
-    size_t tot, avail, ptot, pavail;
-    HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
-    for (int i = 0; i < ARRAY_LOOP; i++) {
-      HIP_ARRAY_DESCRIPTOR desc;
-      desc.NumChannels = 1;
-      if (size == NUM_W) {
-        desc.Width = NUM_W;
-        desc.Height = NUM_H;
-      } else {
-        desc.Width = BIGNUM_W;
-        desc.Height = BIGNUM_H;
-      }
-      desc.Format = HIP_AD_FORMAT_FLOAT;
-      HIP_CHECK(hipArrayCreate(&array[i], &desc));
-    }
-    for (int i = 0; i < ARRAY_LOOP; i++) {
-      ARRAY_DESTROY(array[i]);
-    }
-    HIP_CHECK(hipMemGetInfo(&avail, &tot));
-    if ((pavail != avail)) {
-      HIPASSERT(false);
-    }
-  }
-}
+  HIP_CHECK_THREAD(hipSetDevice(gpu));
+  std::pair<size_t, size_t> size =
+      GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H));
+  std::array<HIP_ARRAY, ARRAY_LOOP> array;
+  size_t pavail, avail;
+  HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
+  HIP_ARRAY_DESCRIPTOR desc;
+  desc.NumChannels = 1;
+  desc.Width = std::get<0>(size);
+  desc.Height = std::get<1>(size);
+  desc.Format = HIP_AD_FORMAT_FLOAT;
 
-/*Thread function*/
-static void ArrayCreateThreadFunc(int gpu) {
-  ArrayCreate_DiffSizes(gpu);
+  for (int i = 0; i < ARRAY_LOOP; i++) {
+    HIP_CHECK_THREAD(hipArrayCreate(&array[i], &desc));
+  }
+  for (int i = 0; i < ARRAY_LOOP; i++) {
+    HIP_CHECK_THREAD(hipArrayDestroy(array[i]));
+  }
+
+  HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
+  REQUIRE_THREAD(pavail == avail);
 }
 
 /* This testcase verifies hipArrayCreate API for small and big chunks data*/
 TEST_CASE("Unit_hipArrayCreate_DiffSizes") {
   ArrayCreate_DiffSizes(0);
+  HIP_CHECK_THREAD_FINALIZE();
 }
 
-
-/* This testcase verifies the negative scenarios of
- * hipArrayCreate API
- */
-TEST_CASE("Unit_hipArrayCreate_Negative") {
-  HIP_ARRAY_DESCRIPTOR desc;
-  HIP_ARRAY array;
-  desc.Format = HIP_AD_FORMAT_FLOAT;
-  desc.NumChannels = 1;
-  desc.Width = NUM_W;
-  desc.Height = NUM_H;
-#if HT_NVIDIA
-  SECTION("NullPointer to Array") {
-    REQUIRE(hipArrayCreate(nullptr, &desc) != hipSuccess);
-  }
-
-  SECTION("NullPointer to Channel Descriptor") {
-    REQUIRE(hipArrayCreate(&array, nullptr) != hipSuccess);
-  }
-#endif
-  SECTION("Width 0 for Array Descriptor") {
-    desc.Width = 0;
-    REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess);
-  }
-
-  SECTION("Invalid NumChannels") {
-    desc.NumChannels = 3;
-    REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess);
-  }
-}
 /*
 This testcase verifies the hipArrayCreate API in multithreaded
 scenario by launching threads in parallel on multiple GPUs
@@ -129,16 +91,16 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") {
 
   devCnt = HipTest::getDeviceCount();
 
-  size_t tot, avail, ptot, pavail;
-  HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
+  const size_t pavail = getFreeMem();
   for (int i = 0; i < devCnt; i++) {
-    threadlist.push_back(std::thread(ArrayCreateThreadFunc, i));
+    threadlist.push_back(std::thread(ArrayCreate_DiffSizes, i));
   }
 
-  for (auto &t : threadlist) {
+  for (auto& t : threadlist) {
     t.join();
   }
-  HIP_CHECK(hipMemGetInfo(&avail, &tot));
+  HIP_CHECK_THREAD_FINALIZE();
+  const size_t avail = getFreeMem();
 
   if (pavail != avail) {
     WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
@@ -146,3 +108,273 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") {
   }
 }
 
+
+// Tests /////////////////////////////////////////
+
+#if HT_AMD
+constexpr auto MemoryTypeHost = hipMemoryTypeHost;
+constexpr auto MemoryTypeArray = hipMemoryTypeArray;
+constexpr auto NORMALIZED_COORDINATES = HIP_TRSF_NORMALIZED_COORDINATES;
+constexpr auto READ_AS_INTEGER = HIP_TRSF_READ_AS_INTEGER;
+#else
+constexpr auto MemoryTypeHost = CU_MEMORYTYPE_HOST;
+constexpr auto MemoryTypeArray = CU_MEMORYTYPE_ARRAY;
+// (EXSWCPHIPT-92) HIP equivalents not defined for CUDA backend.
+constexpr auto NORMALIZED_COORDINATES = CU_TRSF_NORMALIZED_COORDINATES;
+constexpr auto READ_AS_INTEGER = CU_TRSF_READ_AS_INTEGER;
+#endif
+
+// Copy data from host to the hiparray, accounting 1D or 2D arrays
+template <typename T>
+void copyToArray(hiparray dst, const std::vector<T>& src, const size_t height) {
+  const auto sizeInBytes = src.size() * sizeof(T);
+  if (height == 0) {
+    // FIXME(EXSWCPHIPT-64) remove cast when API is fixed (will require major version change)
+    HIP_CHECK(hipMemcpyHtoA(reinterpret_cast<hipArray*>(dst), 0, src.data(), sizeInBytes));
+  } else {
+    const auto pitch = sizeInBytes / height;
+    hip_Memcpy2D copyParams{};
+    copyParams.srcMemoryType = MemoryTypeHost;
+    copyParams.srcXInBytes = 0;  // x offset
+    copyParams.srcY = 0;         // y offset
+    copyParams.srcHost = src.data();
+    copyParams.srcPitch = pitch;
+
+
+    copyParams.dstMemoryType = MemoryTypeArray;
+    copyParams.dstXInBytes = 0;  // x offset
+    copyParams.dstY = 0;         // y offset
+    copyParams.dstArray = dst;
+
+    copyParams.WidthInBytes = pitch;
+    copyParams.Height = height;
+
+    HIP_CHECK(hipMemcpyParam2D(&copyParams));
+  }
+}
+
+// Test the allocated array by generating a texture from it then reading from that texture.
+// Textures are read-only, so write to the array then copy that into normal device memory.
+template <typename T>
+void testArrayAsTexture(hiparray array, const size_t width, const size_t height) {
+  using vec_info = vector_info<T>;
+  using scalar_type = typename vec_info::type;
+  const auto h = height ? height : 1;
+  const auto size = sizeof(T) * width * h;
+
+  // set hip array
+  std::vector<scalar_type> hostData(width * h * vec_info::size);
+  // assigned ascending values to the data array to show indexing is working
+  std::iota(std::begin(hostData), std::end(hostData), 0);
+
+  copyToArray(array, hostData, height);
+
+  // create texture
+  hipTextureObject_t textObj{};
+
+  HIP_RESOURCE_DESC resDesc{};
+  memset(&resDesc, 0, sizeof(HIP_RESOURCE_DESC));
+  resDesc.resType = HIP_RESOURCE_TYPE_ARRAY;
+  resDesc.res.array.hArray = array;
+  resDesc.flags = 0;
+
+  HIP_TEXTURE_DESC texDesc{};
+  memset(&texDesc, 0, sizeof(HIP_TEXTURE_DESC));
+  // use the actual values in the texture, not normalized data
+  texDesc.filterMode = HIP_TR_FILTER_MODE_POINT;
+  // Use normalized coordinates and also read the data in the original data type
+  texDesc.flags |= NORMALIZED_COORDINATES | READ_AS_INTEGER;
+
+  HIP_CHECK(hipTexObjectCreate(&textObj, &resDesc, &texDesc, 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(hipTexObjectDestroy(textObj));
+  HIP_CHECK(hipFree(device_data));
+}
+
+// Selection of types chosen since trying all types would be slow to compile
+// Test the happy path of the hipArrayCreate
+TEMPLATE_TEST_CASE("Unit_hipArrayCreate_happy", "", uint, int, int4, ushort, short2, char, uchar2,
+                   char4, float, float2, float4) {
+#if HT_AMD
+  if (std::is_same<uint, TestType>::value || std::is_same<ushort, TestType>::value ||
+      std::is_same<uchar2, TestType>::value) {
+    HipTest::HIP_SKIP_TEST("Probably EXSWCPHIPT-62");
+    return;
+  }
+#endif
+  using vec_info = vector_info<TestType>;
+  DriverContext ctx;
+
+  HIP_ARRAY_DESCRIPTOR desc;
+  desc.Format = vec_info::format;
+  desc.NumChannels = vec_info::size;
+  desc.Width = 1024;
+  desc.Height = GENERATE(0, 1024);
+
+  size_t initFree = getFreeMem();
+
+  // pointer to the array in device memory
+  hiparray array{};
+
+  HIP_CHECK(hipArrayCreate(&array, &desc));
+
+  testArrayAsTexture<TestType>(array, desc.Width, desc.Height);
+
+  size_t finalFree = getFreeMem();
+
+  const size_t allocSize = sizeof(TestType) * desc.Width * (desc.Height ? desc.Height : 1);
+  // will be aligned to some size, so this is not exact
+  REQUIRE(initFree - finalFree >= allocSize);
+
+  HIP_CHECK(hipArrayDestroy(array));
+}
+
+
+// Only widths and Heights up to the maxTexture size is supported
+TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort, short2, char,
+                   uchar2, char4, float, float2, float4) {
+  using vec_info = vector_info<TestType>;
+  DriverContext ctx;
+
+  HIP_ARRAY_DESCRIPTOR desc;
+  desc.Format = vec_info::format;
+  desc.NumChannels = vec_info::size;
+
+  const Sizes sizes(hipArrayDefault);
+  const size_t s = 64;
+
+  hiparray array{};
+  SECTION("Happy") {
+    SECTION("1D - Max") {
+      desc.Width = sizes.max1D;
+      desc.Height = 0;
+    }
+    SECTION("2D - Max Width") {
+      desc.Width = sizes.max2D[0];
+      desc.Height = s;
+    }
+    SECTION("2D - Max Height") {
+      desc.Width = s;
+      desc.Height = sizes.max2D[1];
+    }
+    SECTION("2D - Max Width and Height") {
+      desc.Width = sizes.max2D[0];
+      desc.Height = sizes.max2D[1];
+    }
+    auto maxArrayCreateError = hipArrayCreate(&array, &desc);
+    // this can try to alloc many GB of memory, so out of memory is acceptable
+    // return to avoid destroy
+    if (maxArrayCreateError == hipErrorOutOfMemory) return;
+    HIP_CHECK(maxArrayCreateError);
+    HIP_CHECK(hipArrayDestroy(array));
+  }
+  SECTION("Negative") {
+    SECTION("1D - More Than Max") {
+      desc.Width = sizes.max1D + 1;
+      desc.Height = 0;
+    }
+    SECTION("2D - More Than Max Width") {
+      desc.Width = sizes.max2D[0] + 1;
+      desc.Height = s;
+    }
+    SECTION("2D - More Than Max Height") {
+      desc.Width = s;
+      desc.Height = sizes.max2D[1] + 1;
+    }
+    SECTION("2D - More Than Max Width and Height") {
+      desc.Width = sizes.max2D[0] + 1;
+      desc.Height = sizes.max2D[1] + 1;
+    }
+    HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue);
+  }
+}
+
+// zero-width array is not supported
+TEST_CASE("Unit_hipArrayCreate_ZeroWidth") {
+  DriverContext ctx;
+  HIP_ARRAY_DESCRIPTOR desc;
+  desc.Format = driverFormats[0];
+  desc.NumChannels = 4;
+  desc.Width = 0;
+  desc.Height = GENERATE(0, 1024);
+
+  // pointer to the array in device memory
+  hiparray array;
+  HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue);
+}
+
+// HipArrayCreate will return an error when nullptr is used as the array argument
+TEST_CASE("Unit_hipArrayCreate_Nullptr") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-130");
+  return;
+#endif
+  DriverContext ctx;
+  SECTION("Null array") {
+    HIP_ARRAY_DESCRIPTOR desc;
+    desc.Format = driverFormats[0];
+    desc.NumChannels = 4;
+    desc.Width = 1024;
+    desc.Height = 1024;
+
+    HIP_CHECK_ERROR(hipArrayCreate(nullptr, &desc), hipErrorInvalidValue);
+  }
+  SECTION("Null Description") {
+    hiparray array;
+    HIP_CHECK_ERROR(hipArrayCreate(&array, nullptr), hipErrorInvalidValue);
+  }
+}
+
+// Only elements with 1,2, or 4 channels is supported
+TEST_CASE("Unit_hipArrayCreate_BadNumberChannelElement") {
+  DriverContext ctx;
+  HIP_ARRAY_DESCRIPTOR desc;
+  desc.Format = GENERATE(from_range(std::begin(driverFormats), std::end(driverFormats)));
+  desc.NumChannels = GENERATE(-1, 0, 3, 5, 8);
+  desc.Width = 1024;
+  desc.Height = GENERATE(0, 1024);
+
+  hiparray array;
+
+  INFO("Format: " << formatToString(desc.Format) << " NumChannels: " << desc.NumChannels
+                  << " Height: " << desc.Height)
+  HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue);
+}
+
+// Only certain channel formats are acceptable.
+TEST_CASE("Unit_hipArrayCreate_BadChannelFormat") {
+  DriverContext ctx;
+  HIP_ARRAY_DESCRIPTOR desc;
+
+  // create a bad format
+  desc.Format =
+      std::accumulate(std::begin(driverFormats), std::end(driverFormats), driverFormats[0],
+                      [](auto i, auto f) { return static_cast<decltype(desc.Format)>(i + f); });
+  for (auto&& format : driverFormats) {
+    REQUIRE(desc.Format != format);
+  }
+
+  desc.NumChannels = 4;
+  desc.Width = 1024;
+  desc.Height = GENERATE(0, 1024);
+
+  hiparray array;
+
+  INFO("Format: " << formatToString(desc.Format) << " Height: " << desc.Height)
+  HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue);
+}

tests/catch/unit/memory/hipHostGetFlags.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
@@ -27,8 +27,34 @@ This testcase verifies the basic scenario of hipHostGetFlags API
 #include <hip_test_common.hh>
 #include <hip_test_checkers.hh>
 #include <hip_test_kernels.hh>
+#include <thread>
+#include <vector>
 
-static constexpr auto LEN{1024*1024};
+std::vector<unsigned int> FlagPart1Vec{hipHostMallocDefault,
+                                       hipHostMallocDefault | hipHostMallocPortable,
+                                       hipHostMallocDefault | hipHostMallocMapped,
+                                       hipHostMallocDefault | hipHostMallocWriteCombined,
+                                       hipHostMallocPortable,
+                                       hipHostMallocPortable | hipHostMallocMapped,
+                                       hipHostMallocPortable | hipHostMallocWriteCombined,
+                                       hipHostMallocMapped,
+                                       hipHostMallocMapped | hipHostMallocWriteCombined,
+                                       hipHostMallocWriteCombined};
+#if HT_AMD
+// For cases where flags from FlagPart1Vec are not used,
+// hipHostMallocDefault is the default on AMD
+// and hipHostMallocMapped on Nvidia
+std::vector<unsigned int> FlagPart2Vec{0x0,
+                                       hipHostMallocNumaUser,
+                                       hipHostMallocNumaUser | hipHostMallocCoherent,
+                                       hipHostMallocNumaUser | hipHostMallocNonCoherent,
+                                       hipHostMallocCoherent,
+                                       hipHostMallocNonCoherent};
+#else
+std::vector<unsigned int> FlagPart2Vec{0x0};
+#endif
+
+static constexpr auto LEN{1024 * 1024};
 
 /*
 This testcase verifies hipHostGetFlags API basic scenario
@@ -38,57 +64,160 @@ This testcase verifies hipHostGetFlags API basic scenario
 3. Validates it with the initial flags used while allocating
    memory
 */
-TEMPLATE_TEST_CASE("Unit_hipHostGetFlags_Basic", "", int,
-                    float, double) {
-  constexpr auto SIZE{LEN * sizeof(TestType)};
-  TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
-  TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
-  unsigned int FlagA, FlagB, FlagC;
 
-  FlagA = hipHostMallocWriteCombined | hipHostMallocMapped;
-  FlagB = hipHostMallocWriteCombined | hipHostMallocMapped;
-  FlagC = hipHostMallocMapped;
+/* Possible host flags
+ * hipHostMallocDefault 0x0
+ * hipHostMallocPortable 0x1
+ * hipHostMallocMapped 0x2
+ * hipHostMallocWriteCombined 0x4
+ * NOT on Nvidia
+ * hipHostMallocNumaUser  0x20000000
+ * hipHostMallocCoherent  0x40000000
+ * hipHostMallocNonCoherent  0x80000000
+ */
+
+inline void checkFlags(unsigned int expected, unsigned int obtained) {
+  // Account for cases where flags from FlagPart1Vec do not include hipHostMallocMapped,
+  // on Nvidia devices it is added by default
+#if HT_NVIDIA
+  expected = expected | hipHostMallocMapped;
+#endif
+  REQUIRE(expected == obtained);
+}
+
+TEST_CASE("Unit_hipHostGetFlags_flagCombos") {
+
+  constexpr auto SIZE{LEN * sizeof(int)};
+  int* A_h{nullptr};
+
+  const unsigned int FlagPart1 = GENERATE(from_range(FlagPart1Vec.begin(), FlagPart1Vec.end()));
+  const unsigned int FlagPart2 = GENERATE(from_range(FlagPart2Vec.begin(), FlagPart2Vec.end()));
+
+  unsigned int FlagComp = FlagPart1 | FlagPart2;
 
   hipDeviceProp_t prop;
-  int device;
+  int device{};
+  HIP_CHECK(hipGetDevice(&device));
+  HIP_CHECK(hipGetDeviceProperties(&prop, device));
+
+  // Skip test if device does not support the property canMapHostMemory
+  if (prop.canMapHostMemory != 1) {
+    HipTest::HIP_SKIP_TEST("Device Property canMapHostMemory is not set");
+    return;
+  } else {
+    // Allocate using the generated flags combos
+    INFO("Flag passed when allocating: 0x" << std::hex << FlagComp << "\n");
+    HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), SIZE, FlagComp));
+    unsigned int flagA{};
+
+    // get the flags from allocations and check if they are the same as the one set
+    HIP_CHECK(hipHostGetFlags(&flagA, A_h));
+
+    checkFlags(FlagComp, flagA);
+    HIP_CHECK(hipHostFree(A_h));
+  }
+}
+
+// Test Allocation with flags and getting flags in another thread
+TEST_CASE("Unit_hipHostGetFlags_DifferentThreads") {
+  constexpr auto SIZE{LEN * sizeof(int)};
+  int* A_h{nullptr};
+
+  const unsigned int FlagPart1 = GENERATE(from_range(FlagPart1Vec.begin(), FlagPart1Vec.end()));
+  const unsigned int FlagPart2 = GENERATE(from_range(FlagPart2Vec.begin(), FlagPart2Vec.end()));
+
+
+  unsigned int FlagComp = FlagPart1 | FlagPart2;
+
+  hipDeviceProp_t prop;
+  int device{};
   HIP_CHECK(hipGetDevice(&device));
   HIP_CHECK(hipGetDeviceProperties(&prop, device));
   if (prop.canMapHostMemory != 1) {
-    SUCCEED("Device Property canMapHostMemory is not set");
+    HipTest::HIP_SKIP_TEST("Device Property canMapHostMemory is not set");
+    return;
   } else {
-    HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), SIZE,
-                            hipHostMallocWriteCombined | hipHostMallocMapped));
-    HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&B_h), SIZE,
-                            hipHostMallocWriteCombined | hipHostMallocMapped));
-    HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&C_h), SIZE,
-                            hipHostMallocMapped));
-
-    unsigned int flagA, flagB, flagC;
-    HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d), A_h, 0));
-    HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&B_d), B_h, 0));
-    HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&C_d), C_h, 0));
+    // Make sure we allocate before trying to get the flags
+    std::thread malloc_thread(
+        [&]() { HIP_CHECK_THREAD(hipHostMalloc(reinterpret_cast<void**>(&A_h), SIZE, FlagComp)); });
+    malloc_thread.join();
+    HIP_CHECK_THREAD_FINALIZE();
+    unsigned int flagA{};
     HIP_CHECK(hipHostGetFlags(&flagA, A_h));
-    HIP_CHECK(hipHostGetFlags(&flagB, B_h));
-    HIP_CHECK(hipHostGetFlags(&flagC, C_h));
 
-    HipTest::setDefaultData<TestType>(LEN, A_h, B_h, C_h);
+    checkFlags(FlagComp, flagA);
 
-    dim3 dimGrid(LEN / 512, 1, 1);
-    dim3 dimBlock(512, 1, 1);
-    hipLaunchKernelGGL(HipTest::vectorADD, dimGrid, dimBlock,
-                       0, 0, static_cast<const TestType*>(A_d),
-                       static_cast<const TestType*>(B_d), C_d, LEN);
-
-    HIP_CHECK(hipMemcpy(C_h, C_d, SIZE, hipMemcpyDeviceToHost));
-    // Note this really HostToHost not
-    // DeviceToHost, since memory is mapped...
-    HipTest::checkVectorADD(A_h, B_h, C_h, LEN);
-
-    REQUIRE(flagA == FlagA);
-    REQUIRE(flagB == FlagB);
-    REQUIRE(flagC == FlagC);
     HIP_CHECK(hipHostFree(A_h));
-    HIP_CHECK(hipHostFree(B_h));
-    HIP_CHECK(hipHostFree(C_h));
+  }
+}
+
+// Test behaviour of hipHostGetFlags with invalid args
+TEST_CASE("Unit_hipHostGetFlags_InvalidArgs") {
+  constexpr auto SIZE{LEN * sizeof(int)};
+  int* A_h{nullptr};
+
+  hipDeviceProp_t prop;
+  int device{};
+  HIP_CHECK(hipGetDevice(&device));
+  HIP_CHECK(hipGetDeviceProperties(&prop, device));
+
+  // Skip test if device does not support the property canMapHostMemory
+  if (prop.canMapHostMemory != 1) {
+    HipTest::HIP_SKIP_TEST("Device Property canMapHostMemory is not set");
+    return;
+  } else {
+    SECTION("Invalid flag ptr being passed to hipHostGetFlags") {
+      // Use default flag
+      unsigned int FlagComp = 0x0;
+
+      // Allocate using the generated flags combos
+      HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), SIZE, FlagComp));
+
+      // use a nullptr to return flags to
+      unsigned int* flagA = nullptr;
+
+      // get the flags from allocations and check if they are the same as the one set
+      HIP_CHECK_ERROR(hipHostGetFlags(flagA, A_h), hipErrorInvalidValue);
+
+      HIP_CHECK(hipHostFree(A_h));
+    }
+
+    SECTION("Device ptr allocated with hipMalloc passed to hipHostGetFlags") {
+      unsigned int FlagComp = 0x4;
+
+      // Allocate memory on device
+      HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_h), SIZE));
+
+      unsigned int flagA{};
+
+      // get the flags from allocations and check if they are the same as the one set
+      HIP_CHECK_ERROR(hipHostGetFlags(&flagA, A_h), hipErrorInvalidValue);
+      INFO("Flag passed when allocating: " << std::hex << FlagComp << " Returned flag: " << std::hex
+                                           << flagA << "\n");
+
+      HIP_CHECK(hipFree(A_h));
+    }
+
+    SECTION("Ptr from hipHostGetDevicePointer passed to hipHostGetFlags") {
+      unsigned int FlagComp = 0x4;
+
+      int* A_d{nullptr};
+      // Allocate memory on device
+      HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), SIZE, FlagComp));
+      HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d), A_h, 0));
+
+      unsigned int flagA;
+
+      // get the flags from allocations and check if they are the same as the one set
+      HIP_CHECK(hipHostGetFlags(&flagA, A_d));
+      INFO("Flag passed when allocating: " << std::hex << FlagComp << " Returned flag: " << std::hex
+                                           << flagA << "\n");
+#if HT_NVIDIA
+      // on Nvidia adjust for cudaHostAllocMapped being set by default
+      FlagComp = FlagComp | hipHostMallocMapped;
+#endif
+      REQUIRE(flagA == FlagComp);
+      HIP_CHECK(hipHostFree(A_h));
+    }
   }
 }

tests/catch/unit/memory/hipHostRegister.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
@@ -28,51 +28,49 @@ This testfile verifies the following scenarios of hipHostRegister API
 */
 
 #include <hip_test_common.hh>
-#include <hip_test_checkers.hh>
-#include <hip_test_kernels.hh>
+#include <hip_test_helper.hh>
 
 #define OFFSET 128
-static constexpr auto LEN{1024*1024};
+static constexpr auto LEN{1024 * 1024};
 
-template<typename T>
-__global__ void Inc(T* Ad) {
-    int tx = threadIdx.x + blockIdx.x * blockDim.x;
-    Ad[tx] = Ad[tx] + static_cast<T>(1);
+template <typename T> __global__ void Inc(T* Ad) {
+  int tx = threadIdx.x + blockIdx.x * blockDim.x;
+  Ad[tx] = Ad[tx] + static_cast<T>(1);
 }
 
 template <typename T>
-void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd,
-               bool internalRegister) {
-    constexpr auto memsetval = 13.0f;
-    A = A + offset;
-    numElements -= offset;
+void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, bool internalRegister) {
+  constexpr auto memsetval = 13.0f;
+  A = A + offset;
+  numElements -= offset;
 
-    size_t sizeBytes = numElements * sizeof(T);
+  size_t sizeBytes = numElements * sizeof(T);
 
-    if (internalRegister) {
-        HIP_CHECK(hipHostRegister(A, sizeBytes, 0));
-    }
+  if (internalRegister) {
+    HIP_CHECK(hipHostRegister(A, sizeBytes, 0));
+  }
 
-    // Reset
-    for (size_t i = 0; i < numElements; i++) {
-        A[i] = static_cast<float>(i);
-        Bh[i] = 0.0f;
-    }
+  // Reset
+  for (size_t i = 0; i < numElements; i++) {
+    A[i] = static_cast<float>(i);
+    Bh[i] = 0.0f;
+  }
 
-    HIP_CHECK(hipMemset(Bd, memsetval, sizeBytes));
+  HIP_CHECK(hipMemset(Bd, memsetval, sizeBytes));
 
-    HIP_CHECK(hipMemcpy(Bd, A, sizeBytes, hipMemcpyHostToDevice));
-    HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(Bd, A, sizeBytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost));
 
-    // Make sure the copy worked
-    for (size_t i = 0; i < numElements; i++) {
-       REQUIRE(Bh[i] == A[i]);
-    }
+  // Make sure the copy worked
+  for (size_t i = 0; i < numElements; i++) {
+    REQUIRE(Bh[i] == A[i]);
+  }
 
-    if (internalRegister) {
-        HIP_CHECK(hipHostUnregister(A));
-    }
+  if (internalRegister) {
+    HIP_CHECK(hipHostUnregister(A));
+  }
 }
+
 /*
 This testcase verifies the hipHostRegister API by
 1. Allocating the memory using malloc
@@ -81,9 +79,7 @@ This testcase verifies the hipHostRegister API by
 4. Launching kernel and access the device pointer variable
 5. performing hipMemset on the device pointer variable
 */
-TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset",
-                   "", int,
-                   float, double) {
+TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", "", int, float, double) {
   size_t sizeBytes{LEN * sizeof(TestType)};
   TestType *A, **Ad;
   int num_devices;
@@ -118,14 +114,14 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset",
   HIP_CHECK(hipHostUnregister(A));
 
   free(A);
-  delete [] Ad;
+  delete[] Ad;
 }
+
 /*
 This testcase verifies hipHostRegister API by
 performing memcpy on the hipHostRegistered variable.
 */
-TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "",
-                   int, float, double) {
+TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", int, float, double) {
   // 1 refers to hipHostRegister
   // 0 refers to malloc
   auto mem_type = GENERATE(0, 1);
@@ -156,5 +152,49 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "",
 
   free(A);
   free(Bh);
-  hipFree(Bd);
+  HIP_CHECK(hipFree(Bd));
+}
+
+template <typename T> __global__ void fill_kernel(T* dataPtr, T value) {
+  size_t tid{blockIdx.x * blockDim.x + threadIdx.x};
+  dataPtr[tid] = value;
+}
+
+TEMPLATE_TEST_CASE("Unit_hipHostRegister_Negative", "", int, float, double) {
+  TestType* hostPtr = nullptr;
+
+  size_t sizeBytes = 1 * sizeof(TestType);
+  SECTION("hipHostRegister Negative Test - nullptr") {
+    HIP_CHECK_ERROR(hipHostRegister(hostPtr, 1, 0), hipErrorInvalidValue);
+  }
+
+  hostPtr = reinterpret_cast<TestType*>(malloc(sizeBytes));
+  SECTION("hipHostRegister Negative Test - zero size") {
+    HIP_CHECK_ERROR(hipHostRegister(hostPtr, 0, 0), hipErrorInvalidValue);
+  }
+
+#if HT_NVIDIA
+  // Flags aren't used for AMD devices currently
+  SECTION("hipHostRegister Negative Test - invalid flag") {
+    HIP_CHECK_ERROR(hipHostRegister(hostPtr, sizeBytes, 0b11111111), hipErrorInvalidValue);
+  }
+#endif
+
+  size_t devMemAvail{0}, devMemFree{0};
+  HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemAvail));
+  auto hostMemFree = HipTest::getMemoryAmount() /* In MB */ * 1024 * 1024;  // In bytes
+  REQUIRE(devMemFree > 0);
+  REQUIRE(devMemAvail > 0);
+  REQUIRE(hostMemFree > 0);
+
+  size_t memFree = (std::min)(devMemFree, hostMemFree);  // which is the limiter cpu or gpu
+
+  SECTION("hipHostRegister Negative Test - invalid memory size") {
+    HIP_CHECK_ERROR(hipHostRegister(hostPtr, memFree, 0), hipErrorInvalidValue);
+  }
+
+  free(hostPtr);
+  SECTION("hipHostRegister Negative Test - freed memory") {
+    HIP_CHECK_ERROR(hipHostRegister(hostPtr, 0, 0), hipErrorInvalidValue);
+  }
 }

tests/catch/unit/memory/hipHostUnregister.cc

@@ -0,0 +1,103 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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.
+*/
+
+#include <hip_test_common.hh>
+
+namespace hipHostUnregisterTests {
+constexpr unsigned int allFlags = hipHostRegisterDefault &  // 0
+    hipHostRegisterPortable &                               // 1
+    hipHostRegisterMapped &                                 // 2
+    hipHostRegisterIoMemory                                 // 4
+#if HT_NVIDIA
+    & cudaHostRegisterReadOnly;  // 8
+#else
+    ;
+#endif
+
+inline bool hipHostRegisterSupported() {
+#if HT_NVIDIA
+  // unable to query for cudaDevAttrHostRegisterSupported equivalent
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-40");
+  HipTest::HIP_SKIP_TEST("hipHostRegister is not supported on this device");
+  return false;
+#else
+  return true;
+#endif
+}
+
+
+TEST_CASE("Unit_hipHostUnregister_MemoryNotAccessableAfterUnregister") {
+  if (!hipHostRegisterSupported()) {
+    return;
+  }
+  // try all combinations of flags
+  for (unsigned int flag = 0; flag <= allFlags; ++flag) {
+    DYNAMIC_SECTION("Using flag: " << flag) {
+      auto x = std::unique_ptr<int>(new int);
+      HIP_CHECK(hipHostRegister(x.get(), sizeof(int), flag));
+
+      void* device_memory;
+      HIP_CHECK(hipHostGetDevicePointer(&device_memory, x.get(), 0));
+
+      HIP_CHECK(hipHostUnregister(x.get()));
+      HIP_CHECK_ERROR(hipHostGetDevicePointer(&device_memory, x.get(), 0), hipErrorInvalidValue);
+    }
+  }
+}
+
+TEST_CASE("Unit_hipHostUnregister_NullPtr") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-90");
+  return;
+#endif
+  HIP_CHECK_ERROR(hipHostUnregister(nullptr), hipErrorInvalidValue);
+}
+
+TEST_CASE("Unit_hipHostUnregister_NotRegisteredPointer") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("TODO-MATCH-ERRORCODE");
+  return;
+#endif
+  auto x = std::unique_ptr<int>(new int);
+  HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered);
+}
+
+TEST_CASE("Unit_hipHostUnregister_AlreadyUnregisteredPointer") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("TODO-MATCH-ERRORCODE");
+  return;
+#endif
+  if (!hipHostRegisterSupported()) {
+    return;
+  }
+  // try all combinations of flags
+  for (unsigned int flag = 0; flag <= allFlags; ++flag) {
+    DYNAMIC_SECTION("Using flag: " << flag) {
+      auto x = std::unique_ptr<int>(new int);
+      HIP_CHECK(hipHostRegister(x.get(), sizeof(int), flag));
+      HIP_CHECK(hipHostUnregister(x.get()));
+      HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered);
+    }
+  }
+}
+
+}  // namespace hipHostUnregisterTests

tests/catch/unit/memory/hipMalloc3D.cc

@@ -64,26 +64,6 @@ static void Malloc3DThreadFunc(int gpu) {
   MemoryAlloc3DDiffSizes(gpu);
 }
 
-/*
- * This verifies the negative scenarios of hipMalloc3D API
- */
-TEST_CASE("Unit_hipMalloc3D_Negative") {
-  size_t width = SMALL_SIZE * sizeof(char);
-  size_t height{SMALL_SIZE}, depth{SMALL_SIZE};
-  hipPitchedPtr devPitchedPtr;
-
-  SECTION("Passing nullptr to device pitched pointer") {
-    hipExtent extent = make_hipExtent(width, height, depth);
-    REQUIRE(hipMalloc3D(nullptr, extent) != hipSuccess);
-  }
-
-  SECTION("Passing Max values to extent") {
-    hipExtent extent = make_hipExtent(std::numeric_limits<size_t>::max(),
-                                      std::numeric_limits<size_t>::max(),
-                                      std::numeric_limits<size_t>::max());
-    REQUIRE(hipMalloc3D(&devPitchedPtr, extent) != hipSuccess);
-  }
-}
 /*
  * This verifies the hipMalloc3D API by
  * assigning width,height and depth as 10

tests/catch/unit/memory/hipMalloc3DArray.cc

@@ -25,14 +25,15 @@ hipMalloc3DArray API test scenarios
 4. Multithreaded scenario
 */
 
-
-
+#include <array>
 #include <hip_test_common.hh>
+#include "hipArrayCommon.hh"
 
 static constexpr auto ARRAY_SIZE{4};
 static constexpr auto BIG_ARRAY_SIZE{100};
 static constexpr auto ARRAY_LOOP{100};
 
+
 /*
  * This API verifies  memory allocations for small and
  * bigger chunks of data.
@@ -47,126 +48,31 @@ static constexpr auto ARRAY_LOOP{100};
  *
  */
 static void Malloc3DArray_DiffSizes(int gpu) {
-  HIP_CHECK(hipSetDevice(gpu));
-  std::vector<int> array_size;
-  array_size.push_back(ARRAY_SIZE);
-  array_size.push_back(BIG_ARRAY_SIZE);
-  for (auto &size : array_size) {
-    int width{size}, height{size}, depth{size};
-    hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0,
-        0, 0, hipChannelFormatKindFloat);
-    hipArray *arr[ARRAY_LOOP];
-    size_t tot, avail, ptot, pavail;
-    HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
-    for (int i = 0; i < ARRAY_LOOP; i++) {
-      HIP_CHECK(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width,
-              height, depth), hipArrayDefault));
-    }
-    for (int i = 0; i < ARRAY_LOOP; i++) {
-      HIP_CHECK(hipFreeArray(arr[i]));
-    }
-    HIP_CHECK(hipMemGetInfo(&avail, &tot));
-    if ((pavail != avail)) {
-      HIPASSERT(false);
-    }
+  HIP_CHECK_THREAD(hipSetDevice(gpu));
+  const int size = GENERATE(ARRAY_SIZE, BIG_ARRAY_SIZE);
+  int width{size}, height{size}, depth{size};
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
+  std::array<hipArray_t, ARRAY_LOOP> arr;
+  size_t pavail, avail;
+  HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
+
+  for (int i = 0; i < ARRAY_LOOP; i++) {
+    HIP_CHECK_THREAD(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, height, depth),
+                                      hipArrayDefault));
   }
-}
-
-/* Thread Function */
-static void Malloc3DArrayThreadFunc(int gpu) {
-  Malloc3DArray_DiffSizes(gpu);
-}
-
-/*
- * Verifies the negative scenarios of hipMalloc3DArray API
- */
-TEST_CASE("Unit_hipMalloc3DArray_Negative") {
-  constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE};
-  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0,
-                                     0, 0, hipChannelFormatKindFloat);
-  hipArray *arr;
-#if HT_NVIDIA
-  SECTION("NullPointer to Array") {
-    REQUIRE(hipMalloc3DArray(nullptr, &channelDesc, make_hipExtent(width,
-            height, depth), hipArrayDefault) != hipSuccess);
+  for (int i = 0; i < ARRAY_LOOP; i++) {
+    HIP_CHECK_THREAD(hipFreeArray(arr[i]));
   }
 
-  SECTION("NullPointer to Channel Descriptor") {
-    REQUIRE(hipMalloc3DArray(&arr, nullptr, make_hipExtent(width,
-            height, depth), hipArrayDefault) != hipSuccess);
-  }
-#endif
-  SECTION("Width 0 in hipExtent") {
-    REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0,
-            height, width), hipArrayDefault) != hipSuccess);
-  }
-
-  SECTION("Height 0 in hipExtent") {
-    REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width,
-            0, width), hipArrayDefault) != hipSuccess);
-  }
-
-  SECTION("Invalid Flag") {
-    REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width,
-            height, depth), 100) != hipSuccess);
-  }
-
-  SECTION("Width,Height & Depth 0 in hipExtent") {
-    REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0,
-            0, 0), hipArrayDefault) != hipSuccess);
-  }
-
-  SECTION("Max int values to extent") {
-    REQUIRE(hipMalloc3DArray(&arr, &channelDesc,
-            make_hipExtent(std::numeric_limits<int>::max(),
-            std::numeric_limits<int>::max(),
-            std::numeric_limits<int>::max()),
-            hipArrayDefault) != hipSuccess);
-  }
-}
-/*
- * Verifies the extent validation scenarios
- * 1. Passing depth as 0 would create 2D array
- * 2. Passing height and depth as 0 would create 1D array
- * from hipMalloc3DArray API
- */
-TEST_CASE("Unit_hipMalloc3DArray_ExtentValidation") {
-  constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE};
-  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0,
-                                     0, 0, hipChannelFormatKindFloat);
-  hipArray *arr;
-
-  SECTION("Depth 0 in hipExtent") {
-    REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width,
-            height, 0), hipArrayDefault) == hipSuccess);
-    HIP_CHECK(hipFreeArray(arr));
-  }
-
-  SECTION("Height & Depth 0 in hipExtent") {
-    REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width,
-            0, 0), hipArrayDefault) == hipSuccess);
-    HIP_CHECK(hipFreeArray(arr));
-  }
-}
-
-/*
- * Verifies hipMalloc3DArray API by passing width,height
- * and depth as 10
- */
-TEST_CASE("Unit_hipMalloc3DArray_Basic") {
-  constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE};
-  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0,
-                                     0, 0, hipChannelFormatKindFloat);
-  hipArray *arr;
-
-  REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width,
-          height, depth), hipArrayDefault) == hipSuccess);
-  HIP_CHECK(hipFreeArray(arr));
+  HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
+  REQUIRE_THREAD(pavail == avail);
 }
 
 TEST_CASE("Unit_hipMalloc3DArray_DiffSizes") {
   Malloc3DArray_DiffSizes(0);
+  HIP_CHECK_THREAD_FINALIZE();
 }
+
 /*
 This testcase verifies the hipMalloc3DArray API in multithreaded
 scenario by launching threads in parallel on multiple GPUs
@@ -176,16 +82,16 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") {
   std::vector<std::thread> threadlist;
   int devCnt = 0;
   devCnt = HipTest::getDeviceCount();
-  size_t tot, avail, ptot, pavail;
-  HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
+  const auto pavail = getFreeMem();
   for (int i = 0; i < devCnt; i++) {
-    threadlist.push_back(std::thread(Malloc3DArrayThreadFunc, i));
+    threadlist.push_back(std::thread(Malloc3DArray_DiffSizes, i));
   }
 
-  for (auto &t : threadlist) {
+  for (auto& t : threadlist) {
     t.join();
   }
-  HIP_CHECK(hipMemGetInfo(&avail, &tot));
+  HIP_CHECK_THREAD_FINALIZE();
+  const auto avail = getFreeMem();
 
   if (pavail != avail) {
     WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
@@ -193,3 +99,324 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") {
   }
 }
 
+namespace {
+void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected_desc,
+                          const hipExtent& expected_extent, const unsigned int expected_flags) {
+// hipArrayGetInfo doesn't currently exist (EXSWCPHIPT-87)
+#if HT_AMD
+  std::ignore = array;
+  std::ignore = expected_desc;
+  std::ignore = expected_extent;
+  std::ignore = expected_flags;
+#else
+  cudaChannelFormatDesc queried_desc;
+  cudaExtent queried_extent;
+  unsigned int queried_flags;
+
+  cudaArrayGetInfo(&queried_desc, &queried_extent, &queried_flags, array);
+
+  REQUIRE(expected_desc.x == queried_desc.x);
+  REQUIRE(expected_desc.y == queried_desc.y);
+  REQUIRE(expected_desc.z == queried_desc.z);
+  REQUIRE(expected_desc.f == queried_desc.f);
+
+  REQUIRE(expected_extent.width == queried_extent.width);
+  REQUIRE(expected_extent.height == queried_extent.height);
+  REQUIRE(expected_extent.depth == queried_extent.depth);
+
+  REQUIRE(expected_flags == queried_flags);
+#endif
+}
+}  // namespace
+
+TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, short4, float,
+                   float2, float4) {
+  hipArray_t array;
+  const auto desc = hipCreateChannelDesc<TestType>();
+#if HT_AMD
+  const unsigned int flags = hipArrayDefault;
+#else
+  const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore);
+#endif
+  constexpr size_t size = 64;
+  hipExtent extent;
+
+  SECTION("1D Array") { extent = make_hipExtent(size, 0, 0); }
+  SECTION("2D Array") { extent = make_hipExtent(size, size, 0); }
+  SECTION("3D Array") { extent = make_hipExtent(size, size, size); }
+
+  HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags));
+  checkArrayIsExpected(array, desc, extent, flags);
+  HIP_CHECK(hipFreeArray(array));
+}
+
+TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_MaxTexture", "", int, uint4, short, ushort2,
+                   unsigned char, float, float4) {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-97");
+  return;
+#endif
+
+  hipArray_t array;
+  const hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
+#if HT_AMD
+  const unsigned int flag = hipArrayDefault;
+#else
+  const unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore);
+#endif
+  if (flag == hipArraySurfaceLoadStore) {
+    HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58");
+    return;
+  }
+  CAPTURE(flag);
+  const Sizes sizes(flag);
+  CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D);
+
+  const size_t s = 64;
+  SECTION("Happy") {
+    // stored in a vector so some values can be ifdef'd out
+    std::vector<hipExtent> extentsToTest{
+        make_hipExtent(sizes.max1D, 0, 0),                              // 1D max
+        make_hipExtent(sizes.max2D[0], s, 0),                           // 2D max width
+        make_hipExtent(s, sizes.max2D[1], 0),                           // 2D max height
+        make_hipExtent(sizes.max2D[0], sizes.max2D[1], 0),              // 2D max
+        make_hipExtent(sizes.max3D[0], s, s),                           // 3D max width
+        make_hipExtent(s, sizes.max3D[1], s),                           // 3D max height
+        make_hipExtent(s, s, sizes.max3D[2]),                           // 3D max depth
+        make_hipExtent(s, sizes.max3D[1], sizes.max3D[2]),              // 3D max height and depth
+        make_hipExtent(sizes.max3D[0], s, sizes.max3D[2]),              // 3D max width and depth
+        make_hipExtent(sizes.max3D[0], sizes.max3D[1], s),              // 3D max width and height
+        make_hipExtent(sizes.max3D[0], sizes.max3D[1], sizes.max3D[2])  // 3D max
+    };
+    const auto extent =
+        GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest)));
+    CAPTURE(extent.width, extent.height, extent.depth);
+    auto maxArrayCreateError = hipMalloc3DArray(&array, &desc, extent, flag);
+    // this can try to alloc many GB of memory, so out of memory is acceptable
+    if (maxArrayCreateError == hipErrorOutOfMemory) return;
+    HIP_CHECK(maxArrayCreateError);
+    checkArrayIsExpected(array, desc, extent, flag);
+    HIP_CHECK(hipFreeArray(array));
+  }
+  SECTION("Negative") {
+    std::vector<hipExtent> extentsToTest {
+      make_hipExtent(sizes.max1D + 1, 0, 0),                          // 1D max
+          make_hipExtent(sizes.max2D[0] + 1, s, 0),                   // 2D max width
+          make_hipExtent(s, sizes.max2D[1] + 1, 0),                   // 2D max height
+          make_hipExtent(sizes.max2D[0] + 1, sizes.max2D[1] + 1, 0),  // 2D max
+          make_hipExtent(sizes.max3D[0] + 1, s, s),                   // 3D max width
+          make_hipExtent(s, sizes.max3D[1] + 1, s),                   // 3D max height
+#if !HT_NVIDIA                                       // leads to hipSuccess on NVIDIA
+          make_hipExtent(s, s, sizes.max3D[2] + 1),  // 3D max depth
+#endif
+          make_hipExtent(s, sizes.max3D[1] + 1, sizes.max3D[2] + 1),  // 3D max height and depth
+          make_hipExtent(sizes.max3D[0] + 1, s, sizes.max3D[2] + 1),  // 3D max width and depth
+          make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, s),  // 3D max width and height
+          make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, sizes.max3D[2] + 1)  // 3D max
+    };
+    const auto extent =
+        GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest)));
+    CAPTURE(extent.width, extent.height, extent.depth);
+    HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, extent, flag), hipErrorInvalidValue);
+  }
+}
+
+
+#if HT_AMD
+constexpr std::array<unsigned int, 1> validFlags{hipArrayDefault};
+#else
+constexpr std::array<unsigned int, 9> validFlags{
+    hipArrayDefault,
+    hipArrayDefault | hipArraySurfaceLoadStore,
+    hipArrayLayered,
+    hipArrayLayered | hipArraySurfaceLoadStore,
+    hipArrayCubemap,
+    hipArrayCubemap | hipArrayLayered,
+    hipArrayCubemap | hipArraySurfaceLoadStore,
+    hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore,
+    hipArrayTextureGather};
+#endif
+
+hipExtent makeExtent(unsigned int flag, size_t s) {
+  if (flag == hipArrayTextureGather) {
+    return make_hipExtent(s, s, 0);
+  }
+  return make_hipExtent(s, s, s);
+}
+
+
+// Providing the array pointer as nullptr should return an error
+TEST_CASE("Unit_hipMalloc3DArray_Negative_NullArrayPtr") {
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
+  constexpr size_t s = 6;
+
+  const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+  HIP_CHECK_ERROR(hipMalloc3DArray(nullptr, &desc, makeExtent(flag, s), flag),
+                  hipErrorInvalidValue);
+}
+
+// Providing the description pointer as nullptr should return an error
+TEST_CASE("Unit_hipMalloc3DArray_Negative_NullDescPtr") {
+  constexpr size_t s = 6;  // 6 to keep cubemap happy
+  hipArray_t array;
+
+  const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+
+  HIP_CHECK_ERROR(hipMalloc3DArray(&array, nullptr, makeExtent(flag, s), flag),
+                  hipErrorInvalidValue);
+}
+
+// Zero width arrays are not allowed
+TEST_CASE("Unit_hipMalloc3DArray_Negative_ZeroWidth") {
+  constexpr size_t s = 6;  // 6 to keep cubemap happy
+  hipArray_t array;
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
+
+  const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+
+  HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, make_hipExtent(0, s, s), flag),
+                  hipErrorInvalidValue);
+}
+
+// Zero height arrays are only allowed for 1D arrays and layered arrays
+TEST_CASE("Unit_hipMalloc3DArray_Negative_ZeroHeight") {
+  constexpr size_t s = 6;  // 6 to keep cubemap happy
+  hipArray_t array;
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
+  std::array<unsigned int, 2> exceptions{hipArrayLayered,
+                                         hipArrayLayered | hipArraySurfaceLoadStore};
+
+  const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+
+  if (std::find(std::begin(exceptions), std::end(exceptions), flag) == std::end(exceptions)) {
+    // flag is not in list of exceptions
+    HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, make_hipExtent(s, 0, s), flag),
+                    hipErrorInvalidValue);
+  }
+}
+
+TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFlags") {
+  constexpr size_t s = 6;  // 6 to keep cubemap happy
+  hipArray_t array;
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
+
+#if HT_AMD
+  const unsigned int flag = 0xDEADBEEF;
+#else
+  const unsigned int flag =
+      GENERATE(0xDEADBEEF, hipArrayTextureGather | hipArraySurfaceLoadStore,
+               hipArrayTextureGather | hipArrayCubemap,
+               hipArrayTextureGather | hipArraySurfaceLoadStore | hipArrayCubemap);
+#endif
+
+  CAPTURE(flag);
+
+  REQUIRE(std::find(std::begin(validFlags), std::end(validFlags), flag) == std::end(validFlags));
+
+  HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, makeExtent(flag, s), flag), hipErrorInvalidValue);
+}
+
+void testInvalidDescription(hipChannelFormatDesc desc){
+  constexpr size_t s = 6;  // 6 to keep cubemap happy
+  hipArray_t array;
+
+#if HT_NVIDIA
+  hipError_t expectedError = hipErrorUnknown;
+#else
+  hipError_t expectedError = hipErrorInvalidValue;
+#endif
+
+  const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+  HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, makeExtent(flag, s), flag), expectedError);
+}
+
+TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFormat") {
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
+  desc.f = GENERATE(hipChannelFormatKindNone, 0xBEEF);
+  testInvalidDescription(desc);
+}
+
+TEST_CASE("Unit_hipMalloc3DArray_Negative_BadChannelLayout") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129");
+  return;
+#endif
+
+  const int bits = GENERATE(8, 16, 32);
+  const hipChannelFormatKind formatKind =
+      GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat);
+  if (bits == 8 && formatKind == hipChannelFormatKindFloat) return;
+
+
+  hipChannelFormatDesc desc = GENERATE_COPY(hipCreateChannelDesc(bits, bits, bits, 0, formatKind),
+                                            hipCreateChannelDesc(0, bits, bits, 0, formatKind),
+                                            hipCreateChannelDesc(0, bits, bits, bits, formatKind),
+                                            hipCreateChannelDesc(bits, 0, bits, 0, formatKind),
+                                            hipCreateChannelDesc(bits, bits, 0, bits, formatKind),
+                                            hipCreateChannelDesc(0, 0, bits, 0, formatKind),
+                                            hipCreateChannelDesc(0, 0, bits, bits, formatKind));
+
+  INFO("kind: " << channelFormatString(formatKind));
+  INFO("x: " << desc.x << ", y: " << desc.y << ", z: " << desc.z << ", w: " << desc.w);
+
+  testInvalidDescription(desc);
+}
+
+TEST_CASE("Unit_hipMalloc3DArray_Negative_8BitFloat") {
+  hipChannelFormatDesc desc = GENERATE(hipCreateChannelDesc(8, 0, 0, 0, hipChannelFormatKindFloat),
+                                       hipCreateChannelDesc(8, 8, 0, 0, hipChannelFormatKindFloat),
+                                       hipCreateChannelDesc(8, 8, 8, 8, hipChannelFormatKindFloat));
+
+  testInvalidDescription(desc);
+}
+
+TEST_CASE("Unit_hipMalloc3DArray_Negative_DifferentChannelSizes") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129");
+  return;
+#endif
+
+  const int bitsX = GENERATE(8, 16, 32);
+  const int bitsY = GENERATE(8, 16, 32);
+  const int bitsZ = GENERATE(8, 16, 32);
+  const int bitsW = GENERATE(8, 16, 32);
+  if (bitsX == bitsY && bitsY == bitsZ && bitsZ == bitsW) return;  // skip when they are equal
+
+  const hipChannelFormatKind channelFormat =
+      GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat);
+
+  if (channelFormat == hipChannelFormatKindFloat &&
+      (bitsX == 8 || bitsY == 8 || bitsZ == 8 || bitsW == 8))
+    return;  // 8 bit floats aren't allowed
+
+  hipChannelFormatDesc desc = hipCreateChannelDesc(bitsX, bitsY, bitsZ, bitsW, channelFormat);
+
+  INFO("format: " << channelFormatString(channelFormat) << ", x bits: " << bitsX
+                  << ", y bits: " << bitsY << ", z bits: " << bitsZ << ", w bits: " << bitsW);
+
+
+  testInvalidDescription(desc);
+}
+
+TEST_CASE("Unit_hipMalloc3DArray_Negative_BadChannelSize") {
+  const int badBits = GENERATE(-1, 0, 10, 100);
+  const hipChannelFormatKind formatKind =
+      GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat);
+  hipChannelFormatDesc desc = hipCreateChannelDesc(badBits, badBits, badBits, badBits, formatKind);
+
+  INFO("Number of bits: " << badBits);
+
+  testInvalidDescription(desc);
+}
+
+
+// hipMalloc3DArray should handle the max numeric value gracefully.
+TEST_CASE("Unit_hipMalloc3DArray_Negative_NumericLimit") {
+  hipArray_t arrayPtr;
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
+
+  size_t size = std::numeric_limits<size_t>::max();
+  const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+  HIP_CHECK_ERROR(hipMalloc3DArray(&arrayPtr, &desc, makeExtent(flag, size), flag),
+                  hipErrorInvalidValue);
+}

tests/catch/unit/memory/hipMallocArray.cc

@@ -27,15 +27,14 @@ hipMallocArray API test scenarios
 
 #include <hip_test_common.hh>
 #include <limits>
-#if defined(_WIN32) || defined(_WIN64)
 #include <numeric>
-#endif
+#include "hipArrayCommon.hh"
 
-static constexpr auto NUM_W{4};
-static constexpr auto BIGNUM_W{100};
-static constexpr auto BIGNUM_H{100};
-static constexpr auto NUM_H{4};
-static constexpr auto ARRAY_LOOP{100};
+static constexpr size_t NUM_W{4};
+static constexpr size_t NUM_H{4};
+static constexpr size_t BIGNUM_W{100};
+static constexpr size_t BIGNUM_H{100};
+static constexpr int ARRAY_LOOP{100};
 
 /*
  * This API verifies  memory allocations for small and
@@ -51,28 +50,30 @@ static constexpr auto ARRAY_LOOP{100};
  *
  */
 static void MallocArray_DiffSizes(int gpu) {
-  HIP_CHECK(hipSetDevice(gpu));
-  std::vector<std::pair<size_t, size_t>> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}};
-  for (auto& size : array_size) {
-    std::array<hipArray_t, ARRAY_LOOP> A_d;
-    size_t tot, avail, ptot, pavail;
-    hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
-    HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
-    for (int i = 0; i < ARRAY_LOOP; i++) {
-      HIP_CHECK(
-          hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault));
-    }
-    for (int i = 0; i < ARRAY_LOOP; i++) {
-      HIP_CHECK(hipFreeArray(A_d[i]));
-    }
-    HIP_CHECK(hipMemGetInfo(&avail, &tot));
-    if ((pavail != avail)) {
-      HIPASSERT(false);
-    }
+  HIP_CHECK_THREAD(hipSetDevice(gpu));
+  std::pair<size_t, size_t> size =
+      GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H));
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
+  std::array<hipArray_t, ARRAY_LOOP> A_d;
+  size_t pavail, avail;
+  HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
+
+  for (int i = 0; i < ARRAY_LOOP; i++) {
+    HIP_CHECK_THREAD(
+        hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault));
   }
+  for (int i = 0; i < ARRAY_LOOP; i++) {
+    HIP_CHECK_THREAD(hipFreeArray(A_d[i]));
+  }
+
+  HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
+  REQUIRE_THREAD(pavail == avail);
 }
 
-TEST_CASE("Unit_hipMallocArray_DiffSizes") { MallocArray_DiffSizes(0); }
+TEST_CASE("Unit_hipMallocArray_DiffSizes") {
+  MallocArray_DiffSizes(0);
+  HIP_CHECK_THREAD_FINALIZE();
+}
 
 /*
 This testcase verifies the hipMallocArray API in multithreaded
@@ -83,81 +84,25 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") {
   std::vector<std::thread> threadlist;
   int devCnt = 0;
   devCnt = HipTest::getDeviceCount();
-  size_t tot, avail, ptot, pavail;
-  HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
+  const auto pavail = getFreeMem();
   for (int i = 0; i < devCnt; i++) {
-    // TODO the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken.
     threadlist.push_back(std::thread(MallocArray_DiffSizes, i));
   }
 
   for (auto& t : threadlist) {
     t.join();
   }
-  HIP_CHECK(hipMemGetInfo(&avail, &tot));
+  HIP_CHECK_THREAD_FINALIZE();
 
+  const auto avail = getFreeMem();
   if (pavail != avail) {
     WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
     REQUIRE(false);
   }
 }
 
-
-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);
@@ -190,30 +135,6 @@ template <typename T> size_t getAllocSize(const size_t width, const size_t heigh
   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);
-}
-
-const char* channelFormatString(hipChannelFormatKind formatKind) noexcept {
-  switch (formatKind) {
-    case hipChannelFormatKindFloat:
-      return "float";
-    case hipChannelFormatKindSigned:
-      return "signed";
-    case hipChannelFormatKindUnsigned:
-      return "unsigned";
-    default:
-      return "error";
-  }
-}
-
 // Tests /////////////////////////////////////////
 
 // Test the default array by generating a texture from it then reading from that texture.
@@ -458,12 +379,6 @@ void testArrayAsSurface(hipArray_t arrayPtr, const size_t width, const size_t he
   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,
@@ -480,7 +395,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho
   // pointer to the array in device memory
   hipArray_t arrayPtr{};
   size_t width = 1024;
-  size_t height;
+  size_t height{};
 
   SECTION("hipArrayDefault") {
     height = GENERATE(0, 1024);
@@ -522,54 +437,55 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho
 // EXSWCPHIPT-71 - no equivalent value for maxSurface and maxTexture2DGather.
 TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ushort, short2, char,
                    char4, float2, float4) {
-  int device;
-  HIP_CHECK(hipGetDevice(&device));
-  hipDeviceProp_t prop;
-  HIP_CHECK(hipGetDeviceProperties(&prop, device));
   size_t width, height;
   hipArray_t array{};
   hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
   const unsigned int flag = hipArrayDefault;
 
+  const Sizes sizes(flag);
+  CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D);
+
+  const size_t s = 64;
+
   SECTION("Happy") {
     SECTION("1D - Max") {
-      width = prop.maxTexture1D;
+      width = sizes.max1D;
       height = 0;
     }
     SECTION("2D - Max Width") {
-      width = prop.maxTexture2D[0];
-      height = 64;
+      width = sizes.max2D[0];
+      height = s;
     }
     SECTION("2D - Max Height") {
-      width = 64;
-      height = prop.maxTexture2D[1];
+      width = s;
+      height = sizes.max2D[1];
     }
     SECTION("2D - Max Width and Height") {
-      width = prop.maxTexture2D[0];
-      height = prop.maxTexture2D[1];
+      width = sizes.max2D[0];
+      height = sizes.max2D[1];
     }
     auto maxArrayCreateError = hipMallocArray(&array, &desc, width, height, flag);
-    // this can try to alloc many GB of memory, so out of memory is fair
+    // this can try to alloc many GB of memory, so out of memory is acceptable
     if (maxArrayCreateError == hipErrorOutOfMemory) return;
     HIP_CHECK(maxArrayCreateError);
     HIP_CHECK(hipFreeArray(array));
   }
   SECTION("Negative") {
     SECTION("1D - More Than Max") {
-      width = prop.maxTexture1D + 1;
+      width = sizes.max1D + 1;
       height = 0;
     }
     SECTION("2D - More Than Max Width") {
-      width = prop.maxTexture2D[0] + 1;
-      height = 64;
+      width = sizes.max2D[0] + 1;
+      height = s;
     }
     SECTION("2D - More Than Max Height") {
-      width = 64;
-      height = prop.maxTexture2D[1] + 1;
+      width = s;
+      height = sizes.max2D[1] + 1;
     }
     SECTION("2D - More Than Max Width and Height") {
-      width = prop.maxTexture2D[0] + 1;
-      height = prop.maxTexture2D[1] + 1;
+      width = sizes.max2D[0] + 1;
+      height = sizes.max2D[1] + 1;
     }
     HIP_CHECK_ERROR(hipMallocArray(&array, &desc, width, height, flag), hipErrorInvalidValue);
   }
@@ -579,7 +495,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush
 // Arrays with channels of different size are not allowed.
 TEST_CASE("Unit_hipMallocArray_Negative_DifferentChannelSizes") {
 #if HT_AMD
-  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-59");
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129");
   return;
 #endif
   const int bitsX = GENERATE(8, 16, 32);
@@ -634,10 +550,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_ZeroWidth") {
 
 // Providing the array pointer as nullptr should return an error
 TEST_CASE("Unit_hipMallocArray_Negative_NullArrayPtr") {
-#if HT_AMD
-  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-45");
-  return;
-#endif
   hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
 
   HIP_CHECK_ERROR(hipMallocArray(nullptr, &desc, 1024, 0, hipArrayDefault), hipErrorInvalidValue);
@@ -645,10 +557,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_NullArrayPtr") {
 
 // Providing the desc pointer as nullptr should return an error
 TEST_CASE("Unit_hipMallocArray_Negative_NullDescPtr") {
-#if HT_AMD
-  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-83");
-  return;
-#endif
   hipArray_t arrayPtr;
   HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, nullptr, 1024, 0, hipArrayDefault),
                   hipErrorInvalidValue);
@@ -656,10 +564,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_NullDescPtr") {
 
 // Inappropriate but related flags should still return an error
 TEST_CASE("Unit_hipMallocArray_Negative_BadFlags") {
-#if HT_AMD
-  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-72");
-  return;
-#endif
   hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
 
   hipArray_t arrayPtr;
@@ -757,7 +661,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_3ChannelElement") {
 // The bit channel description should not allow any channels after a zero channel
 TEST_CASE("Unit_hipMallocArray_Negative_ChannelAfterZeroChannel") {
 #if HT_AMD
-  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-59");
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129");
   return;
 #endif
   const int bits = GENERATE(8, 16, 32);

tests/catch/unit/memory/hipMallocPitch.cc

@@ -1,13 +1,16 @@
 /*
 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 WARRANTY OF ANY KIND, EXPRESS OR
 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
@@ -17,20 +20,294 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
-/*
-Test Scenarios of hipMallocPitch API
-1. Negative Scenarios
-2. Basic Functionality Scenario
-3. Allocate memory using hipMallocPitch API, Launch Kernel validate result.
-4. Allocate Memory in small chunks and large chunks and check for possible memory leaks
-5. Allocate Memory using hipMallocPitch API, Memcpy2D on the allocated variables.
-6. Multithreaded scenario
-*/
-
 #include <hip_test_common.hh>
+#include <initializer_list>
+#include <memory>
+#include "hip/driver_types.h"
+#include <cstring>
+#include <vector>
+#include <limits>
 #include <hip_test_checkers.hh>
 #include <hip_test_kernels.hh>
 
+/**
+ * @brief Test hipMalloc3D, hipMallocPitch and hipMemAllocPitch with multiple input values.
+ *        Checks that the memory has been allocated with the specified pitch and extent sizes.
+ */
+
+struct MemoryInfo {
+  size_t freeMem;
+  size_t totalMem;
+};
+
+inline static MemoryInfo createMemoryInfo() {
+  MemoryInfo memoryInfo{};
+  HIP_CHECK(hipMemGetInfo(&memoryInfo.freeMem, &memoryInfo.totalMem));
+  return memoryInfo;
+}
+
+static void validateMemory(void* devPtr, hipExtent extent, size_t pitch,
+                           MemoryInfo memBeforeAllocation) {
+  INFO("Width: " << extent.width << " Height: " << extent.height << " Depth: " << extent.depth);
+
+  MemoryInfo memAfterAllocation{createMemoryInfo()};
+  const size_t theoreticalAllocatedMemory{pitch * extent.height * extent.depth};
+  const size_t allocatedMemory = memBeforeAllocation.freeMem - memAfterAllocation.freeMem;
+
+  if (theoreticalAllocatedMemory == 0) {
+    REQUIRE(theoreticalAllocatedMemory == allocatedMemory);
+    return; /* If there was no memory allocated then we don't need to do further checks. */
+  } else {
+    REQUIRE(theoreticalAllocatedMemory <= allocatedMemory);
+  }
+
+  std::unique_ptr<char[]> hostPtr{new char[theoreticalAllocatedMemory]};
+  std::memset(hostPtr.get(), 2, theoreticalAllocatedMemory);
+
+  hipPitchedPtr devPitchedPtr{devPtr, pitch, extent.width, extent.height};
+  hipPitchedPtr hostPitchedPtr{hostPtr.get(), pitch, extent.width, extent.height};
+
+  HIP_CHECK(hipMemset3D(devPitchedPtr, 1, extent));
+
+  hipMemcpy3DParms params{};
+  params.srcPtr = devPitchedPtr;
+  params.kind = hipMemcpyKind::hipMemcpyDeviceToHost;
+  params.dstPtr = hostPitchedPtr;
+  params.extent = extent;
+  HIP_CHECK(hipMemcpy3D(&params))
+
+  bool mismatch = false;
+  for (size_t width = 0; width < extent.width; ++width) {
+    for (size_t height = 0; height < extent.height; ++height) {
+      for (size_t depth = 0; depth < extent.depth; ++depth) {
+        char* reinterpretedPtr = reinterpret_cast<char*>(hostPtr.get());
+        size_t index = (pitch * extent.height * depth) + (pitch * height) + width;
+        if (*(reinterpretedPtr + index) != 1) {
+          mismatch = true;
+        }
+      }
+    }
+  }
+  REQUIRE(!mismatch);
+}
+
+class ExtentGenerator {
+ public:
+  static constexpr size_t totalRandomValues{20};
+  static constexpr size_t seed{1337};
+
+  std::uniform_int_distribution<size_t> width_distribution{1, 1024};
+  std::uniform_int_distribution<size_t> height_distribution{1, 100};
+  std::uniform_int_distribution<size_t> depth_distribution{1, 100};
+
+  std::vector<hipExtent> extents2D{};
+  std::vector<hipExtent> extents3D{};
+
+  static ExtentGenerator& getInstance() {
+    static ExtentGenerator instance;
+    return instance;
+  }
+
+ private:
+  ExtentGenerator() {
+    std::mt19937 randomGenerator{seed};
+    extents3D = std::vector<hipExtent>{hipExtent{0, 0, 0}, hipExtent{1, 0, 0}, hipExtent{0, 1, 0},
+                                       hipExtent{0, 0, 1}};
+
+    for (size_t i = 0; i < totalRandomValues; ++i) {
+      extents3D.push_back(hipExtent{width_distribution(randomGenerator),
+                                    height_distribution(randomGenerator),
+                                    depth_distribution(randomGenerator)});
+    }
+
+    extents2D = std::vector<hipExtent>{hipExtent{0, 0, 1}, hipExtent{1, 0, 1}, hipExtent{0, 1, 1}};
+
+    for (size_t i = 0; i < totalRandomValues; ++i) {
+      extents2D.push_back(
+          hipExtent{width_distribution(randomGenerator), height_distribution(randomGenerator), 1});
+    }
+  }
+};
+
+enum class AllocationApi { hipMalloc3D, hipMallocPitch, hipMemAllocPitch };
+
+hipExtent generateExtent(AllocationApi api) {
+  hipExtent extent;
+  if (api == AllocationApi::hipMalloc3D) {
+    auto& extents3D = ExtentGenerator::getInstance().extents3D;
+    extent = GENERATE_REF(from_range(extents3D.begin(), extents3D.end()));
+  } else {
+    auto& extents2D = ExtentGenerator::getInstance().extents2D;
+    extent = GENERATE_REF(from_range(extents2D.begin(), extents2D.end()));
+  }
+
+  return extent;
+}
+
+
+TEST_CASE("Unit_hipMalloc3D_ValidatePitch") {
+  hipPitchedPtr hipPitchedPtr;
+  hipExtent validExtent{generateExtent(AllocationApi::hipMalloc3D)};
+
+  MemoryInfo memBeforeAllocation{createMemoryInfo()};
+  HIP_CHECK(hipMalloc3D(&hipPitchedPtr, validExtent));
+  validateMemory(hipPitchedPtr.ptr, validExtent, hipPitchedPtr.pitch, memBeforeAllocation);
+  HIP_CHECK(hipFree(hipPitchedPtr.ptr));
+}
+
+TEST_CASE("Unit_hipMemAllocPitch_ValidatePitch") {
+  size_t pitch;
+  hipDeviceptr_t ptr;
+  hipExtent validExtent{generateExtent(AllocationApi::hipMemAllocPitch)};
+  MemoryInfo memBeforeAllocation{createMemoryInfo()};
+  unsigned int elementSizeBytes = GENERATE(4, 8, 16);
+
+  if (validExtent.width == 0 || validExtent.height == 0) {
+    return;
+  }
+
+  HIP_CHECK(
+      hipMemAllocPitch(&ptr, &pitch, validExtent.width, validExtent.height, elementSizeBytes));
+  validateMemory(reinterpret_cast<void*>(ptr), validExtent, pitch, memBeforeAllocation);
+  HIP_CHECK(hipFree(reinterpret_cast<void*>(ptr)));
+}
+
+TEST_CASE("Unit_hipMallocPitch_ValidatePitch") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("TODO-FIX-EXTENT-GENERATOR");
+  return;
+#endif
+  size_t pitch;
+  void* ptr;
+  hipExtent validExtent{generateExtent(AllocationApi::hipMemAllocPitch)};
+  MemoryInfo memBeforeAllocation{createMemoryInfo()};
+  HIP_CHECK(hipMallocPitch(&ptr, &pitch, validExtent.width, validExtent.height));
+  validateMemory(ptr, validExtent, pitch, memBeforeAllocation);
+  HIP_CHECK(hipFree(ptr));
+}
+
+TEST_CASE("Unit_hipMalloc3D_Negative") {
+  SECTION("Invalid ptr") {
+    hipExtent validExtent{1, 1, 1};
+    HIP_CHECK_ERROR(hipMalloc3D(nullptr, validExtent), hipErrorInvalidValue);
+  }
+
+  hipPitchedPtr ptr;
+  constexpr size_t maxSizeT = std::numeric_limits<size_t>::max();
+
+#if HT_NVIDIA
+  //TODO-MATCH-ERRORCODE
+  SECTION("Max size_t width") {
+    hipExtent validExtent{maxSizeT, 1, 1};
+    HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorInvalidValue);
+  }
+#endif
+
+  SECTION("Max size_t height") {
+    hipExtent validExtent{1, maxSizeT, 1};
+    HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorOutOfMemory);
+  }
+
+  SECTION("Max size_t depth") {
+    hipExtent validExtent{1, 1, maxSizeT};
+    HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorOutOfMemory);
+  }
+
+#if HT_NVIDIA
+  //TODO-MATCH-ERRORCODE
+  SECTION("Max size_t all dimensions") {
+    hipExtent validExtent{maxSizeT, maxSizeT, maxSizeT};
+    HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorInvalidValue);
+  }
+#endif
+}
+
+TEST_CASE("Unit_hipMallocPitch_Negative") {
+  size_t pitch;
+  void* ptr;
+  constexpr size_t maxSizeT = std::numeric_limits<size_t>::max();
+
+  SECTION("Invalid ptr") {
+    HIP_CHECK_ERROR(hipMallocPitch(nullptr, &pitch, 1, 1), hipErrorInvalidValue);
+  }
+
+  SECTION("Invalid pitch") {
+    HIP_CHECK_ERROR(hipMallocPitch(&ptr, nullptr, 1, 1), hipErrorInvalidValue);
+  }
+
+#if HT_NVIDIA
+  //TODO-MATCH-ERRORCODE
+  SECTION("Max size_t width") {
+    HIP_CHECK_ERROR(hipMallocPitch(&ptr, &pitch, maxSizeT, 1), hipErrorInvalidValue);
+  }
+#endif
+
+  SECTION("Max size_t height") {
+    HIP_CHECK_ERROR(hipMallocPitch(&ptr, &pitch, 1, maxSizeT), hipErrorOutOfMemory);
+  }
+}
+
+TEST_CASE("Unit_hipMemAllocPitch_Negative") {
+  size_t pitch;
+  hipDeviceptr_t ptr{};
+  unsigned int validElementSizeBytes{4};
+  constexpr size_t maxSizeT = std::numeric_limits<size_t>::max();
+
+#if HT_NVIDIA
+  /* Device synchronize is used here to initialize the device.
+   * Nvidia does not implicitly do it for this Api. And hipInit(0) does not work either.
+   */
+  HIP_CHECK(hipDeviceSynchronize());
+
+  SECTION("Invalid elementSizeBytes") {
+    unsigned int invalidElementSizeBytes = GENERATE(0, 7, 12, 17);
+    HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 1, 1, invalidElementSizeBytes),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Zero width") {
+    HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 0, 1, validElementSizeBytes),
+                    hipErrorInvalidValue);
+  }
+  SECTION("Zero height") {
+    HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 1, 0, validElementSizeBytes),
+                    hipErrorInvalidValue);
+  }
+#endif
+
+  SECTION("Invalid dptr") {
+    HIP_CHECK_ERROR(hipMemAllocPitch(nullptr, &pitch, 1, 1, validElementSizeBytes),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Invalid pitch") {
+    HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, nullptr, 1, 1, validElementSizeBytes),
+                    hipErrorInvalidValue);
+  }
+
+#if HT_NVIDIA
+  //TODO-MATCH-ERRORCODE
+  SECTION("Max size_t width") {
+    HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, maxSizeT, 1, validElementSizeBytes),
+                    hipErrorInvalidValue);
+  }
+#endif
+
+  SECTION("Max size_t height") {
+    HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 1, maxSizeT, validElementSizeBytes),
+                    hipErrorOutOfMemory);
+  }
+}
+
+/*
+Test Scenarios of hipMallocPitch API
+1. Basic Functionality Scenario
+2. Allocate memory using hipMallocPitch API, Launch Kernel validate result.
+3. Allocate Memory in small chunks and large chunks and check for possible memory leaks
+4. Allocate Memory using hipMallocPitch API, Memcpy2D on the allocated variables.
+5. Multithreaded scenario
+*/
+
 static constexpr auto SMALLCHUNK_NUMW{4};
 static constexpr auto SMALLCHUNK_NUMH{4};
 static constexpr auto LARGECHUNK_NUMW{1025};
@@ -295,4 +572,3 @@ TEMPLATE_TEST_CASE("Unit_hipMallocPitch_KernelLaunch", ""
                                 A_h, B_h, C_h, false);
 }
 
-

tests/catch/unit/memory/hipMemVmm.cc

@@ -0,0 +1,85 @@
+/*
+   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.
+ */
+
+/* Test Case Description:
+   1) This testcase verifies the  basic scenario - supported on
+     all devices
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_kernels.hh>
+#include <hip_test_checkers.hh>
+#include <cstdio>
+#include <cstdint>
+#include <algorithm>
+#include <thread>
+#include <chrono>
+#include <vector>
+
+/*
+    This testcase verifies HIP Mem VMM API basic scenario - supported on all devices
+ */
+
+TEST_CASE("Unit_hipMemVmm_Basic") {
+  int vmm = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&vmm, hipDeviceAttributeVirtualMemoryManagementSupported, 0));
+  INFO("hipDeviceAttributeVirtualMemoryManagementSupported: " << vmm);
+
+  if (vmm == 0) {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeVirtualMemoryManagement "
+           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+
+  size_t granularity = 0;
+
+  hipMemAllocationProp memAllocationProp;
+  memAllocationProp.type = hipMemAllocationTypePinned;
+  memAllocationProp.location.id = 0;
+  memAllocationProp.location.type = hipMemLocationTypeDevice;
+
+  HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &memAllocationProp, hipMemAllocationGranularityRecommended));
+
+  size_t size = 4 * 1024;
+  void* reservedAddress{nullptr};
+  HIP_CHECK(hipMemAddressReserve(&reservedAddress, size, granularity, nullptr, 0));
+
+  hipMemGenericAllocationHandle_t gaHandle{nullptr};
+  HIP_CHECK(hipMemCreate(&gaHandle, size, &memAllocationProp, 0));
+
+  HIP_CHECK(hipMemMap(reservedAddress, size, 0, gaHandle, 0));
+
+  hipMemAccessDesc desc;
+  std::vector<char> values(size);
+  const char value = 1;
+
+  HIP_CHECK(hipMemSetAccess(reservedAddress, size, &desc, 1));
+  HIP_CHECK(hipMemset(reservedAddress, value, size));
+  HIP_CHECK(hipMemcpy(&values[0], reservedAddress, size, hipMemcpyDeviceToHost));
+
+  for (size_t i=0; i < size; ++i) {
+    REQUIRE(values[i] == value);
+  }
+
+  HIP_CHECK(hipMemUnmap(reservedAddress, size));
+
+  HIP_CHECK(hipMemRelease(gaHandle));
+  HIP_CHECK(hipMemAddressFree(reservedAddress, size));
+}
+

tests/catch/unit/memory/hipMemsetFunctional.cc

@@ -0,0 +1,560 @@
+/*
+ * 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.
+ */
+
+/**
+ * Testcase Scenarios:
+ * For hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32, hipMemset2D, hipMemset3D and all async
+ * counterparts
+ * 1) (ZeroValue)  - Test setting a specified range to zero.
+ * 2) (SmallSize)  - Test setting a unique memset value for small sizes.
+ * 3) (ZeroSize)   - Test that trying to set memory with a zero dimension does not fail and doesn't
+ *                   affect the memory.
+ * 4) (PartialSet) - Test setting a partial range of total allocated memory and
+ *                   ensure the full range isn't affected.
+ */
+#include <hip_test_common.hh>
+
+constexpr size_t FULL_DIM = 10;
+
+// Enum used to determine which 1D memset function to use.
+enum MemsetType {
+  hipMemsetTypeDefault = 1,
+  hipMemsetTypeD8 = 2,
+  hipMemsetTypeD16 = 3,
+  hipMemsetTypeD32 = 4
+};
+
+// Macro used to assert all elements in a flat vector range is equal to a specified value
+#define HIP_ASSERT_VEC_EQ(ptr, value, N)                                                           \
+  for (size_t i = 0; i < N; i++) {                                                                 \
+    CAPTURE(N, i, ptr[i], value);                                                                  \
+    HIP_ASSERT(ptr[i] == value);                                                                   \
+  }
+
+// Copies device data to host and checks that each element is equal to the
+// specified value
+template <typename T> void check_device_data(T* devPtr, T value, size_t numElems) {
+  std::unique_ptr<T[]> hostPtr(new T[numElems]);
+  HIP_CHECK(hipMemcpy(hostPtr.get(), devPtr, numElems * sizeof(T), hipMemcpyDeviceToHost));
+  HIP_ASSERT_VEC_EQ(hostPtr.get(), value, numElems);
+}
+
+// Macro to assist calling and then checking the result of the 1D memset API with the necessary
+// manipulation to the arguments.
+#define HIP_MEMSET_CHECK(hipMemsetFunc, devPtr, value, count, async)                               \
+  using scalar_t = decltype(value);                                                                \
+  size_t sizeBytes = count * sizeof(scalar_t);                                                     \
+  HIP_CHECK(hipMemsetFunc(devPtr, value, sizeBytes));                                              \
+  if (async) {                                                                                     \
+    HIP_CHECK(hipStreamSynchronize(stream));                                                       \
+  }                                                                                                \
+  check_device_data(devPtr, value, count);
+
+#define HIP_MEMSET_CHECK_DTYPE(hipMemsetFunc, devPtr, value, count, async)                         \
+  HIP_CHECK(hipMemsetFunc(reinterpret_cast<hipDeviceptr_t>(devPtr), value, count));                \
+  if (async) {                                                                                     \
+    HIP_CHECK(hipStreamSynchronize(stream));                                                       \
+  }                                                                                                \
+  check_device_data(devPtr, value, count);
+
+// Enum for specifying wether to allocate the data using hipMalloc, hipHostMalloc or not at all.
+enum MemsetMallocType { hipDeviceMalloc_t = 1, hipHostMalloc_t = 2, hipNoMalloc_t };
+
+// Helper function for allocating memory, setting data with the specified 1D memset API and then
+// checking result of operation.
+template <typename T>
+void checkMemset(T value, size_t count, MemsetType memsetType, bool async = false,
+                 MemsetMallocType mallocType = hipDeviceMalloc_t, T* devPtr = nullptr) {
+  hipStream_t stream{nullptr};
+  if (async) {
+    hipStreamCreate(&stream);
+  }
+
+  // Allocate Memory
+  if (mallocType == hipDeviceMalloc_t) {
+    HIP_CHECK(hipMalloc(&devPtr, count * sizeof(T)));
+  } else if (mallocType == hipHostMalloc_t) {
+    HIP_CHECK(hipHostMalloc(&devPtr, count * sizeof(T)));
+  }
+
+  // memset API calls
+  switch (memsetType) {
+    case hipMemsetTypeDefault:
+      if (!async) {
+        INFO("Testing hipMemset call")
+        HIP_MEMSET_CHECK(hipMemset, devPtr, value, count, false);
+      } else {
+        INFO("Testing hipMemsetAsync call")
+        HIP_MEMSET_CHECK(hipMemsetAsync, devPtr, value, count, true);
+      }
+      break;
+    case hipMemsetTypeD8:
+      if (!async) {
+        INFO("Testing hipMemsetD8 call")
+        HIP_MEMSET_CHECK_DTYPE(hipMemsetD8, devPtr, value, count, false);
+      } else {
+        INFO("Testing hipMemsetD8Async call")
+        HIP_MEMSET_CHECK_DTYPE(hipMemsetD8Async, devPtr, value, count, true);
+      }
+      break;
+    case hipMemsetTypeD16:
+      if (!async) {
+        INFO("Testing hipMemsetD16 call")
+        HIP_MEMSET_CHECK_DTYPE(hipMemsetD16, devPtr, value, count, false);
+      } else {
+        INFO("Testing hipMemsetD16Async call")
+        HIP_MEMSET_CHECK_DTYPE(hipMemsetD16Async, devPtr, value, count, true);
+      }
+      break;
+    case hipMemsetTypeD32:
+      if (!async) {
+        INFO("Testing hipMemsetD32 call")
+        HIP_MEMSET_CHECK_DTYPE(hipMemsetD32, devPtr, value, count, false);
+      } else {
+        INFO("Testing hipMemsetD32Async call")
+        HIP_MEMSET_CHECK_DTYPE(hipMemsetD32Async, devPtr, value, count, true);
+      }
+      break;
+  }
+
+  // Cleanup
+  if (async) {
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+
+  // Free memory
+  if (mallocType == hipDeviceMalloc_t) {
+    HIP_CHECK(hipFree(devPtr));
+  } else if (mallocType == hipHostMalloc_t) {
+    HIP_CHECK(hipHostFree(devPtr));
+  }
+}
+
+// Macro which defines a TEST_CASE which calls and then checks the result of the 1D memset macros
+// for all combinations of sync/async and hipMalloc/hipHostMalloc, given the value and memory range.
+#define DEFINE_1D_BASIC_TEST_CASE(suffix, memsetType, T, value, count)                             \
+  TEST_CASE("Unit_hipMemsetFunctional_" + std::string(suffix)) {                                   \
+    const std::string memsetStr = std::string(suffix);                                             \
+    SECTION(memsetStr + " - Device Malloc") {                                                      \
+      checkMemset(static_cast<T>(value), count, memsetType, false, hipDeviceMalloc_t);             \
+    }                                                                                              \
+    SECTION(memsetStr + " - Host Malloc") {                                                        \
+      checkMemset(static_cast<T>(value), count, memsetType, false, hipHostMalloc_t);               \
+    }                                                                                              \
+    SECTION(memsetStr + "Async - Device Malloc") {                                                 \
+      checkMemset(static_cast<T>(value), count, memsetType, true, hipDeviceMalloc_t);              \
+    }                                                                                              \
+    SECTION(memsetStr + "Async - Host Malloc") {                                                   \
+      checkMemset(static_cast<T>(value), count, memsetType, true, hipHostMalloc_t);                \
+    }                                                                                              \
+  }
+
+DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemset", hipMemsetTypeDefault, float, 0, 1024)
+DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemsetD32", hipMemsetTypeD32, uint32_t, 0, 1024)
+DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemsetD16", hipMemsetTypeD16, int16_t, 0, 1024)
+DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemsetD8", hipMemsetTypeD8, int8_t, 0, 1024)
+
+DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemset", hipMemsetTypeDefault, char, 0x42, 1)
+DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemsetD32", hipMemsetTypeD32, uint32_t, 0x101, 1)
+DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemsetD16", hipMemsetTypeD16, int16_t, 0x10, 1)
+DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemsetD8", hipMemsetTypeD8, int8_t, 0x1, 1)
+
+DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemset", hipMemsetTypeDefault, char, 0x42, 0)
+DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemsetD32", hipMemsetTypeD32, uint32_t, 0x101, 0)
+DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemsetD16", hipMemsetTypeD16, int16_t, 0x10, 0)
+DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemsetD8", hipMemsetTypeD8, int8_t, 0x1, 0)
+
+// Helper function that sets a full region of memory with an initial value, sets a smaller subregion
+// with another value and check that the memset API do not write outside of the subregion of data.
+template <typename T>
+void partialMemsetTest(T valA, T valB, size_t count, size_t offset, MemsetType memsetType,
+                       bool async) {
+  T* devPtr;
+  size_t subSize{count - offset};
+  HIP_CHECK(hipMalloc(&devPtr, count * sizeof(T)));
+
+  // Set entire region to be first value.
+  INFO("Setting full region");
+  checkMemset(valA, count, memsetType, async, hipNoMalloc_t, devPtr);
+
+  // Set partial region to be second value.
+  INFO("Setting partial region");
+  checkMemset(valB, subSize, memsetType, async, hipNoMalloc_t, devPtr + offset);
+
+  // Ensure the first section remains unchanged
+  check_device_data(devPtr, valA, offset);
+  HIP_CHECK(hipFree(devPtr));
+}
+
+TEST_CASE("Unit_hipMemsetFunctional_PartialSet_1D") {
+  for (auto widthOffset = 8; widthOffset <= 8; widthOffset *= 2) {
+    SECTION("hipMemset - Partial Set") {
+      partialMemsetTest<char>(0x1, 0x42, 1024, widthOffset, hipMemsetTypeDefault, false);
+    }
+    SECTION("hipMemsetAsync - Partial Set") {
+      partialMemsetTest<char>(0x1, 0x42, 1024, widthOffset, hipMemsetTypeDefault, true);
+    }
+    SECTION("hipMemsetD8 - Partial Set") {
+      partialMemsetTest<int8_t>(0x1, 0xDE, 1024, widthOffset, hipMemsetTypeD8, false);
+    }
+    SECTION("hipMemsetD8Async - Partial Set") {
+      partialMemsetTest<int8_t>(0x1, 0xDE, 1024, widthOffset, hipMemsetTypeD8, true);
+    }
+    SECTION("hipMemsetD16 - Partial Set") {
+      partialMemsetTest<int16_t>(0x1, 0xDEAD, 1024, widthOffset, hipMemsetTypeD16, false);
+    }
+    SECTION("hipMemsetD16Async - Partial Set") {
+      partialMemsetTest<int16_t>(0x1, 0xDEAD, 1024, widthOffset, hipMemsetTypeD16, true);
+    }
+    SECTION("hipMemsetD32 - Partial Set") {
+      partialMemsetTest<uint32_t>(0x1, 0xDEADBEEF, 1024, widthOffset, hipMemsetTypeD32, false);
+    }
+    SECTION("hipMemsetD32Async - Partial Set") {
+      partialMemsetTest<uint32_t>(0x1, 0xDEADBEEF, 1024, widthOffset, hipMemsetTypeD32, true);
+    }
+  }
+}
+
+// Helper function that copies the device data to the host and returns a unique_ptr to that data.
+template <typename T>
+std::unique_ptr<T[]> get_device_data_2D(T* devPtr, size_t pitch, size_t width, size_t height) {
+  std::unique_ptr<T[]> hostPtr(new T[width * height]);
+  constexpr size_t elementSize = sizeof(T);
+  HIP_CHECK(hipMemcpy2D(hostPtr.get(), width * elementSize, devPtr, pitch, width, height,
+                        hipMemcpyDeviceToHost));
+  return hostPtr;
+}
+
+// Copies device data to host and checks that each element is equal to the
+// specified value
+template <typename T>
+void check_device_data_2D(T* devPtr, T value, size_t pitch, size_t width, size_t height) {
+  auto hostPtr = get_device_data_2D<T>(devPtr, pitch, width, height);
+  HIP_ASSERT_VEC_EQ(hostPtr.get(), value, width * height);
+}
+
+// Helper function for allocating memory, setting data with the specified 2D memset API and then
+// checking result of operation.
+template <typename T>
+void checkMemset2D(T value, size_t width, size_t height, bool async = false, size_t pitch = 0,
+                   T* devPtr = nullptr) {
+  hipStream_t stream{nullptr};
+  hipStreamCreate(&stream);
+  constexpr size_t elementSize = sizeof(T);
+  bool freeDevPtr = false;
+  if (devPtr == nullptr) {
+    freeDevPtr = true;
+    HIP_CHECK(
+        hipMallocPitch(reinterpret_cast<void**>(&devPtr), &pitch, width * elementSize, height));
+  }
+
+  if (!async) {
+    INFO("Testing hipMemset2D call")
+    HIP_CHECK(hipMemset2D(devPtr, pitch, value, width * elementSize, height));
+  } else {
+    INFO("Testing hipMemset2DAsync call")
+    HIP_CHECK(hipMemset2DAsync(devPtr, pitch, value, width * elementSize, height, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+  }
+  if (width * height > 0) {
+    check_device_data_2D(devPtr, value, pitch, width, height);
+  }
+  if (freeDevPtr) {
+    HIP_CHECK(hipFree(devPtr));
+  }
+  hipStreamDestroy(stream);
+}
+
+TEST_CASE("Unit_hipMemsetFunctional_ZeroValue_2D") {
+  constexpr size_t width{128};
+  constexpr size_t height{128};
+  constexpr char memsetVal = 0;
+  SECTION("hipMemset2D - Zero Value") { checkMemset2D(memsetVal, width, height, false); }
+  SECTION("hipMemset2DAsync - Zero Value") { checkMemset2D(memsetVal, width, height, true); }
+}
+
+TEST_CASE("Unit_hipMemsetFunctional_SmallSize_2D") {
+  constexpr char memsetVal = 0x42;
+  SECTION("hipMemset2D - Small Size") { checkMemset2D(memsetVal, 1, 1, false); }
+  SECTION("hipMemset2DAsync - Small Size") { checkMemset2D(memsetVal, 1, 1, true); }
+}
+
+TEST_CASE("Unit_hipMemsetFunctional_ZeroSize_2D") {
+  size_t pitch{0};
+  size_t width{10};
+  size_t height{10};
+  char* devPtr{nullptr};
+  HIP_CHECK(
+      hipMallocPitch(reinterpret_cast<void**>(&devPtr), &pitch, width * sizeof(char), height));
+
+  const char initValue = 0x1;
+  const char testValue = 0x11;
+  // Set full region to initial value
+  checkMemset2D(initValue, width, height, false, pitch, devPtr);
+
+  SECTION("hipMemset2D - Zero Width") {
+    checkMemset2D(testValue, 0, height, false, pitch, devPtr);
+    check_device_data_2D(devPtr, initValue, pitch, width, height);
+  }
+  SECTION("hipMemset2DAsync - Zero Width") {
+    checkMemset2D(testValue, 0, height, true, pitch, devPtr);
+    check_device_data_2D(devPtr, initValue, pitch, width, height);
+  }
+  SECTION("hipMemset2D - Zero Height") {
+    checkMemset2D(testValue, width, 0, false, pitch, devPtr);
+    check_device_data_2D(devPtr, initValue, pitch, width, height);
+  }
+  SECTION("hipMemset2DAsync - Zero Height") {
+    checkMemset2D(testValue, width, 0, true, pitch, devPtr);
+    check_device_data_2D(devPtr, initValue, pitch, width, height);
+  }
+  SECTION("hipMemset2D - Zero Width and Height") {
+    checkMemset2D(testValue, 0, 0, false, pitch, devPtr);
+    check_device_data_2D(devPtr, initValue, pitch, width, height);
+  }
+  SECTION("hipMemset2DAsync - Zero Width and Height") {
+    checkMemset2D(testValue, 0, 0, true, pitch, devPtr);
+    check_device_data_2D(devPtr, initValue, pitch, width, height);
+  }
+  HIP_CHECK(hipFree(devPtr));
+}
+
+// Helper function that sets a full region of memory with an initial value, sets a smaller subregion
+// with another value and check that the memset API do not write outside of the subregion of data.
+template <typename T>
+void partialMemsetTest2D(T valA, T valB, size_t width, size_t height, size_t widthOffset,
+                         size_t heightOffset, bool async) {
+  T* devPtr{nullptr};
+  size_t pitch{0};
+  size_t subWidth{width - widthOffset};
+  size_t subHeight{height - heightOffset};
+  constexpr size_t elementSize = sizeof(T);
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&devPtr), &pitch, width * elementSize, height));
+
+  // Set entire region to be first value.
+  INFO("Setting full square region");
+  checkMemset2D(valA, width, height, async, pitch, devPtr);
+
+  // Set partial region to be second value.
+  INFO("Setting partial square region")
+  checkMemset2D(valB, subWidth, subHeight, async, pitch, devPtr);
+
+  auto hostPtr = get_device_data_2D<T>(devPtr, pitch, width, height);
+  T comparVal{0};
+  size_t idx{0};
+  for (size_t i = 0; i < width; i++) {
+    for (size_t j = 0; j < height; j++) {
+      if (i < subWidth && j < subHeight) {
+        // Compare subregion value
+        comparVal = valB;
+      } else {
+        // Compare full region value
+        comparVal = valA;
+      }
+      idx = i * height + j;
+      CAPTURE(width, height, subWidth, subHeight, i, j, idx, hostPtr[idx], comparVal);
+      HIP_ASSERT(hostPtr[idx] == comparVal);
+    }
+  }
+  HIP_CHECK(hipFree(devPtr));
+}
+
+TEST_CASE("Unit_hipMemsetFunctional_PartialSet_2D") {
+  for (auto widthOffset = 8; widthOffset <= 128; widthOffset *= 2) {
+    for (auto heightOffset = 8; heightOffset <= 128; heightOffset *= 2) {
+      SECTION("hipMemset2D - Partial Set") {
+        partialMemsetTest2D('a', 'b', 200, 200, widthOffset, heightOffset, false);
+      }
+      SECTION("hipMemset2DAsync - Partial Set") {
+        partialMemsetTest2D('a', 'b', 200, 200, widthOffset, heightOffset, true);
+      }
+    }
+  }
+}
+
+// Helper function that copies the device data to the host and returns a unique_ptr to that data.
+template <typename T>
+std::unique_ptr<T[]> get_device_data_3D(hipPitchedPtr& devPitchedPtr, hipExtent extent) {
+  constexpr size_t elementSize = sizeof(T);
+  std::unique_ptr<T[]> hostPtr(
+      new T[devPitchedPtr.pitch * extent.width * extent.height / elementSize]);
+  hipMemcpy3DParms myparms{};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.dstPtr = make_hipPitchedPtr(hostPtr.get(), devPitchedPtr.pitch,
+                                      extent.width / elementSize, extent.height);
+  myparms.srcPtr = devPitchedPtr;
+  myparms.extent = extent;
+  myparms.kind = hipMemcpyDeviceToHost;
+  HIP_CHECK(hipMemcpy3D(&myparms));
+  return hostPtr;
+}
+
+// Copies device data to host and checks that each element is equal to the
+// specified value
+template <typename T>
+void check_device_data_3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent extent) {
+  auto hostPtr = get_device_data_3D<T>(devPitchedPtr, extent);
+  size_t width = extent.width / sizeof(T);
+  size_t height = extent.height;
+  size_t depth = extent.depth;
+  size_t idx;
+  for (size_t k = 0; k < depth; k++) {
+    for (size_t j = 0; j < height; j++) {
+      for (size_t i = 0; i < width; i++) {
+        idx = devPitchedPtr.pitch * height * k + devPitchedPtr.pitch * j + i;
+        INFO("idx=" << idx << " hostPtr[idx]=" << hostPtr[idx] << " value=" << value)
+        HIP_ASSERT(hostPtr[idx] == value);
+      }
+    }
+  }
+}
+
+// Helper function for allocating memory, setting data with the specified 3D memset API and then
+// checking result of operation.
+template <typename T>
+void checkMemset3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent extent, bool async = false) {
+  hipStream_t stream{nullptr};
+  hipStreamCreate(&stream);
+  if (devPitchedPtr.ptr == nullptr) {
+    HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
+  }
+  if (!async) {
+    INFO("Testing hipMemset3D call")
+    HIP_CHECK(hipMemset3D(devPitchedPtr, value, extent));
+  } else {
+    INFO("Testing hipMemset3DAsync call")
+    HIP_CHECK(hipMemset3DAsync(devPitchedPtr, value, extent, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+  }
+  if (extent.width * extent.height * extent.depth > 0) {
+    check_device_data_3D(devPitchedPtr, value, extent);
+  }
+  hipStreamDestroy(stream);
+}
+
+void check_memset_3D(std::string sectionStr, size_t width, size_t height, size_t depth,
+                     char value) {
+  hipPitchedPtr devPitchedPtr;
+  hipExtent fullExtent;
+  constexpr char fullVal = 0x21;
+  hipExtent extent = make_hipExtent(width, height, depth);
+  // Check if any of the dimensions are zero
+  bool anyZero = width * height * depth == 0;
+  if (anyZero) {
+    // If they are zero then set a full region with memset value to later check if it's changed.
+    devPitchedPtr.ptr = nullptr;
+    fullExtent = make_hipExtent(FULL_DIM, FULL_DIM, FULL_DIM);
+    checkMemset3D(devPitchedPtr, fullVal, fullExtent, false);
+  }
+  SECTION("hipMemset3D - " + sectionStr) {
+    if (!anyZero) {
+      devPitchedPtr.ptr = nullptr;
+    }
+    checkMemset3D(devPitchedPtr, value, extent, false);
+    if (anyZero) {
+      // Check to make sure memsets with a zero dimension did not affect above set region.
+      check_device_data_3D(devPitchedPtr, fullVal, fullExtent);
+    }
+    HIP_CHECK(hipFree(devPitchedPtr.ptr));
+  }
+  SECTION("hipMemset3DAsync - " + sectionStr) {
+    if (!anyZero) {
+      devPitchedPtr.ptr = nullptr;
+    }
+    checkMemset3D(devPitchedPtr, value, extent, true);
+    if (anyZero) {
+      // Check to make sure memsets with a zero dimension did not affect above set region.
+      check_device_data_3D(devPitchedPtr, fullVal, fullExtent);
+    }
+    HIP_CHECK(hipFree(devPitchedPtr.ptr));
+  }
+}
+
+TEST_CASE("Unit_hipMemsetFunctional_ZeroValue_3D") {
+  check_memset_3D("Zero Value", 128, 128, 10, 0);
+}
+
+TEST_CASE("Unit_hipMemsetFunctional_SmallSize_3D") { check_memset_3D("Small Size", 1, 1, 1, 0x42); }
+
+TEST_CASE("Unit_hipMemsetFunctional_ZeroSize_3D") {
+  constexpr size_t elementSize = sizeof(char);
+  check_memset_3D("Zero Width", 0, FULL_DIM, FULL_DIM, 0x23);
+  check_memset_3D("Zero Height", FULL_DIM * elementSize, 0, FULL_DIM, 0x23);
+  check_memset_3D("Zero Depth", FULL_DIM * elementSize, FULL_DIM, 0, 0x23);
+  check_memset_3D("Zero Width and Height", 0 * elementSize, 0, FULL_DIM, 0x23);
+  check_memset_3D("Zero Width and Depth", 0 * elementSize, FULL_DIM, 0, 0x23);
+  check_memset_3D("Zero Height and Depth", FULL_DIM * elementSize, 0, 0, 0x23);
+  check_memset_3D("Zero Width, Height and Depth", 0 * elementSize, 0, 0, 0x23);
+}
+
+// Helper function that sets a full region of memory with an initial value, sets a smaller subregion
+// with another value and check that the memset API do not write outside of the subregion of data.
+template <typename T>
+void partialMemsetTest3D(T valA, T valB, size_t width, size_t height, size_t depth,
+                         size_t widthOffset, size_t heightOffset, size_t depthOffset, bool async) {
+  size_t subWidth{width - widthOffset};
+  size_t subHeight{height - heightOffset};
+  size_t subDepth{depth - depthOffset};
+  hipPitchedPtr devPitchedPtr;
+  devPitchedPtr.ptr = nullptr;
+  hipExtent extent = make_hipExtent(width * sizeof(T), height, depth);
+  hipExtent subExtent = make_hipExtent(subWidth * sizeof(T), subHeight, subDepth);
+
+  // Set entire region to be first value.
+  INFO("Setting full cuboid region") { checkMemset3D(devPitchedPtr, valA, extent, async); }
+  // Set partial region to be second value.
+  INFO("Setting partial cuboid region") { checkMemset3D(devPitchedPtr, valB, subExtent, async); }
+  auto pitch = devPitchedPtr.pitch;
+  auto hostPtr = get_device_data_3D<T>(devPitchedPtr, extent);
+  T comparVal{0};
+  size_t idx{0};
+  for (size_t k = 0; k < depth; k++) {
+    for (size_t j = 0; j < height; j++) {
+      for (size_t i = 0; i < width; i++) {
+        if (i < subWidth && j < subHeight && k < subDepth) {
+          comparVal = valB;
+        } else {
+          comparVal = valA;
+        }
+        idx = devPitchedPtr.pitch * height * k + devPitchedPtr.pitch * j + i;
+        CAPTURE(width, height, depth, pitch, subWidth, subHeight, subDepth, i, j, k, idx,
+                hostPtr[idx], comparVal);
+        HIP_ASSERT(hostPtr[idx] == comparVal);
+      }
+    }
+  }
+  HIP_CHECK(hipFree(devPitchedPtr.ptr));
+}
+
+TEST_CASE("Unit_hipMemsetFunctional_PartialSet_3D") {
+  for (auto widthOffset = 8; widthOffset <= 128; widthOffset *= 2) {
+    for (auto heightOffset = 8; heightOffset <= 128; heightOffset *= 2) {
+      for (auto depthOffset = 2; depthOffset <= 5; depthOffset++) {
+        SECTION("hipMemset3D - Partial Set") {
+          partialMemsetTest3D('a', 'b', 200, 200, 10, widthOffset, heightOffset, depthOffset,
+                              false);
+        }
+        SECTION("hipMemset3DAsync - Partial Set") {
+          partialMemsetTest3D('a', 'b', 200, 200, 10, widthOffset, heightOffset, depthOffset, true);
+        }
+      }
+    }
+  }
+}

tests/catch/unit/printf/CMakeLists.txt

@@ -16,3 +16,10 @@ elseif (HIP_PLATFORM MATCHES "nvidia")
                     TEST_TARGET_NAME build_tests
                     COMPILE_OPTIONS -std=c++17)
 endif()
+
+# Standalone exes
+add_executable(printfFlags EXCLUDE_FROM_ALL printfFlags_exe.cc)
+add_executable(printfSpecifiers EXCLUDE_FROM_ALL printfSpecifiers_exe.cc)
+
+add_dependencies(printfTests printfFlags)
+add_dependencies(printfTests printfSpecifiers)

tests/catch/unit/printf/printfFlags.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
@@ -23,20 +23,6 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip_test_process.hh>
 
-__global__ void test_kernel() {
-  printf("%08d\n", 42);
-  printf("%08i\n", -42);
-  printf("%08u\n", 42);
-  printf("%08g\n", 123.456);
-  printf("%0+8d\n", 42);
-  printf("%+d\n", -42);
-  printf("%+08d\n", 42);
-  printf("%-8s\n", "xyzzy");
-  printf("% i\n", -42);
-  printf("%-16.8d\n", 42);
-  printf("%16.8d\n", 42);
-}
-
 TEST_CASE("Unit_printf_flags") {
   std::string reference(R"here(00000042
 -0000042
@@ -51,7 +37,7 @@ xyzzy
         00000042
 )here");
 
-  hip::SpawnProc proc("printfExe/printfFlags", true);
+  hip::SpawnProc proc("printfFlags", true);
   REQUIRE(proc.run() == 0);
   REQUIRE(proc.getOutput() == reference);
 }

tests/catch/unit/printf/printfFlags_exe.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

tests/catch/unit/printf/printfSpecifiers.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
@@ -89,7 +89,7 @@ x
 )here");
 #endif
 
-  hip::SpawnProc proc("printfExe/printfSepcifiers", true);
+  hip::SpawnProc proc("printfSpecifiers", true);
   REQUIRE(0 == proc.run());
   REQUIRE(proc.getOutput() == reference);
 }

tests/catch/unit/printf/printfSpecifiers_exe.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

tests/catch/unit/printfExe/CMakeLists.txt

@@ -1,5 +0,0 @@
-add_executable(printfFlags EXCLUDE_FROM_ALL printfFlags.cc)
-add_executable(printfSepcifiers EXCLUDE_FROM_ALL printfSepcifiers.cc)
-
-add_dependencies(build_tests printfFlags)
-add_dependencies(build_tests printfSepcifiers)

tests/catch/unit/stream/CMakeLists.txt

@@ -25,9 +25,13 @@ set(TEST_SRC
     hipStreamCreateWithFlags.cc
     hipStreamCreateWithPriority.cc
     hipAPIStreamDisable.cc
+    # hipStreamAttachMemAsync.cc   # Disabling it on nvidia due to issue in function definition of hipStreamAttachMemAsync
+                                   # Fixing would break ABI, to be re-enabled when the fix is made.
     streamCommon.cc
     hipStreamValue.cc
     )
+
+    set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17)
 endif()
 
 hip_add_exe_to_target(NAME StreamTest

tests/catch/unit/stream/hipStreamAttachMemAsync.cc

@@ -0,0 +1,143 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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.
+*/
+
+// TODO Enable it after hipStreamAttachMemAsync is feature complete on HIP
+
+#include <hip_test_common.hh>
+#include <memory>
+
+__device__ __managed__ int var = 0;
+
+enum class StreamAttachTestType { NullStream = 0, StreamPerThread, CreatedStream };
+
+TEST_CASE("Unit_hipStreamAttachMemAsync_Negative") {
+  hipStream_t stream{nullptr};
+
+  auto streamType =
+      GENERATE(StreamAttachTestType::NullStream, StreamAttachTestType::StreamPerThread,
+               StreamAttachTestType::CreatedStream);
+
+  if (streamType == StreamAttachTestType::StreamPerThread) {
+    stream = hipStreamPerThread;
+  } else if (streamType == StreamAttachTestType::CreatedStream) {
+    HIP_CHECK(hipStreamCreate(&stream));
+    REQUIRE(stream != nullptr);
+  }
+
+  SECTION("Invalid Resource Handle") {
+    int definitelyNotAManagedVariable = 0;
+    HIP_CHECK_ERROR(
+        hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(&definitelyNotAManagedVariable),
+                                sizeof(int), hipMemAttachSingle),
+        hipErrorInvalidValue);
+  }
+
+  SECTION("Invalid devptr") {
+    HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream, nullptr, sizeof(int), hipMemAttachSingle),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Invalid Resource Size") {
+    HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(&var), sizeof(int) - 1,
+                                            hipMemAttachSingle),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Invalid Flags") {
+    HIP_CHECK_ERROR(
+        hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(&var), sizeof(int) - 1,
+                                hipMemAttachSingle | hipMemAttachHost | hipMemAttachGlobal),
+        hipErrorInvalidValue);
+  }
+
+  if (streamType == StreamAttachTestType::CreatedStream) {
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+
+__global__ void kernel(int* ptr, size_t size) {
+  auto i = threadIdx.x;
+  if (i < size) {
+    ptr[i] = 1024;
+  }
+}
+
+constexpr size_t size = 1024;
+__device__ __managed__ int m_memory[size];
+
+TEST_CASE("Unit_hipStreamAttachMemAsync_UseCase") {
+  hipStream_t stream{nullptr};
+
+  auto streamType =
+      GENERATE(StreamAttachTestType::NullStream, StreamAttachTestType::StreamPerThread,
+               StreamAttachTestType::CreatedStream);
+
+  if (streamType == StreamAttachTestType::CreatedStream) {
+    HIP_CHECK(hipStreamCreate(&stream));
+    REQUIRE(stream != nullptr);
+  }
+
+  SECTION("Size zero is valid") {
+    int* d_memory{nullptr};
+    HIP_CHECK(hipMallocManaged(&d_memory, sizeof(int) * size, hipMemAttachHost));
+    HIP_CHECK(
+        hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(d_memory), 0, hipMemAttachHost));
+    HIP_CHECK(hipStreamSynchronize(stream));  // Wait for command to complete
+    HIP_CHECK(hipFree(d_memory));
+  }
+
+  SECTION("Access from device and host") {
+    int* d_memory{nullptr};
+
+    HIP_CHECK(hipMallocManaged(&d_memory, sizeof(int) * size, hipMemAttachHost));
+    HIP_CHECK(hipMemset(d_memory, 0, sizeof(int) * size));
+    HIP_CHECK(
+        hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(d_memory), 0, hipMemAttachHost));
+    HIP_CHECK(hipStreamSynchronize(stream));  // Wait for the command to complete
+
+    kernel<<<1, size, 0, stream>>>(d_memory, size);
+    HIP_CHECK(hipStreamSynchronize(stream));  // Wait for the kernel to complete
+
+    auto ptr = std::make_unique<int[]>(size);
+    std::copy(d_memory, d_memory + size, ptr.get());
+
+    HIP_CHECK(hipFree(d_memory));
+
+    REQUIRE(std::all_of(ptr.get(), ptr.get() + size, [](int n) { return n == size; }));
+  }
+
+  SECTION("Access ManagedMemory") {
+    HIP_CHECK(hipMemset(m_memory, 0, sizeof(int) * size));
+    HIP_CHECK(
+        hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(m_memory), 0, hipMemAttachHost));
+    HIP_CHECK(hipStreamSynchronize(stream));  // Wait for the command to complete
+
+    kernel<<<1, size, 0, stream>>>(m_memory, size);
+    HIP_CHECK(hipStreamSynchronize(stream));  // Wait for the kernel to complete
+
+    auto ptr = std::make_unique<int[]>(size);
+    std::copy(m_memory, m_memory + size, ptr.get());
+
+    REQUIRE(std::all_of(ptr.get(), ptr.get() + size, [](int n) { return n == size; }));
+  }
+
+  if (streamType == StreamAttachTestType::CreatedStream) {
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}

tests/catch/unit/streamperthread/CMakeLists.txt

@@ -4,6 +4,7 @@ set(TEST_SRC
   hipStreamPerThread_Event.cc
   hipStreamPerThread_MultiThread.cc
   hipStreamPerThread_DeviceReset.cc
+  hipStreamPerThrdTsts.cc
 )
 
 hip_add_exe_to_target(NAME StreamPerThreadTest

tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc

@@ -0,0 +1,597 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 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.
+*/
+
+
+/* Test Description:
+   Scenario-1: Launch a kernel in hipStreamPerThread, while it is in flight
+   check for hipStreamQuery(hipStreamPerThread) it should return
+   hipErrorNotReady.
+   Scenario-2: Testing hipStreamPerThread stream object with hipMallocManaged()
+   memory
+   Scenario-3: To check the working of hipStreamPerThread in forked process
+   Scenario-4: The following test case tests the working of hipEventSynchronize
+   in multiple threads which are launched in quick succession
+   Scenario-5: The following test case checks the working of
+   hipStreamWaitEvent() with hipStreamWaitEvent()
+   Scenario-6: Testing hipLaunchCooperativeKernel() api with hipStreamPerThread
+   Scenario-7: Testing hipLaunchCooperativeKernelMultiDevice() with
+   hipStreamPerThread
+*/
+#include <vector>
+#include <thread>
+#include <chrono>
+#ifdef _WIN32
+  #include <Windows.h>
+  #define sleep(x) _sleep(x)
+#endif
+#ifdef __linux__
+  #include <unistd.h>
+  #include <sys/mman.h>
+  #include <sys/wait.h>
+#endif
+
+#include <hip_test_common.hh>
+#ifdef HT_AMD
+  #include "hip/hip_cooperative_groups.h"
+#endif
+using namespace std::chrono;
+using namespace cooperative_groups;
+#if HT_AMD
+#define HIPRT_CB
+#endif
+
+
+static bool IfTestPassed = false;
+// kernel
+__global__ void StreamPerThrd(int *Ad, int *Ad1, size_t n, int Pk_Clk,
+                              int Wait, int WaitEvnt  = 0) {
+  size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+  if (index < n) {
+    Ad[index] = Ad[index] + 10;
+  }
+  if (Wait) {
+    int64_t GpuFrq = (Pk_Clk * 1000);
+    int64_t StrtTck = clock64();
+    if (index == 0) {
+      // The following while loop checks the value in ptr for around 4 seconds
+      while ((clock64() - StrtTck) <= (6 * GpuFrq)) {
+      }
+      if (WaitEvnt == 1) {
+        *Ad1 = 1;
+      }
+    }
+  }
+}
+
+
+__global__ void StreamPerThrd1(int *A, int Pk_Clk) {
+  int64_t GpuFrq = (Pk_Clk * 1000);
+  int64_t StrtTck = clock64();
+  // The following while loop checks the value in ptr for around 3-4 seconds
+  while ((clock64() - StrtTck) <= (3 * GpuFrq)) {
+  }
+  *A = 1;
+}
+
+__global__ void MiniKernel(int *A) {
+  if (*A == 0) {
+    *A = 2;  //  Fail condition
+  } else if (*A == 1) {
+    *A = 3;  //  Pass condition
+  } else {
+     *A = 4;  //  Garbage value found in A
+  }
+}
+
+__global__ void StreamPerThrdCoopKrnl(int *Ad, int *n) {
+  int NumElms = (*n);
+  int index = blockIdx.x * blockDim.x + threadIdx.x;
+  if (index < NumElms) {
+    Ad[index] = Ad[index] + 10;
+  }
+}
+
+#if HT_AMD
+__global__ void test_gwsPerThrd(uint* buf, uint bufSize, int64_t* tmpBuf,
+                                int64_t* result) {
+    extern __shared__ int64_t tmp[];
+    uint groups = gridDim.x;
+    uint group_id = blockIdx.x;
+    uint local_id = threadIdx.x;
+    uint chunk = gridDim.x * blockDim.x;
+
+    uint i = group_id * blockDim.x + local_id;
+    int64_t sum = 0;
+    while (i < bufSize) {
+      sum += buf[i];
+      i += chunk;
+    }
+    tmp[local_id] = sum;
+    __syncthreads();
+    i = 0;
+    if (local_id == 0) {
+        sum = 0;
+        while (i < blockDim.x) {
+          sum += tmp[i];
+          i++;
+        }
+        tmpBuf[group_id] = sum;
+    }
+
+    // wait
+    cooperative_groups::this_grid().sync();
+
+    if (((blockIdx.x * blockDim.x) + threadIdx.x) == 0) {
+        for (uint i = 1; i < groups; ++i) {
+          sum += tmpBuf[i];
+       }
+       // *result = sum;
+       result[1 + cooperative_groups::this_multi_grid().grid_rank()] = sum;
+    }
+    cooperative_groups::this_multi_grid().sync();
+    if (cooperative_groups::this_multi_grid().grid_rank() == 0) {
+      sum = 0;
+      for (uint i = 1; i <= cooperative_groups::this_multi_grid().num_grids();
+           ++i) {
+        sum += result[i];
+      }
+      *result = sum;
+    }
+}
+#endif
+static const uint BufferSizeInDwords = 256 * 1024 * 1024;
+static constexpr uint NumKernelArgs = 4;
+static constexpr uint MaxGPUs = 8;
+// callback function
+static void HIPRT_CB CallBackFunctn(hipStream_t strm, hipError_t err,
+                                    void *ChkVal) {
+  // The following HIPASSERT() is just to satisfy catch2 framework.
+  // As it ensures the use of all the variables.
+  HIPASSERT(strm);
+  HIPCHECK(err);
+  if (*(reinterpret_cast<int*>(ChkVal)) == 1) {
+    IfTestPassed = true;
+  } else {
+    IfTestPassed = false;
+  }
+}
+
+static void EventSync() {
+  int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST_NUM = 123;
+  int blockSize = 32, peak_clk;
+  HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int)));
+  Ah = new int[NumElms];
+  for (int i = 0; i < NumElms; ++i) {
+    Ah[i] = CONST_NUM;
+  }
+  // creating event objects
+  hipEvent_t start, end;
+  HIP_CHECK(hipEventCreate(&start));
+  HIP_CHECK(hipEventCreate(&end));
+  HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice));
+  HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
+  dim3 dimBlock(blockSize, 1, 1);
+  dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
+  HIP_CHECK(hipEventRecord(start, hipStreamPerThread));
+  StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL, NumElms,
+                                                              peak_clk, 0);
+  HIP_CHECK(hipEventRecord(end, hipStreamPerThread));
+  HIP_CHECK(hipEventSynchronize(end));
+  HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost));
+  int MisMatch = 0;
+  for (int i = 0; i < NumElms; ++i) {
+    if (Ah[i] != (CONST_NUM + 10)) {
+      MisMatch++;
+    }
+  }
+  delete[] Ah;
+  HIP_CHECK(hipFree(Ad));
+  if (MisMatch) {
+    WARN("Data Mismatch observed!!\n");
+    IfTestPassed = false;
+  } else {
+    IfTestPassed = true;
+  }
+}
+
+/* Launch a kernel in hipStreamPerThread, while it is in flight check for
+   hipStreamQuery(hipStreamPerThread) it should return hipErrorNotReady.*/
+TEST_CASE("Unit_hipStreamPerThreadTst_StrmQuery") {
+  int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST_NUM = 123;
+  int blockSize = 32, peak_clk;
+  hipError_t err;
+  HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int)));
+  Ah = new int[NumElms];
+  for (int i = 0; i < NumElms; ++i) {
+    Ah[i] = CONST_NUM;
+  }
+  HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice));
+  HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
+  dim3 dimBlock(blockSize, 1, 1);
+  dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
+  SECTION("Test working of hipStreamQuery") {
+    StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL,
+                    NumElms, peak_clk, 1);
+    err = hipStreamQuery(hipStreamPerThread);
+    if (err != hipErrorNotReady) {
+      WARN("hipStreamQuery on hipStreamPerThread didnt return expected error!");
+      IfTestPassed = false;
+    } else {
+      IfTestPassed = true;
+    }
+  }
+  SECTION("check working of hipStreamAddCallback() with hipStreamPerThread") {
+    int *Hptr = nullptr, *A_d = nullptr;
+    HIP_CHECK(hipHostMalloc(&Hptr, sizeof(int)));
+    *Hptr = 0;
+    HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d), Hptr, 0));
+    StreamPerThrd1<<<1, 1, 0, hipStreamPerThread>>>(A_d, peak_clk);
+    HIP_CHECK(hipStreamAddCallback(hipStreamPerThread, CallBackFunctn, A_d, 0));
+    HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+    HIP_CHECK(hipHostFree(Hptr));
+  }
+  HIP_CHECK(hipFree(Ad));
+  delete[] Ah;
+  REQUIRE(IfTestPassed);
+}
+
+/* Testing hipStreamPerThread stream object with hipMallocManaged() memory*/
+TEST_CASE("Unit_hipStreamPerThread_MangdMem") {
+  int managed = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  if (managed == 1) {
+    int *Hmm = nullptr, NumElms = 4096, CONST_NUM = 123, blockSize = 32;
+    SECTION("Using Managed memory") {
+      HIP_CHECK(hipMallocManaged(&Hmm, NumElms * sizeof(int)));
+      for (int i = 0; i < NumElms; ++i) {
+        Hmm[i] = CONST_NUM;
+      }
+    }
+    SECTION("Prefetching Managed memory to device") {
+      HIP_CHECK(hipMallocManaged(&Hmm, NumElms * sizeof(int)));
+      for (int i = 0; i < NumElms; ++i) {
+        Hmm[i] = CONST_NUM;
+      }
+      HIP_CHECK(hipMemPrefetchAsync(Hmm, NumElms * sizeof(int), 0,
+                hipStreamPerThread));
+    }
+    int peak_clk;
+    HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
+    dim3 dimBlock(blockSize, 1, 1);
+    dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
+    StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Hmm, NULL,
+                    NumElms, peak_clk, 0);
+    HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+    // Validating the result
+    int MisMatch = 0;
+    for (int i = 0; i < NumElms; ++i) {
+      if (Hmm[i] != (CONST_NUM + 10)) {
+        MisMatch++;
+      }
+    }
+    HIP_CHECK(hipFree(Hmm));
+    if (MisMatch) {
+      WARN("Data mismatch observed!!\n");
+      REQUIRE(false);
+    }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+            "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+/*  To check the working of hipStreamPerThread in forked process*/
+#ifdef __linux__
+TEST_CASE("Unit_hipStreamPerThread_ChildProc") {
+  if (fork() == 0) {  //  child process
+    int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST_NUM = 123;
+    int blockSize = 32, peak_clk;
+    HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int)));
+    Ah = new int[NumElms];
+    for (int i = 0; i < NumElms; ++i) {
+      Ah[i] = CONST_NUM;
+    }
+    HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice));
+    HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
+    dim3 dimBlock(blockSize, 1, 1);
+    dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
+    StreamPerThrd<<<dimGrid, dimBlock, 0, hipStreamPerThread>>>(Ad, NULL,
+                    NumElms, peak_clk, 0);
+    HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+    HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost));
+    int MisMatch = 0;
+    for (int i = 0; i < NumElms; ++i) {
+      if (Ah[i] != (CONST_NUM + 10)) {
+        MisMatch++;
+      }
+    }
+    delete[] Ah;
+    HIP_CHECK(hipFree(Ad));
+    if (MisMatch) {
+      WARN("Data Mismatch observed!!\n");
+      exit(9);
+    } else {
+      exit(10);
+    }
+  } else {  //  Parent process
+    int stat;
+    wait(&stat);
+    int Result = WEXITSTATUS(stat);
+    if (Result != 10) {
+      REQUIRE(false);
+    }
+  }
+}
+#endif
+
+/* The following test case tests the working of hipEventSynchronize in
+   multiple threads which are launched in quick succession*/
+TEST_CASE("Unit_hipStreamPerThread_EvtRcrdMThrd") {
+  IfTestPassed = true;
+  int MAX_THREAD_CNT = 20;
+  std::vector<std::thread> threads(MAX_THREAD_CNT);
+  for (auto &th : threads) {
+    th = std::thread(EventSync);
+  }
+  for (auto& th : threads) {
+    th.join();
+  }
+  REQUIRE(IfTestPassed);
+}
+
+/* The following test case checks the working of hipStreamWaitEvent() with
+   hipStreamWaitEvent()*/
+TEST_CASE("Unit_hipStreamPerThread_StrmWaitEvt") {
+  IfTestPassed = true;
+  int *Ad = nullptr, NumElms = 4096, CONST_NUM = 123, blockSize = 32, *Ah = nullptr;
+  int *Ad1 = nullptr, *Ah1 = nullptr;
+  Ah = new int[NumElms];
+  Ah1 = new int;
+  hipStream_t Strm;
+  HIP_CHECK(hipStreamCreate(&Strm));
+  for (int i = 0; i < NumElms; ++i) {
+    Ah[i] = CONST_NUM;
+  }
+  Ah1[0] = 0;
+  HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int)));
+  HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice));
+  memset(Ah, 0, NumElms * sizeof(int));
+  HIP_CHECK(hipMalloc(&Ad1, sizeof(int)));
+  HIP_CHECK(hipMemset(Ad1, 0, sizeof(int)));
+  int peak_clk;
+  HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
+  dim3 dimBlock(blockSize, 1, 1);
+  dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
+  hipEvent_t e1;
+  HIPCHECK(hipEventCreate(&e1));
+  StreamPerThrd<<<dimGrid, dimBlock, 0, Strm>>>(Ad, Ad1, NumElms,
+  peak_clk, 1, 1);
+  HIP_CHECK(hipEventRecord(e1, Strm));
+  HIP_CHECK(hipStreamWaitEvent(hipStreamPerThread, e1, 0 /*flags*/));
+  MiniKernel<<<1, 1, 0, hipStreamPerThread>>>(Ad1);
+  sleep(1);
+  HIP_CHECK(hipMemcpy(Ah1, Ad1, sizeof(int), hipMemcpyDeviceToHost));
+  if (*Ah1  != 3) {
+    IfTestPassed = false;
+    if (*Ah1 == 2) {
+      WARN("hipStreamPerThread didn't honour hipStreamWaitEvent()");
+    } else if (*Ah1 == 4) {
+      WARN("Unexpected behavior observed with hipStreamPerThread");
+    }
+  }
+  // Validating the result
+  HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost));
+  int MisMatch = 0;
+  for (int i = 0; i < NumElms; ++i) {
+    if (Ah[i] != (CONST_NUM + 10)) {
+      MisMatch++;
+    }
+  }
+  HIP_CHECK(hipFree(Ad));
+  HIP_CHECK(hipFree(Ad1));
+  HIP_CHECK(hipStreamDestroy(Strm));
+  delete[] Ah;
+  delete Ah1;
+  if (MisMatch) {
+    WARN("Data mismatch observed!!\n");
+    IfTestPassed = false;
+  }
+  REQUIRE(IfTestPassed);
+}
+
+
+/* Testing hipLaunchCooperativeKernel() api with hipStreamPerThread*/
+TEST_CASE("Unit_hipStreamPerThread_CoopLaunch") {
+  hipDeviceProp_t device_properties;
+  HIPCHECK(hipGetDeviceProperties(&device_properties, 0));
+  /* Test whether target device supports cooperative groups ****************/
+  if (device_properties.cooperativeLaunch == 0) {
+    SUCCEED("Cooperative group support not available...");
+  } else {
+    /* We will launch enough waves to fill up all of the GPU *****************/
+    int warp_size = device_properties.warpSize;
+    int num_sms = device_properties.multiProcessorCount;
+    // long long totalTicks = device_properties.clockRate ;
+    int max_blocks_per_sm = 0;
+    // Calculate the device occupancy to know how many blocks can be run.
+    HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm,
+                                                          StreamPerThrdCoopKrnl,
+                                                          warp_size, 0));
+    int max_active_blocks = max_blocks_per_sm * num_sms;
+    int *Ad = nullptr, *Ah = nullptr, *DNumElms = nullptr, NumElms = 4096;
+    int Const = 123;
+    Ah = new int[NumElms];
+    for (int i = 0; i < NumElms; ++i) {
+      Ah[i] = Const;
+    }
+    HIP_CHECK(hipMalloc(&Ad, sizeof(int) * NumElms));
+    HIP_CHECK(hipMalloc(&DNumElms, sizeof(int)));
+    HIP_CHECK(hipMemcpyAsync(Ad, Ah, sizeof(int) * NumElms,
+                             hipMemcpyHostToDevice, hipStreamPerThread));
+    HIP_CHECK(hipMemcpyAsync(DNumElms, &NumElms, sizeof(int),
+                             hipMemcpyHostToDevice, hipStreamPerThread));
+    HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+
+    void *coop_params[2];
+    coop_params[0] = reinterpret_cast<void*>(&Ad);
+    coop_params[1] = reinterpret_cast<void*>(&DNumElms);
+    HIP_CHECK(hipLaunchCooperativeKernel(
+              reinterpret_cast<void*>(StreamPerThrdCoopKrnl),
+              max_active_blocks, warp_size,
+              coop_params, 0, hipStreamPerThread));
+    HIP_CHECK(hipMemcpy(Ah, Ad, sizeof(int) * NumElms, hipMemcpyDeviceToHost));
+    // Verifying the result
+    int DataMismatch = 0;
+    for (int i = 0; i < NumElms; ++i) {
+      if (Ah[i] != (Const + 10)) {
+        DataMismatch++;
+      }
+    }
+    HIP_CHECK(hipFree(Ad));
+    HIP_CHECK(hipFree(DNumElms));
+    delete[] Ah;
+    if (DataMismatch > 0) {
+      REQUIRE(false);
+    }
+  }
+}
+
+/* Testing hipLaunchCooperativeKernelMultiDevice() with hipStreamPerThread*/
+#if HT_AMD
+TEST_CASE("Unit_hipStreamPerThread_CoopLaunchMDev") {
+  uint* dA[MaxGPUs];
+  int64_t* dB[MaxGPUs];
+  int64_t* dC;
+
+  uint32_t* init = new uint32_t[BufferSizeInDwords];
+  for (uint32_t i = 0; i < BufferSizeInDwords; ++i) {
+    init[i] = i;
+  }
+
+  int nGpu = 0;
+  HIPCHECK(hipGetDeviceCount(&nGpu));
+  size_t copySizeInDwords = BufferSizeInDwords / nGpu;
+  hipDeviceProp_t deviceProp[MaxGPUs];
+
+  for (int i = 0; i < nGpu; i++) {
+    HIPCHECK(hipSetDevice(i));
+
+    // Calculate the device occupancy to know how many blocks can be
+    // run concurrently
+    hipGetDeviceProperties(&deviceProp[i], 0);
+    if (!deviceProp[i].cooperativeMultiDeviceLaunch) {
+      WARN("Device doesn't support cooperative launch!");
+      SUCCEED("");
+    }
+    size_t SIZE = copySizeInDwords * sizeof(uint);
+
+    HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dA[i]), SIZE));
+    HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dB[i]),
+             64 * deviceProp[i].multiProcessorCount * sizeof(int64_t)));
+    if (i == 0) {
+      HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(&dC),
+      (nGpu + 1) * sizeof(int64_t)));
+    }
+    HIPCHECK(hipMemcpy(dA[i], &init[i * copySizeInDwords] , SIZE,
+             hipMemcpyHostToDevice));
+    hipDeviceSynchronize();
+  }
+
+  dim3 dimBlock;
+  dim3 dimGrid;
+  dimGrid.x = 1;
+  dimGrid.y = 1;
+  dimGrid.z = 1;
+  dimBlock.x = 64;
+  dimBlock.y = 1;
+  dimBlock.z = 1;
+
+  int numBlocks = 0;
+  uint workgroups[3] = {64, 128, 256};
+
+  hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu];
+  std::time_t end_time;
+  double time = 0;
+  for (uint set = 0; set < 3; ++set) {
+    void* args[MaxGPUs * NumKernelArgs];
+    WARN("---------- Test#" << set << ", size: "<< BufferSizeInDwords <<
+      " dwords ---------------\n");
+    for (int i = 0; i < nGpu; i++) {
+      HIPCHECK(hipSetDevice(i));
+      dimBlock.x = workgroups[set];
+      HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
+      test_gwsPerThrd, dimBlock.x * dimBlock.y * dimBlock.z,
+      dimBlock.x * sizeof(int64_t)));
+
+      WARN("GPU(" << i << ") Block size: " << dimBlock.x <<
+        " Num blocks per CU: " << numBlocks << "\n");
+
+      dimGrid.x = deviceProp[i].multiProcessorCount * (std::min)(numBlocks, 32);
+
+      args[i * NumKernelArgs]     = reinterpret_cast<void*>(&dA[i]);
+      args[i * NumKernelArgs + 1] = reinterpret_cast<void*>(&copySizeInDwords);
+      args[i * NumKernelArgs + 2] = reinterpret_cast<void*>(&dB[i]);
+      args[i * NumKernelArgs + 3] = reinterpret_cast<void*>(&dC);
+
+      launchParamsList[i].func = reinterpret_cast<void*>(test_gwsPerThrd);
+      launchParamsList[i].gridDim = dimGrid;
+      launchParamsList[i].blockDim = dimBlock;
+      launchParamsList[i].sharedMem = dimBlock.x * sizeof(int64_t);
+
+      launchParamsList[i].stream = hipStreamPerThread;
+      launchParamsList[i].args = &args[i * NumKernelArgs];
+    }
+
+    system_clock::time_point start = system_clock::now();
+    hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0);
+    for (int i = 0; i < nGpu; i++) {
+      HIP_CHECK(hipSetDevice(i));
+      HIP_CHECK(hipDeviceSynchronize());
+    }
+    system_clock::time_point end = system_clock::now();
+    std::chrono::duration<double> elapsed_seconds = end - start;
+    end_time = std::chrono::system_clock::to_time_t(end);
+
+    time += elapsed_seconds.count();
+
+    size_t processedDwords = copySizeInDwords * nGpu;
+    if (*dC != (((int64_t)(processedDwords) * (processedDwords - 1)) / 2)) {
+      WARN("Data validation failed ("<< *dC << " != " <<
+        (((int64_t)(BufferSizeInDwords) * (BufferSizeInDwords - 1)) / 2) <<
+        ") for grid size = " << dimGrid.x << " and block size = " <<
+        dimBlock.x << "\n");
+      WARN("Test failed!");
+    }
+  }
+
+  delete [] launchParamsList;
+
+  WARN("finished computation at " << std::ctime(&end_time));
+  WARN("elapsed time: " << time << "s\n");
+
+  hipSetDevice(0);
+  hipFree(dC);
+  for (int i = 0; i < nGpu; i++) {
+    hipFree(dA[i]);
+    hipFree(dB[i]);
+  }
+  delete [] init;
+}
+#endif

tests/performance/compute/hipPerfDotProduct.cpp

@@ -38,7 +38,7 @@ __global__ void vectors_not_equal(int n,
                                  const double* __restrict__ x,
                                  const double* __restrict__ y,
                                  double* __restrict__ workspace) {
-  int gid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
+  int gid = blockIdx.x * blockDim.x + threadIdx.x;
 
   double sum = 0.0;
   for(int idx = gid; idx < n; idx += hipGridDim_x * hipBlockDim_x) {
@@ -46,51 +46,51 @@ __global__ void vectors_not_equal(int n,
   }
 
   __shared__ double sdata[BLOCKSIZE];
-  sdata[hipThreadIdx_x] = sum;
+  sdata[threadIdx.x] = sum;
 
   __syncthreads();
 
-  if(hipThreadIdx_x < 128) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 128];
+  if(threadIdx.x < 128) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 128];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 64){
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 64];
+  if(threadIdx.x < 64){
+    sdata[threadIdx.x] += sdata[threadIdx.x + 64];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 32){
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 32];
+  if(threadIdx.x < 32){
+    sdata[threadIdx.x] += sdata[threadIdx.x + 32];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 16) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 16];
+  if(threadIdx.x < 16) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 16];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 8) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 8];
+  if(threadIdx.x < 8) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 8];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 4) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 4];
+  if(threadIdx.x < 4) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 4];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 2) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 2];
+  if(threadIdx.x < 2) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 2];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 1) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 1];
+  if(threadIdx.x < 1) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 1];
   }
 
-  if(hipThreadIdx_x == 0) {
-    workspace[hipBlockIdx_x] = sdata[0];
+  if(threadIdx.x == 0) {
+    workspace[blockIdx.x] = sdata[0];
   }
 
 }
@@ -99,59 +99,59 @@ template <unsigned int BLOCKSIZE>
 __launch_bounds__(BLOCKSIZE)
 __global__ void vectors_equal(int n, const double* __restrict__ x,
                                   double* __restrict__ workspace) {
-  int gid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
+  int gid = blockIdx.x * blockDim.x + threadIdx.x;
 
   double sum = 0.0;
-  for(int idx = gid; idx < n; idx += hipGridDim_x * hipBlockDim_x) {
+  for(int idx = gid; idx < n; idx += hipGridDim_x * blockDim.x) {
     sum = fma(x[idx], x[idx], sum);
   }
 
   __shared__ double sdata[BLOCKSIZE];
-  sdata[hipThreadIdx_x] = sum;
+  sdata[threadIdx.x] = sum;
 
   __syncthreads();
 
-  if(hipThreadIdx_x < 128) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 128];
+  if(threadIdx.x < 128) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 128];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 64) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 64];
+  if(threadIdx.x < 64) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 64];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 32) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 32];
+  if(threadIdx.x < 32) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 32];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 16) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 16];
+  if(threadIdx.x < 16) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 16];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 8) {
-          sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 8];
+  if(threadIdx.x < 8) {
+          sdata[threadIdx.x] += sdata[threadIdx.x + 8];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 4) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 4];
+  if(threadIdx.x < 4) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 4];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 2) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 2];
+  if(threadIdx.x < 2) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 2];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 1) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 1];
+  if(threadIdx.x < 1) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 1];
   }
 
-  if(hipThreadIdx_x == 0) {
-    workspace[hipBlockIdx_x] = sdata[0];
+  if(threadIdx.x == 0) {
+    workspace[blockIdx.x] = sdata[0];
     }
 }
 
@@ -161,49 +161,49 @@ __global__ void dot_reduction(double* __restrict__ workspace) {
 
   __shared__ double sdata[BLOCKSIZE];
 
-  sdata[hipThreadIdx_x] = workspace[hipThreadIdx_x];
+  sdata[threadIdx.x] = workspace[threadIdx.x];
 
   __syncthreads();
 
-  if(hipThreadIdx_x < 128) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 128];
+  if(threadIdx.x < 128) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 128];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 64) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 64];
+  if(threadIdx.x < 64) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 64];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 32) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 32];
+  if(threadIdx.x < 32) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 32];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 16) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 16];
+  if(threadIdx.x < 16) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 16];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 8) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 8];
+  if(threadIdx.x < 8) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 8];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 4) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 4];
+  if(threadIdx.x < 4) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 4];
   } __syncthreads();
 
-  if(hipThreadIdx_x < 2) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 2];
+  if(threadIdx.x < 2) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 2];
   }
   __syncthreads();
 
-  if(hipThreadIdx_x < 1) {
-    sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 1];
+  if(threadIdx.x < 1) {
+    sdata[threadIdx.x] += sdata[threadIdx.x + 1];
   }
 
-  if(hipThreadIdx_x == 0) {
+  if(threadIdx.x == 0) {
     workspace[0] = sdata[0];
   }
 

tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp

@@ -83,12 +83,12 @@ int main(int argc, char* argv[]) {
     defaultCUMask.push_back(temp);
   }
 
-  str_out = hipGetErrorString(hipExtStreamGetCUMask(0, cuMask.size(), 0));
+  str_out = hipGetErrorName(hipExtStreamGetCUMask(0, cuMask.size(), 0));
   if ((str_err.compare(str_out)) != 0) {
     failed("hipExtStreamGetCUMask returned wrong error code!");
   }
 
-  str_out = hipGetErrorString(hipExtStreamGetCUMask(0, 0, &cuMask[0]));
+  str_out = hipGetErrorName(hipExtStreamGetCUMask(0, 0, &cuMask[0]));
   if ((str_err.compare(str_out)) != 0) {
     failed("hipExtStreamGetCUMask returned wrong error code!");
   }