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

Change-Id: Icf46c5a6f83e40caab2d2fcda77320dbbf200c55

.jenkins/Jenkinsfile

@@ -75,6 +75,7 @@ def hipBuildTest(String backendLabel) {
                     cd build
                     # Check if backend label contains string "amd" or backend host is a server with amd gpu
                     if [[ $backendLabel =~ amd ]]; then
+                        export HT_CONFIG_FILE="$HIP_DIR/tests/catch/hipTestMain/config/config_amd_linux.json"
                         cmake -DHIP_CATCH_TEST=1 -DHIP_PATH=\$PWD/install -DHIP_COMMON_DIR=$HIP_DIR -DAMD_OPENCL_PATH=\$OPENCL_DIR -DROCCLR_PATH=\$ROCclr_DIR -DCMAKE_PREFIX_PATH="/opt/rocm/" -DCMAKE_INSTALL_PREFIX=\$PWD/install ..
                     else
                         export HIP_PLATFORM=nvidia
@@ -98,7 +99,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 HT_CONFIG_FILE="$HIP_DIR/tests/catch/hipTestMain/config/config_amd_linux.json" ctest -E 'Unit_hipGraphChildGraphNodeGetGraph_Functional|Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative'
+                            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'
                         else
                             make test
                         fi

bin/hipcc.pl

@@ -155,6 +155,7 @@ if ($HIP_PLATFORM eq "amd") {
     # to avoid using dk linker or MSVC linker
     if($isWindows) {
         $HIPLDFLAGS .= " -fuse-ld=lld";
+        $HIPLDFLAGS .= " --ld-path=$HIP_CLANG_PATH/lld-link.exe";
     }
     $HIP_CLANG_VERSION = `$HIPCC --version`;
     $HIP_CLANG_VERSION=~/.*clang version (\S+).*/;

docs/markdown/hip_faq.md

@@ -142,7 +142,7 @@ The tools also struggle with more complex CUDA applications, in particular, thos
 - For Nvidia platforms, HIP requires Unified Memory and should run on any device supporting CUDA SDK 6.0 or newer. We have tested the Nvidia Titan and Tesla K40.
 
 ### Do HIPIFY tools automatically convert all source code?
-Typically, HIPIFY tools can automatically convert almost all run-time code, and the coordinate indexing device code ( threadIdx.x -> hipThreadIdx_x ).
+Typically, HIPIFY tools can automatically convert almost all run-time code.
 Most device code needs no additional conversion since HIP and CUDA have similar names for math and built-in functions.
 The hipify-clang tool will automatically modify the kernel signature as needed (automating a step that used to be done manually).
 Additional porting may be required to deal with architecture feature queries or with CUDA capabilities that HIP doesn't support.

docs/markdown/hip_kernel_language.md

@@ -500,7 +500,7 @@ Returns the value of counter that is incremented every clock cycle on device. Di
 
 Atomic functions execute as read-modify-write operations residing in global or shared memory. No other device or thread can observe or modify the memory location during an atomic operation. If multiple instructions from different devices or threads target the same memory location, the instructions are serialized in an undefined order.
 
-HIP adds new APIs with _system as suffix to support system scope atomic operations. For example,  atomicAnd atomic is dedicated to the GPU device, atomicAnd_system will allow developers to extend the atomic operation to system scope, from the GPU device to other CPUs and GPU devices in the system.  
+HIP adds new APIs with _system as suffix to support system scope atomic operations. For example, the `atomicAnd` function is meant to be atomic and coherent within the GPU device executing the function. `atomicAnd_system` will allow developers to extend the atomic operation to system scope, from the GPU device to other CPUs and GPU devices in the system.
 
 HIP supports the following atomic operations.
 
@@ -516,6 +516,10 @@ HIP supports the following atomic operations.
 | float atomicAdd_system(float* address, float val)                                                                    |  ✓                    |  ✓                     |
 | double atomicAdd(double* address, double val)                                                                        |  ✓                    |  ✓                     |
 | double atomicAdd_system(double* address, double val)                                                                 |  ✓                    |  ✓                     |
+| float unsafeAtomicAdd(float* address, float val)                                                                     |  ✓                    |  ✗                     |
+| float safeAtomicAdd(float* address, float val)                                                                       |  ✓                    |  ✗                     |
+| double unsafeAtomicAdd(double* address, double val)                                                                  |  ✓                    |  ✗                     |
+| double safeAtomicAdd(double* address, double val)                                                                    |  ✓                    |  ✗                     |
 | int atomicSub(int* address, int val)                                                                                 |  ✓                    |  ✓                     |
 | int atomicSub_system(int* address, int val)                                                                          |  ✓                    |  ✓                     |
 | unsigned int atomicSub(unsigned int* address,unsigned int val)                                                       |  ✓                    |  ✓                     |
@@ -566,12 +570,35 @@ HIP supports the following atomic operations.
 | unsigned long long atomicXor(unsigned long long* address,unsigned long long val))                                    |  ✓                    |  ✓                     |
 | unsigned long long atomicXor_system(unsigned long long* address, unsigned long long val)                             |  ✓                    |  ✓                     |
 
-Note, in order to keep backwards compitability in float/double atomicAdd functions, in ROCm4.4 release, we introduce a new compilation flag as an option in CMake file,
-__HIP_USE_CMPXCHG_FOR_FP_ATOMICS
+### Unsafe Floating-Point Atomic RMW Operations
 
-By default, this compilation flag is not set("0"), so hip runtime will use current float/double atomicAdd functions.
-If this compilation flag is set to "1", that is, with the cmake option "-D__HIP_USE_CMPXCHG_FOR_FP_ATOMICS=1", the old float/double atomicAdd functions will be used instead, for compatibility with compilers not supporting floating point atomics.
-For details steps how to build hip runtime, please refer to the section "build HIPAMD" (https://github.com/ROCm-Developer-Tools/hipamd/blob/develop/INSTALL.md).
+Some HIP devices support fast atomic read-modify-write (RMW) operations on floating-point values.
+For example, `atomicAdd` on single- or double-precision floating-point values may generate a hardware RMW instruction that is faster than emulating the atomic operation using an atomic compare-and-swap (CAS) loop.
+
+On some devices, these fast atomic RMW instructions can produce different results when compared with the same functions implemented with atomic CAS loops.
+For example, some devices will produce incorrect answers if a fast atomic floating-point RMW instruction targets fine-grained memory allocations.
+As another example, some devices will use different rounding or denormal modes when using fast atomic floating-point RMW instructions.
+
+As such, the HIP-Clang compiler offers a compile-time option for users to choose whether their code will use the fast, potentially unsafe, atomic instructions.
+On devices that support these fast, but unsafe, floating-point atomic RMW instructions, the compiler option `-munsafe-fp-atomics` will allow the compiler to generate them when it sees appropriate atomic RMW function calls.
+By passing the `-munsafe-fp-atomics` flag to the compiler, the user is indicating that all floating-point atomic function calls are allowed to use an unsafe version if one exists.
+For instance, on some devices, this flag indicates to the compiler that that no floating-point `atomicAdd` function targets fine-grained memory.
+
+If the user instead compiles with `-mno-unsafe-fp-atomics`, the user is telling the compiler to never use a floating-point atomic RMW that may not be safe.
+The compiler will default to not producing unsafe floating-point atomic RMW instructions, so the `-mno-unsafe-fp-atomics` compilation option is not strictly necessary.
+Explicitly passing this flag to the compiler is good practice, however.
+
+Whenever either of the two options described above, `-munsafe-fp-atomics` and `-mno-unsafe-fp-atomics` are passed to the compiler's command line, they are applied globally for that entire compilation.
+If only a subset of the atomic RMW function calls could safely use the faster floating-point atomic RMW instructions, the developer would instead need to compile with `-mno-unsafe-fp-atomics` in order to ensure the remaining atomic RMW function calls produce correct results.
+Towards this end, HIP has four extra functions to help developers more precisely control which floating-point atomic RMW functions produce unsafe atomic RMW instructions:
+
+- `float unsafeAtomicAdd(float* address, float val)`
+- `double unsafeAtomicAdd(double* address, double val)`
+   - These functions will always produce fast atomic RMW instructions on devices that have them, even when `-mno-unsafe-fp-atomics` is set
+
+- `float safeAtomicAdd(float* address, float val)`
+- `double safeAtomicAdd(double* address, double val)`
+   - These functions will always produce safe atomic RMW operations, even when `-munsafe-fp-atomics` is set
 
 ## Warp Cross-Lane Functions
 

include/hip/hip_runtime_api.h

@@ -165,32 +165,6 @@ typedef enum hipMemoryType {
     hipMemoryTypeUnified  ///< Not used currently
 } hipMemoryType;
 
-/**
- * @brief hipKernelNodeAttrID
- * @enum
- *
- */
-typedef enum hipKernelNodeAttrID {
-    hipKernelNodeAttributeAccessPolicyWindow = 1,
-    hipKernelNodeAttributeCooperative = 2,
-} hipKernelNodeAttrID;
-typedef enum hipAccessProperty {
-    hipAccessPropertyNormal = 0,
-    hipAccessPropertyStreaming  = 1,
-    hipAccessPropertyPersisting = 2,
-} hipAccessProperty;
-typedef struct hipAccessPolicyWindow {
-    void* base_ptr;
-    hipAccessProperty hitProp;
-    float hitRatio;
-    hipAccessProperty missProp;
-    size_t num_bytes;
-} hipAccessPolicyWindow;
-typedef union hipKernelNodeAttrValue {
-    hipAccessPolicyWindow accessPolicyWindow;
-    int cooperative;
-} hipKernelNodeAttrValue;
-
 /**
  * Pointer attributes
  */
@@ -1141,6 +1115,32 @@ typedef struct hipMemsetParams {
   size_t width;
 } hipMemsetParams;
 
+/**
+ * @brief hipKernelNodeAttrID
+ * @enum
+ *
+ */
+typedef enum hipKernelNodeAttrID {
+    hipKernelNodeAttributeAccessPolicyWindow = 1,
+    hipKernelNodeAttributeCooperative = 2,
+} hipKernelNodeAttrID;
+typedef enum hipAccessProperty {
+    hipAccessPropertyNormal = 0,
+    hipAccessPropertyStreaming  = 1,
+    hipAccessPropertyPersisting = 2,
+} hipAccessProperty;
+typedef struct hipAccessPolicyWindow {
+    void* base_ptr;
+    hipAccessProperty hitProp;
+    float hitRatio;
+    hipAccessProperty missProp;
+    size_t num_bytes;
+} hipAccessPolicyWindow;
+typedef union hipKernelNodeAttrValue {
+    hipAccessPolicyWindow accessPolicyWindow;
+    int cooperative;
+} hipKernelNodeAttrValue;
+
 /**
  * @brief hipGraphExecUpdateResult
  * @enum
@@ -1178,6 +1178,13 @@ typedef enum hipStreamUpdateCaptureDependenciesFlags {
   hipStreamSetCaptureDependencies,      ///< Replace the dependency set with the new nodes
 } hipStreamUpdateCaptureDependenciesFlags;
 
+typedef enum hipGraphMemAttributeType {
+  hipGraphMemAttrUsedMemCurrent = 0, ///< Amount of memory, in bytes, currently associated with graphs
+  hipGraphMemAttrUsedMemHigh,        ///< High watermark of memory, in bytes, associated with graphs since the last time.
+  hipGraphMemAttrReservedMemCurrent, ///< Amount of memory, in bytes, currently allocated for graphs.
+  hipGraphMemAttrReservedMemHigh,    ///< High watermark of memory, in bytes, currently allocated for graphs
+}hipGraphMemAttributeType;
+
 typedef enum hipGraphInstantiateFlags {
   hipGraphInstantiateFlagAutoFreeOnLaunch =
       1,  ///< Automatically free memory allocated in a graph before relaunching.
@@ -4904,6 +4911,7 @@ hipError_t hipBindTextureToMipmappedArray(
  * @returns hipSuccess, hipErrorInvalidValue
  *
  */
+DEPRECATED(DEPRECATED_MSG)
  hipError_t hipGetTextureReference(
     const textureReference** texref,
     const void* symbol);
@@ -4990,20 +4998,25 @@ hipError_t hipGetTextureObjectTextureDesc(
 /**
  *
  */
+DEPRECATED(DEPRECATED_MSG)
 hipError_t hipTexRefSetAddressMode(
     textureReference* texRef,
     int dim,
     enum hipTextureAddressMode am);
+DEPRECATED(DEPRECATED_MSG)
 hipError_t hipTexRefSetArray(
     textureReference* tex,
     hipArray_const_t array,
     unsigned int flags);
+DEPRECATED(DEPRECATED_MSG)
 hipError_t hipTexRefSetFilterMode(
     textureReference* texRef,
     enum hipTextureFilterMode fm);
+DEPRECATED(DEPRECATED_MSG)
 hipError_t hipTexRefSetFlags(
     textureReference* texRef,
     unsigned int Flags);
+DEPRECATED(DEPRECATED_MSG)
 hipError_t hipTexRefSetFormat(
     textureReference* texRef,
     hipArray_Format fmt,
@@ -5136,16 +5149,20 @@ DEPRECATED(DEPRECATED_MSG)
 hipError_t hipTexRefSetBorderColor(
     textureReference* texRef,
     float* pBorderColor);
+DEPRECATED(DEPRECATED_MSG)
 hipError_t hipTexRefSetMipmapFilterMode(
     textureReference* texRef,
     enum hipTextureFilterMode fm);
+DEPRECATED(DEPRECATED_MSG)
 hipError_t hipTexRefSetMipmapLevelBias(
     textureReference* texRef,
     float bias);
+DEPRECATED(DEPRECATED_MSG)
 hipError_t hipTexRefSetMipmapLevelClamp(
     textureReference* texRef,
     float minMipMapLevelClamp,
     float maxMipMapLevelClamp);
+DEPRECATED(DEPRECATED_MSG)
 hipError_t hipTexRefSetMipmappedArray(
     textureReference* texRef,
     struct hipMipmappedArray* mipmappedArray,
@@ -6168,6 +6185,38 @@ hipError_t hipGraphEventWaitNodeSetEvent(hipGraphNode_t node, hipEvent_t event);
 hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
                                              hipEvent_t event);
 
+/**
+ * @brief Get the mem attribute for graphs.
+ *
+ * @param [in] device - device the attr is get for.
+ * @param [in] attr - attr to get.
+ * @param [out] value - value for specific attr.
+ * @returns #hipSuccess, #hipErrorInvalidDevice
+ * @warning : This API is marked as beta, meaning, while this is feature complete,
+ * it is still open to changes and may have outstanding issues.
+ */
+hipError_t hipDeviceGetGraphMemAttribute(int device, hipGraphMemAttributeType attr, void* value);
+
+/**
+ * @brief Set the mem attribute for graphs.
+ *
+ * @param [in] device - device the attr is set for.
+ * @param [in] attr - attr to set.
+ * @param [in] value - value for specific attr.
+ * @returns #hipSuccess, #hipErrorInvalidDevice
+ * @warning : This API is marked as beta, meaning, while this is feature complete,
+ * it is still open to changes and may have outstanding issues.
+ */
+hipError_t hipDeviceSetGraphMemAttribute(int device, hipGraphMemAttributeType attr, void* value);
+
+/**
+ * @brief Free unused memory on specific device used for graph back to OS.
+ *
+ * @param [in] device - device the memory is used for graphs
+ * @warning : This API is marked as beta, meaning, while this is feature complete,
+ * it is still open to changes and may have outstanding issues.
+ */
+hipError_t hipDeviceGraphMemTrim(int device);
 // doxygen end graph API
 /**
  * @}

include/hip/hiprtc.h

@@ -57,13 +57,59 @@ typedef enum hiprtcResult {
     HIPRTC_ERROR_INVALID_PROGRAM = 4,
     HIPRTC_ERROR_INVALID_OPTION = 5,
     HIPRTC_ERROR_COMPILATION = 6,
-    HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7,
-    HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8,
-    HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9,
-    HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10,
-    HIPRTC_ERROR_INTERNAL_ERROR = 11
+    HIPRTC_ERROR_LINKING = 7,
+    HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 8,
+    HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 9,
+    HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 10,
+    HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 11,
+    HIPRTC_ERROR_INTERNAL_ERROR = 12
 } hiprtcResult;
 
+typedef enum hiprtcJIT_option {
+  HIPRTC_JIT_MAX_REGISTERS = 0,
+  HIPRTC_JIT_THREADS_PER_BLOCK,
+  HIPRTC_JIT_WALL_TIME,
+  HIPRTC_JIT_INFO_LOG_BUFFER,
+  HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES,
+  HIPRTC_JIT_ERROR_LOG_BUFFER,
+  HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES,
+  HIPRTC_JIT_OPTIMIZATION_LEVEL,
+  HIPRTC_JIT_TARGET_FROM_HIPCONTEXT,
+  HIPRTC_JIT_TARGET,
+  HIPRTC_JIT_FALLBACK_STRATEGY,
+  HIPRTC_JIT_GENERATE_DEBUG_INFO,
+  HIPRTC_JIT_LOG_VERBOSE,
+  HIPRTC_JIT_GENERATE_LINE_INFO,
+  HIPRTC_JIT_CACHE_MODE,
+  HIPRTC_JIT_NEW_SM3X_OPT,
+  HIPRTC_JIT_FAST_COMPILE,
+  HIPRTC_JIT_GLOBAL_SYMBOL_NAMES,
+  HIPRTC_JIT_GLOBAL_SYMBOL_ADDRESS,
+  HIPRTC_JIT_GLOBAL_SYMBOL_COUNT,
+  HIPRTC_JIT_LTO,
+  HIPRTC_JIT_FTZ,
+  HIPRTC_JIT_PREC_DIV,
+  HIPRTC_JIT_PREC_SQRT,
+  HIPRTC_JIT_FMA,
+  HIPRTC_JIT_NUM_OPTIONS,
+} hiprtcJIT_option;
+
+typedef enum hiprtcJITInputType {
+  HIPRTC_JIT_INPUT_CUBIN = 0,
+  HIPRTC_JIT_INPUT_PTX,
+  HIPRTC_JIT_INPUT_FATBINARY,
+  HIPRTC_JIT_INPUT_OBJECT,
+  HIPRTC_JIT_INPUT_LIBRARY,
+  HIPRTC_JIT_INPUT_NVVM,
+  HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES,
+  HIPRTC_JIT_INPUT_LLVM_BITCODE = 100,
+  HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE = 101,
+  HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE = 102,
+  HIPRTC_JIT_NUM_INPUT_TYPES = (HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES + 3)
+} hiprtcJITInputType;
+
+typedef struct ihiprtcLinkState* hiprtcLinkState;
+
  /**
  * @brief Returns text string message to explain the error which occurred
  *
@@ -223,6 +269,83 @@ hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code);
  */
 hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet);
 
+/**
+ * @brief Creates the link instance via hiprtc APIs.
+ *
+ * @param [in] hip_jit_options
+ * @param [out] hiprtc link state instance
+ * @return HIPRTC_SUCCESS
+ *
+ * @see hiprtcResult
+ */
+hiprtcResult hiprtcLinkCreate(unsigned int num_options, hiprtcJIT_option* option_ptr,
+                              void** option_vals_pptr, hiprtcLinkState* hip_link_state_ptr);
+
+/**
+ * @brief Adds a file with bit code to be linked with options
+ *
+ * @param [in] hiprtc link state, jit input type, file path,
+ *        option reated parameters.
+ * @param [out] None.
+ * @return HIPRTC_SUCCESS
+ *
+ * If input values are invalid, it will
+ * @return HIPRTC_ERROR_INVALID_INPUT
+ *
+ * @see hiprtcResult
+ */
+
+hiprtcResult hiprtcLinkAddFile(hiprtcLinkState hip_link_state, hiprtcJITInputType input_type,
+                               const char* file_path, unsigned int num_options,
+                               hiprtcJIT_option* options_ptr, void** option_values);
+
+/**
+ * @brief Completes the linking of the given program.
+ *
+ * @param [in] hiprtc link state, jit input type, image_ptr ,
+ *        option reated parameters.
+ * @param [out] None.
+ * @return HIPRTC_SUCCESS
+ *
+ * If adding the file fails, it will
+ * @return HIPRTC_ERROR_PROGRAM_CREATION_FAILURE
+ *
+ * @see hiprtcResult
+ */
+
+hiprtcResult hiprtcLinkAddData(hiprtcLinkState hip_link_state, hiprtcJITInputType input_type, 
+                               void* image, size_t image_size, const char* name,
+                               unsigned int num_options, hiprtcJIT_option* options_ptr,
+                               void** option_values);
+
+/**
+ * @brief Completes the linking of the given program.
+ *
+ * @param [in] hiprtc link state instance
+ * @param [out] linked_binary, linked_binary_size.
+ * @return HIPRTC_SUCCESS
+ *
+ * If adding the data fails, it will
+ * @return HIPRTC_ERROR_PROGRAM_CREATION_FAILURE
+ *
+ * @see hiprtcResult
+ */
+hiprtcResult hiprtcLinkComplete(hiprtcLinkState hip_link_state, void** bin_out, size_t* size_out);
+
+/**
+ * @brief Deletes the link instance via hiprtc APIs.
+ *
+ * @param [in] hiprtc link state instance
+ * @param [out] code  the size of binary.
+ * @return HIPRTC_SUCCESS
+ *
+ * If linking fails, it will
+ * @return HIPRTC_ERROR_LINKING
+ *
+ * @see hiprtcResult
+ */
+hiprtcResult hiprtcLinkDestroy(hiprtcLinkState hip_link_state);
+
 #if !defined(_WIN32)
 #pragma GCC visibility pop
 #endif

samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp

@@ -130,7 +130,7 @@ bool runTest(int argc, char** argv) {
 inline bool isImageSupported() {
     int imageSupport = 1;
 #ifdef __HIP_PLATFORM_AMD__
-    HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
+    HIP_CHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
                               0));
 #endif
   return imageSupport != 0;

tests/catch/hipTestMain/config/config_amd_windows.json

@@ -66,7 +66,14 @@
 	   "Unit_hipHostMalloc_Coherent",
 	   "Unit_hipHostMalloc_Default",
 	   "Unit_hipStreamCreate_MultistreamBasicFunctionalities",
-	   "Unit_hipEventIpc"
+	   "Unit_hipEventIpc",
+	   "Unit_hipGraphAddDependencies_NegTest",
+	   "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags",
+	   "Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime",
+	   "Unit_hipGraphAddEventRecordNode_Functional_WithFlags",
+	   "Unit_hipGraphAddEventRecordNode_MultipleRun",
+	   "Unit_hipMalloc3D_Negative",
+	   "Unit_hipPointerGetAttribute_MappedMem"
 	]
 
 }

tests/catch/include/hip_test_common.hh

@@ -71,7 +71,27 @@ THE SOFTWARE.
         printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__);                        \
         abort();                                                                                   \
     }
-
+#if HT_NVIDIA
+#define CTX_CREATE() \
+  hipCtx_t context;\
+  initHipCtx(&context);
+#define CTX_DESTROY() HIPCHECK(hipCtxDestroy(context));
+#define ARRAY_DESTROY(array) HIPCHECK(hipArrayDestroy(array));
+#define HIP_TEX_REFERENCE hipTexRef
+#define HIP_ARRAY hiparray
+static void initHipCtx(hipCtx_t *pcontext) {
+  HIPCHECK(hipInit(0));
+  hipDevice_t device;
+  HIPCHECK(hipDeviceGet(&device, 0));
+  HIPCHECK(hipCtxCreate(pcontext, 0, device));
+}
+#else
+#define CTX_CREATE()
+#define CTX_DESTROY()
+#define ARRAY_DESTROY(array) HIPCHECK(hipFreeArray(array));
+#define HIP_TEX_REFERENCE textureReference*
+#define HIP_ARRAY hipArray*
+#endif
 
 
 // Utility Functions

tests/catch/multiproc/hipMemCoherencyTstMProc.cc

@@ -113,9 +113,8 @@ TEST_CASE("Unit_malloc_CoherentTst") {
     REQUIRE(false);
   }
 
-  /* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
-     GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
-  int GpuId[2] = {0, 0};
+  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no*/
+  int GpuId[1] = {0};
   p = fork();
 
   if (p < 0) {
@@ -128,8 +127,8 @@ TEST_CASE("Unit_malloc_CoherentTst") {
     //  Read string from child and close reading end.
     read(fd1[0], GpuId, 2 * sizeof(int));
     close(fd1[0]);
-    if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
-      WARN("This test is not applicable on MI60 & MI100."
+    if (GpuId[0] == 0) {
+      WARN("This test is applicable for MI200."
               "Skipping the test!!");
       exit(0);
     }
@@ -138,16 +137,11 @@ TEST_CASE("Unit_malloc_CoherentTst") {
     hipDeviceProp_t prop;
     HIPCHECK(hipGetDeviceProperties(&prop, 0));
     char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx906");
+    p = strstr(prop.gcnArchName, "gfx90a");
     if (p) {
-      WARN("gfx906 gpu found on this system!!");
+      WARN("gfx90a gpu found on this system!!");
       GpuId[0] = 1;
     }
-    p = strstr(prop.gcnArchName, "gfx908");
-    if (p) {
-      WARN("gfx908 gpu found on this system!!");
-      GpuId[1] = 1;
-    }
     // Write concatenated string and close writing end
     write(fd1[1], GpuId, 2 * sizeof(int));
     close(fd1[1]);
@@ -208,9 +202,8 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
     REQUIRE(false);
   }
 
-  /* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
-     GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
-  int GpuId[2] = {0, 0};
+  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
+  int GpuId[1] = {0};
   p = fork();
 
   if (p < 0) {
@@ -223,8 +216,8 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
     //  Read string from child and close reading end.
     read(fd1[0], GpuId, 2 * sizeof(int));
     close(fd1[0]);
-    if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
-      WARN("This test is not applicable on MI60 & MI100."
+    if (GpuId[0] == 0) {
+      WARN("This test is applicable for MI200."
               "Skipping the test!!");
       exit(0);
     }
@@ -233,16 +226,11 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
     hipDeviceProp_t prop;
     HIPCHECK(hipGetDeviceProperties(&prop, 0));
     char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx906");
+    p = strstr(prop.gcnArchName, "gfx90a");
     if (p) {
-      WARN("gfx906 gpu found on this system!!");
+      WARN("gfx90a gpu found on this system!!");
       GpuId[0] = 1;
     }
-    p = strstr(prop.gcnArchName, "gfx908");
-    if (p) {
-      WARN("gfx908 gpu found on this system!!");
-      GpuId[1] = 1;
-    }
     // Write concatenated string and close writing end
     write(fd1[1], GpuId, 2 * sizeof(int));
     close(fd1[1]);
@@ -305,9 +293,8 @@ TEST_CASE("Unit_mmap_CoherentTst") {
     REQUIRE(false);
   }
 
-  /* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
-     GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
-  int GpuId[2] = {0, 0};
+  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
+  int GpuId[1] = {0};
   p = fork();
 
   if (p < 0) {
@@ -320,8 +307,8 @@ TEST_CASE("Unit_mmap_CoherentTst") {
     //  Read string from child and close reading end.
     read(fd1[0], GpuId, 2 * sizeof(int));
     close(fd1[0]);
-    if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
-      WARN("This test is not applicable on MI60 & MI100."
+    if (GpuId[0] == 0) {
+      WARN("This test is not applicable for MI200."
               "Skipping the test!!");
       exit(0);
     }
@@ -330,16 +317,11 @@ TEST_CASE("Unit_mmap_CoherentTst") {
     hipDeviceProp_t prop;
     HIPCHECK(hipGetDeviceProperties(&prop, 0));
     char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx906");
+    p = strstr(prop.gcnArchName, "gfx90a");
     if (p) {
-      WARN("gfx906 gpu found on this system!!");
+      WARN("gfx90a gpu found on this system!!");
       GpuId[0] = 1;
     }
-    p = strstr(prop.gcnArchName, "gfx908");
-    if (p) {
-      WARN("gfx908 gpu found on this system!!");
-      GpuId[1] = 1;
-    }
     // Write concatenated string and close writing end
     write(fd1[1], GpuId, 2 * sizeof(int));
     close(fd1[1]);
@@ -403,9 +385,8 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
     REQUIRE(false);
   }
 
-  /* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
-     GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
-  int GpuId[2] = {0, 0};
+  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
+  int GpuId[1] = {0};
   p = fork();
 
   if (p < 0) {
@@ -418,8 +399,8 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
     //  Read string from child and close reading end.
     read(fd1[0], GpuId, 2 * sizeof(int));
     close(fd1[0]);
-    if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
-      WARN("This test is not applicable on MI60 & MI100."
+    if (GpuId[0] == 0) {
+      WARN("This test is applicable for MI200."
               "Skipping the test!!");
       exit(0);
     }
@@ -428,16 +409,11 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
     hipDeviceProp_t prop;
     HIPCHECK(hipGetDeviceProperties(&prop, 0));
     char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx906");
+    p = strstr(prop.gcnArchName, "gfx90a");
     if (p) {
-      WARN("gfx906 gpu found on this system!!");
+      WARN("gfx90a gpu found on this system!!");
       GpuId[0] = 1;
     }
-    p = strstr(prop.gcnArchName, "gfx908");
-    if (p) {
-      WARN("gfx908 gpu found on this system!!");
-      GpuId[1] = 1;
-    }
     // Write concatenated string and close writing end
     write(fd1[1], GpuId, 2 * sizeof(int));
     close(fd1[1]);
@@ -497,7 +473,7 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") {
   if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
   int stat = 0;
@@ -538,7 +514,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") {
   if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
   int stat = 0;
@@ -579,7 +555,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") {
   if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
   int stat = 0;
@@ -587,7 +563,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") {
     int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int);
     YES_COHERENT = false;
     // Allocating hipHostMalloc() memory
-     HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
+    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
     *Ptr = 4;
     hipStream_t strm;
     HIP_CHECK(hipStreamCreate(&strm));
@@ -620,7 +596,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") {
   if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
   int stat = 0;
@@ -662,33 +638,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv1") {
   if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
-  int stat = 0;
-  if (fork() == 0) {  // child process
-    int *Ptr = nullptr, SIZE = sizeof(int);
-    bool HmmMem = false;
-    YES_COHERENT = false;
-    // Allocating hipHostMalloc() memory
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
-    *Ptr = 4;
-    TstCoherency(Ptr, HmmMem);
-    if (YES_COHERENT) {
-      // exit() with code 10 which indicates pass
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(10);
-    } else {
-      // exit() with code 9 which indicates fail
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(9);
-    }
-  } else {  // parent process
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
+  int stat = 0, Pageable = 0;
+
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  if (Pageable) {
+    if (fork() == 0) {  // child process
+      int *Ptr = nullptr, SIZE = sizeof(int);
+      bool HmmMem = false;
+      YES_COHERENT = false;
+      // Allocating hipHostMalloc() memory
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
+      *Ptr = 4;
+      TstCoherency(Ptr, HmmMem);
+      if (YES_COHERENT) {
+        // exit() with code 10 which indicates pass
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(10);
+      } else {
+        // exit() with code 9 which indicates fail
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(9);
+      }
+    } else {  // parent process
+      wait(&stat);
+      int Result = WEXITSTATUS(stat);
+      if (Result != 10) {
+        REQUIRE(false);
+      }
     }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
+               "attribute. Hence skipping the test with Pass result.\n"); 
   }
 }
 #endif
@@ -700,33 +686,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") {
   if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
-  int stat = 0;
-  if (fork() == 0) {  // child process
-    int *Ptr = nullptr, SIZE = sizeof(int);
-    bool HmmMem = false;
-    YES_COHERENT = false;
-    // Allocating hipHostMalloc() memory
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable));
-    *Ptr = 1;
-    TstCoherency(Ptr, HmmMem);
-    if (YES_COHERENT) {
-      // exit() with code 10 which indicates pass
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(10);
-    } else {
-      // exit() with code 9 which indicates fail
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(9);
-    }
-  } else {  // parent process
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
+  int stat = 0, Pageable = 0;
+
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  if (Pageable) {
+    if (fork() == 0) {  // child process
+      int *Ptr = nullptr, SIZE = sizeof(int);
+      bool HmmMem = false;
+      YES_COHERENT = false;
+      // Allocating hipHostMalloc() memory
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable));
+      *Ptr = 1;
+      TstCoherency(Ptr, HmmMem);
+      if (YES_COHERENT) {
+        // exit() with code 10 which indicates pass
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(10);
+      } else {
+        // exit() with code 9 which indicates fail
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(9);
+      }
+    } else {  // parent process
+      wait(&stat);
+      int Result = WEXITSTATUS(stat);
+      if (Result != 10) {
+        REQUIRE(false);
+      }
     }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
+               "attribute. Hence skipping the test with Pass result.\n"); 
   }
 }
 #endif
@@ -737,33 +733,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") {
   if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
-  int stat = 0;
-  if (fork() == 0) {  // child process
-    int *Ptr = nullptr, SIZE = sizeof(int);
-    bool HmmMem = false;
-    YES_COHERENT = false;
-    // Allocating hipHostMalloc() memory
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined));
-    *Ptr = 4;
-    TstCoherency(Ptr, HmmMem);
-    if (YES_COHERENT) {
-      // exit() with code 10 which indicates pass
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(10);
-    } else {
-      // exit() with code 9 which indicates fail
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(9);
-    }
-  } else {  // parent process
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
+  int stat = 0, Pageable = 0;
+
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  if (Pageable) {
+    if (fork() == 0) {  // child process
+      int *Ptr = nullptr, SIZE = sizeof(int);
+      bool HmmMem = false;
+      YES_COHERENT = false;
+      // Allocating hipHostMalloc() memory
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined));
+      *Ptr = 4;
+      TstCoherency(Ptr, HmmMem);
+      if (YES_COHERENT) {
+        // exit() with code 10 which indicates pass
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(10);
+      } else {
+        // exit() with code 9 which indicates fail
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(9);
+      }
+    } else {  // parent process
+      wait(&stat);
+      int Result = WEXITSTATUS(stat);
+      if (Result != 10) {
+        REQUIRE(false);
+      }
     }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
+               "attribute. Hence skipping the test with Pass result.\n"); 
   }
 }
 #endif
@@ -775,33 +781,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") {
   if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
-  int stat = 0;
-  if (fork() == 0) {  // child process
-    int *Ptr = nullptr, SIZE = sizeof(int);
-    bool HmmMem = false;
-    YES_COHERENT = false;
-    // Allocating hipHostMalloc() memory
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
-    *Ptr = 1;
-    TstCoherency(Ptr, HmmMem);
-    if (YES_COHERENT) {
-      // exit() with code 10 which indicates pass
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(10);
-    } else {
-      // exit() with code 9 which indicates fail
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(9);
-    }
-  } else {  // parent process
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
+  int stat = 0, Pageable = 0;
+
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  if (Pageable) {
+    if (fork() == 0) {  // child process
+      int *Ptr = nullptr, SIZE = sizeof(int);
+      bool HmmMem = false;
+      YES_COHERENT = false;
+      // Allocating hipHostMalloc() memory
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
+      *Ptr = 1;
+      TstCoherency(Ptr, HmmMem);
+      if (YES_COHERENT) {
+        // exit() with code 10 which indicates pass
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(10);
+      } else {
+        // exit() with code 9 which indicates fail
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(9);
+      }
+    } else {  // parent process
+      wait(&stat);
+      int Result = WEXITSTATUS(stat);
+      if (Result != 10) {
+        REQUIRE(false);
+      }
     }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
+               "attribute. Hence skipping the test with Pass result.\n"); 
   }
 }
 #endif

tests/catch/unit/device/hipGetDeviceAttribute.cc

@@ -231,6 +231,40 @@ TEST_CASE("Unit_hipGetDeviceAttribute_CheckAttrValues") {
                                     props.texturePitchAlignment));
 }
 
+/**
+ * Validate the hipDeviceAttributeFineGrainSupport property in AMD.
+ */
+#if HT_AMD
+// This is AMD specific property test
+TEST_CASE("Unit_hipGetDeviceAttribute_CheckFineGrainSupport") {
+  int deviceId;
+  int deviceCount = 0;
+  HIP_CHECK(hipGetDeviceCount(&deviceCount));
+  REQUIRE(deviceCount > 0);
+  // Check hipDeviceAttributeFineGrainSupport for all available device
+  // in system.
+  for (int dev = 0; dev < deviceCount; dev++) {
+    HIP_CHECK(hipSetDevice(dev));
+    HIP_CHECK(hipGetDevice(&deviceId));
+    hipDeviceProp_t props;
+    HIP_CHECK(hipGetDeviceProperties(&props, deviceId));
+    int value = 0;
+    HIP_CHECK(hipDeviceGetAttribute(&value,
+        hipDeviceAttributeFineGrainSupport, deviceId));
+    std::string gpu_arch_name(props.gcnArchName);
+    if (std::string::npos != gpu_arch_name.find("gfx90a")) {
+      // Current GPU is gfx90a architecture
+      REQUIRE(value == 1);
+    } else if (std::string::npos != gpu_arch_name.find("gfx906")) {
+      // Current GPU is gfx906 architecture
+      REQUIRE(value == 0);
+    } else if (std::string::npos != gpu_arch_name.find("gfx908")) {
+      // Current GPU is gfx908 architecture
+      REQUIRE(value == 0);
+    }
+  }
+}
+#endif
 /**
  * Validates negative scenarios for hipDeviceGetAttribute
  * scenario1: pi = nullptr

tests/catch/unit/graph/CMakeLists.txt

@@ -22,6 +22,7 @@
 set(TEST_SRC
   hipGraphAddEmptyNode.cc
   hipGraphAddDependencies.cc
+  hipGraphAddEventRecordNode.cc
   hipGraph.cc
   hipSimpleGraphWithKernel.cc
   hipGraphAddMemcpyNode.cc
@@ -51,6 +52,10 @@ set(TEST_SRC
   hipGraphMemsetNodeGetParams.cc
   hipGraphMemsetNodeSetParams.cc
   hipGraphExecMemcpyNodeSetParamsFromSymbol.cc
+  hipGraphEventRecordNodeGetEvent.cc
+  hipGraphEventRecordNodeSetEvent.cc
+  hipGraphEventWaitNodeGetEvent.cc
+  hipGraphExecMemcpyNodeSetParams.cc
 )
 
 hip_add_exe_to_target(NAME GraphsTest

tests/catch/unit/graph/hipGraphAddDependencies.cc

@@ -21,6 +21,7 @@ THE SOFTWARE.
 Testcase Scenarios :
  1) Add different kinds of nodes to graph and add dependencies to nodes.
  Verify sequence of graph execution is based on dependencies created.
+ 2) Negative Scenarios
 */
 
 #include <hip_test_common.hh>
@@ -125,3 +126,124 @@ TEST_CASE("Unit_hipGraphAddDependencies_Functional") {
   HIP_CHECK(hipGraphDestroy(graph));
   HIP_CHECK(hipStreamDestroy(streamForGraph));
 }
+
+/**
+ * Negative Tests for hipGraphAddDependencies.
+ */
+TEST_CASE("Unit_hipGraphAddDependencies_NegTest") {
+  // Initialize
+  constexpr size_t Nbytes = 1024;
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  char *A_d;
+  hipGraphNode_t memset_A;
+  hipMemsetParams memsetParams{};
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(A_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  hipGraphNode_t memcpyH2D_A;
+  char *A_h;
+  A_h = reinterpret_cast<char*>(malloc(Nbytes));
+
+  SECTION("Null Graph") {
+    // Create dependencies
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                   &memsetParams));
+    HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0,
+                                   A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+    REQUIRE(hipErrorInvalidValue == hipGraphAddDependencies(nullptr, &memset_A,
+                                   &memcpyH2D_A, 1));
+  }
+
+  SECTION("numDependencies is zero") {
+    REQUIRE(hipSuccess == hipGraphAddDependencies(graph, nullptr,
+                                   nullptr, 0));
+  }
+
+  SECTION("One Null Graph Node") {
+    // Create dependencies
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                    &memsetParams));
+    REQUIRE(hipErrorInvalidValue == hipGraphAddDependencies(graph, &memset_A,
+                                    nullptr, 1));
+    REQUIRE(hipErrorInvalidValue == hipGraphAddDependencies(graph, nullptr,
+                                    &memset_A, 1));
+  }
+
+  SECTION("Both Null Graph Node") {
+    REQUIRE(hipErrorInvalidValue == hipGraphAddDependencies(graph, nullptr,
+                                                  nullptr, 1));
+  }
+
+  // The following tests fail on AMD.
+  SECTION("from belongs different graph") {
+    hipGraph_t graph1;
+    HIP_CHECK(hipGraphCreate(&graph1, 0));
+    // Create dependencies
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph1, nullptr, 0,
+                                    &memsetParams));
+    HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0,
+                            A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+    REQUIRE(hipErrorInvalidValue == hipGraphAddDependencies(graph, &memset_A,
+                                    &memcpyH2D_A, 1));
+    HIP_CHECK(hipGraphDestroy(graph1));
+  }
+
+  SECTION("To belongs different graph") {
+    hipGraph_t graph1;
+    HIP_CHECK(hipGraphCreate(&graph1, 0));
+    // Create dependencies
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                &memsetParams));
+    HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph1, nullptr,
+                                0, A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+    REQUIRE(hipErrorInvalidValue == hipGraphAddDependencies(graph, &memset_A,
+                                &memcpyH2D_A, 1));
+    HIP_CHECK(hipGraphDestroy(graph1));
+  }
+
+  SECTION("From is uninitialized") {
+    // Create dependencies
+    HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr,
+                                0, A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+    REQUIRE(hipErrorInvalidValue == hipGraphAddDependencies(graph, &memset_A,
+                                &memcpyH2D_A, 1));
+  }
+
+  SECTION("To is uninitialized") {
+    // Create dependencies
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                    &memsetParams));
+    REQUIRE(hipErrorInvalidValue == hipGraphAddDependencies(graph, &memset_A,
+                                    &memcpyH2D_A, 1));
+  }
+
+  SECTION("Duplicate Dependencies") {
+    // Create dependencies
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                          &memsetParams));
+    HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0,
+                          A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+    HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1));
+    REQUIRE(hipErrorInvalidValue == hipGraphAddDependencies(graph, &memset_A,
+            &memcpyH2D_A, 1));
+  }
+
+  SECTION("Same Node Dependencies") {
+    // Create dependencies
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                    &memsetParams));
+    REQUIRE(hipErrorInvalidValue == hipGraphAddDependencies(graph, &memset_A,
+                                    &memset_A, 1));
+  }
+
+  // Destroy
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipGraphDestroy(graph));
+  free(A_h);
+}

tests/catch/unit/graph/hipGraphAddEventRecordNode.cc

@@ -0,0 +1,346 @@
+/*
+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 :
+ 1) Simple Scenario: Create an event node and add it to graph.
+Instantiate and Launch the Graph. Wait for the event to complete.
+The operation must succeed without any failures.
+ 2) Add different kinds of nodes to graph and add dependencies to nodes.
+Create an event record node at the end. Instantiate and Launch the Graph.
+Wait for the event to complete. Verify the results. Event is created using
+hipEventCreate.
+ 3) Add different kinds of nodes to graph and add dependencies to nodes.
+Create event record nodes at the beginning and end. Instantiate and Launch
+the Graph. Wait for the event to complete. Verify the results. Also verify
+the elapsed time. Events are created using hipEventCreate.
+ 4) Add different kinds of nodes to graph and add dependencies to nodes.
+Create an event record node at the end.  Instantiate and Launch graph.
+Wait for the event to complete. Verify the results. Event is created
+using hipEventCreateWithFlags (for different flag values).
+ 5) Create event record node at the beginning with
+flag = hipEventDisableTiming, a memset node and event record nodes at the
+end. Instantiate and Launch the Graph. Wait for the event to complete.
+Verify that hipEventElapsedTime() returns error.
+ 6) Validate scenario 2 by running the graph multiple times in a loop
+(100 times) after instantiation.
+ 7) Negative Scenarios
+    - Output node is a nullptr.
+    - Input graph is a nullptr.
+    - Input dependencies is a nullptr.
+    - Input event is a nullptr.
+    - Input graph is uninitialized.
+    - Input event is uninitialized.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+/**
+ * Scenario 1: Create s simple graph with just one event record
+ * node and instantiate and launch the graph.
+ */
+TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_Simple") {
+  hipGraph_t graph;
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipEvent_t event;
+  HIP_CHECK(hipEventCreate(&event));
+  hipGraphNode_t eventrec;
+  HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
+                                                            event));
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  // Wait for event
+  HIP_CHECK(hipEventSynchronize(event));
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}
+
+/**
+ * Local test function
+ */
+static void validateAddEventRecordNode(bool measureTime, bool withFlags,
+                            int nstep, unsigned flag = 0) {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  hipGraph_t graph;
+  hipGraphNode_t memset_A, memset_B, memsetKer_C;
+  hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C;
+  hipGraphNode_t ker_vecAdd;
+  hipKernelNodeParams kernelNodeParams{};
+  hipStream_t streamForGraph;
+  int *A_d, *B_d, *C_d;
+  int *A_h, *B_h, *C_h;
+  hipGraphExec_t graphExec;
+  hipMemsetParams memsetParams{};
+  int memsetVal{};
+  size_t NElem{N};
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+  unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(A_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                  &memsetParams));
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(B_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0,
+                                  &memsetParams));
+
+  void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func =
+                       reinterpret_cast<void *>(HipTest::memsetReverse<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0,
+                                  &kernelNodeParams));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d,
+                                  A_h, Nbytes, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d,
+                                  B_h, Nbytes, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h,
+                                  C_d, Nbytes, hipMemcpyDeviceToHost));
+
+  void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&ker_vecAdd, graph, nullptr, 0,
+                                                        &kernelNodeParams));
+  hipEvent_t eventstart, eventend;
+  if (withFlags) {
+    HIP_CHECK(hipEventCreateWithFlags(&eventstart, flag));
+    HIP_CHECK(hipEventCreateWithFlags(&eventend, flag));
+  } else {
+    HIP_CHECK(hipEventCreate(&eventstart));
+    HIP_CHECK(hipEventCreate(&eventend));
+  }
+  hipGraphNode_t event_start, event_final;
+  HIP_CHECK(hipGraphAddEventRecordNode(&event_start, graph, nullptr, 0,
+                                                            eventstart));
+  HIP_CHECK(hipGraphAddEventRecordNode(&event_final, graph, nullptr, 0,
+                                                            eventend));
+  // Create dependencies
+  HIP_CHECK(hipGraphAddDependencies(graph, &event_start, &memset_A, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &event_start, &memset_B, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &ker_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &ker_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &ker_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &ker_vecAdd, &memcpyD2H_C, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_C, &event_final, 1));
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  for (int istep = 0; istep < nstep; istep++) {
+    HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+    // Wait for eventend
+    HIP_CHECK(hipEventSynchronize(eventend));
+    // Verify graph execution result
+    HipTest::checkVectorADD(A_h, B_h, C_h, N);
+    if (measureTime) {
+      // Verify event record time difference_type
+      float t = 0.0f;
+      HIP_CHECK(hipEventElapsedTime(&t, eventstart, eventend));
+      REQUIRE(t > 0.0f);
+    }
+  }
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(eventstart));
+  HIP_CHECK(hipEventDestroy(eventend));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}
+
+/**
+ * Scenario 2: Validate event record nodes created without flags.
+ */
+TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags") {
+  // Create events without flags using hipEventCreate and
+  // elapsed time is not validated
+  validateAddEventRecordNode(false, false, 1);
+}
+
+/**
+ * Scenario 3: Validate elapsed time between 2 recorded events.
+ */
+TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime") {
+  // Create events without flags using hipEventCreate and
+  // elapsed time is validated
+  validateAddEventRecordNode(true, false, 1);
+}
+
+/**
+ * Scenario 4: Validate event record nodes created with different
+ * event flags.
+ */
+TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_WithFlags") {
+  // Create events with different flags using hipEventCreate and
+  // elapsed time is not validated
+  SECTION("Flag = hipEventDefault") {
+    validateAddEventRecordNode(false, true, 1, hipEventDefault);
+  }
+
+  SECTION("Flag = hipEventBlockingSync") {
+    validateAddEventRecordNode(false, true, 1, hipEventBlockingSync);
+  }
+
+  SECTION("Flag = hipEventDisableTiming") {
+    validateAddEventRecordNode(false, true, 1, hipEventDisableTiming);
+  }
+}
+
+/**
+ * Scenario 5: Validate hipGraphAddEventRecordNode by executing graph
+ * 100 times in a loop.
+ */
+TEST_CASE("Unit_hipGraphAddEventRecordNode_MultipleRun") {
+  validateAddEventRecordNode(false, false, 100);
+}
+
+/**
+ * Scenario 6: Validate hipGraphAddEventRecordNode with time disabled events.
+ */
+TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_TimingDisabled") {
+  constexpr size_t Nbytes = 1024;
+  hipGraph_t graph;
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  hipEvent_t event_start, event_end;
+  HIP_CHECK(hipEventCreateWithFlags(&event_start, hipEventDisableTiming));
+  HIP_CHECK(hipEventCreateWithFlags(&event_end, hipEventDisableTiming));
+  // memset node
+  char *A_d;
+  hipGraphNode_t memset_A;
+  hipMemsetParams memsetParams{};
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(A_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                   &memsetParams));
+
+  hipGraphNode_t event_node_start, event_node_end;
+  HIP_CHECK(hipGraphAddEventRecordNode(&event_node_start, graph, nullptr, 0,
+                                                            event_start));
+  HIP_CHECK(hipGraphAddEventRecordNode(&event_node_end, graph, nullptr, 0,
+                                                            event_end));
+  // Add dependencies between nodes
+  HIP_CHECK(hipGraphAddDependencies(graph, &event_node_start, &memset_A, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &event_node_end, 1));
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  // Wait for event
+  HIP_CHECK(hipEventSynchronize(event_end));
+  // Validate hipEventElapsedTime returns error code because timing is
+  // disabled for start and end event nodes.
+  float t;
+  REQUIRE(hipSuccess != hipEventElapsedTime(&t, event_start, event_end));
+
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event_end));
+  HIP_CHECK(hipEventDestroy(event_start));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}
+
+/**
+ * Scenario 7: All negative tests
+ */
+TEST_CASE("Unit_hipGraphAddEventRecordNode_Negative") {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipEvent_t event;
+  HIP_CHECK(hipEventCreate(&event));
+  hipGraphNode_t eventwait;
+  SECTION("pGraphNode = nullptr") {
+    REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(nullptr,
+                                    graph, nullptr, 0, event));
+  }
+
+  SECTION("graph = nullptr") {
+    REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(&eventwait,
+                                    nullptr, nullptr, 0, event));
+  }
+
+  SECTION("pDependencies = nullptr and numDependencies != 0") {
+    REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(&eventwait,
+                                    graph, nullptr, 1, event));
+  }
+
+  SECTION("event = nullptr") {
+    REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(&eventwait,
+                                    graph, nullptr, 0, nullptr));
+  }
+
+  SECTION("graph is uninitialized") {
+    hipGraph_t graph_uninit{};
+    REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(&eventwait,
+                                    graph_uninit, nullptr, 0, nullptr));
+  }
+
+  SECTION("event is uninitialized") {
+    hipEvent_t event_uninit{};
+    REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(&eventwait,
+                                    graph, nullptr, 0, event_uninit));
+  }
+
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event));
+}

tests/catch/unit/graph/hipGraphAddMemcpyNode.cc

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
 in the Software without restriction, including without limitation the rights
@@ -18,13 +18,36 @@ THE SOFTWARE.
 */
 
 /**
-Testcase Scenarios :
-1) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D).
-Memcpy nodes are added and assigned to default device.
-2) Perform memcpy operation for 1D, 2D and 3D arrays on default device and verify the results.
-3) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D).
-Memcpy nodes are added and assigned to Peer device.
-4) Perform memcpy operation for 1D, 2D and 3D arrays on Peer device and verify the results.
+Testcase Scenarios : Negative
+1) Pass pGraphNode as nullptr and check if api returns error.
+2) When graph is un-initialized argument(skipping graph creation),
+   api should return error code.
+3) Passing pDependencies as nullptr, api should return success.
+4) When numDependencies is max(size_t) and pDependencies is not valid ptr,
+   api expected to return error code.
+5) When pDependencies is nullptr, but numDependencies is non-zero,
+   api expected to return error.
+6) When pCopyParams is nullptr, api expected to return error code.
+7) API expects atleast one memcpy src pointer to be set.
+   When hipMemcpy3DParms::srcArray and hipMemcpy3DParms::srcPtr.ptr both
+   are nullptr, api expected to return error code.
+8) API expects atleast one memcpy dst pointer to be set.
+   When hipMemcpy3DParms::dstArray and hipMemcpy3DParms::dstPtr.ptr both
+   are nullptr, api expected to return error code.
+9) Passing different element size for hipMemcpy3DParms::srcArray and
+   hipMemcpy3DParms::dstArray is expected to return error code.
+
+Testcase Scenarios : Functional
+1) Add memcpy node to graph and verify memcpy operation is success for all
+   memcpy kinds(H2D, D2H and D2D).
+   Memcpy nodes are added and assigned to default device.
+2) Perform memcpy operation for 1D, 2D and 3D arrays on default device and
+   verify the results.
+3) Add memcpy node to graph and verify memcpy operation is success for all
+   memcpy kinds(H2D, D2H and D2D).
+   Memcpy nodes are added and assigned to Peer device.
+4) Perform memcpy operation for 1D, 2D and 3D arrays on Peer device and
+   verify the results.
 */
 
 #include <hip_test_common.hh>
@@ -34,7 +57,121 @@ Memcpy nodes are added and assigned to Peer device.
 #define YSIZE 32
 #define XSIZE 32
 
-void validateMemcpyNode3DArray(bool peerAccess = false) {
+/* Test verifies hipGraphAddMemcpyNode API Negative scenarios.
+ */
+
+TEST_CASE("Unit_hipGraphAddMemcpyNode_Negative") {
+  constexpr int width{10}, height{10}, depth{10};
+  hipArray *devArray1;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparams;
+  uint32_t size = width * height * depth * sizeof(int);
+  hipGraph_t graph;
+  hipGraphNode_t memcpyNode;
+  hipStream_t streamForGraph;
+  hipError_t ret;
+
+  int *hData = reinterpret_cast<int*>(malloc(size));
+  int *hOutputData = reinterpret_cast<int *>(malloc(size));
+
+  REQUIRE(hData != nullptr);
+  REQUIRE(hOutputData != nullptr);
+  memset(hData, 0, size);
+  memset(hOutputData, 0,  size);
+
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // Initialize host buffer
+  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(&devArray1, &channelDesc,
+                       make_hipExtent(width, height, depth), hipArrayDefault));
+
+  // Host to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width , height, depth);
+  myparams.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int),
+                                      width, height);
+  myparams.dstArray = devArray1;
+  myparams.kind = hipMemcpyHostToDevice;
+
+  SECTION("Pass pGraphNode as nullptr") {
+    ret = hipGraphAddMemcpyNode(nullptr, graph, nullptr, 0, &myparams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("When graph is nullptr") {
+    ret = hipGraphAddMemcpyNode(&memcpyNode, nullptr,  nullptr, 0, &myparams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Passing pDependencies as nullptr") {
+    ret = hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams);
+    REQUIRE(hipSuccess == ret);
+  }
+  SECTION("When numDependencies is max and pDependencies is not valid ptr") {
+    ret = hipGraphAddMemcpyNode(&memcpyNode, graph,
+                                nullptr, INT_MAX, &myparams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("When pDependencies is nullptr, but numDependencies is non-zero") {
+    ret = hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 11, &myparams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass pCopyParams as nullptr") {
+    ret = hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, nullptr);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("API expects atleast one memcpy src pointer to be set") {
+    memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+    myparams.srcPos = make_hipPos(0, 0, 0);
+    myparams.dstPos = make_hipPos(0, 0, 0);
+    myparams.extent = make_hipExtent(width , height, depth);
+    myparams.dstArray = devArray1;
+    myparams.kind = hipMemcpyHostToDevice;
+    ret = hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("API expects atleast one memcpy dst pointer to be set") {
+    memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+    myparams.srcPos = make_hipPos(0, 0, 0);
+    myparams.dstPos = make_hipPos(0, 0, 0);
+    myparams.extent = make_hipExtent(width , height, depth);
+    myparams.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int),
+                                      width, height);
+    myparams.kind = hipMemcpyHostToDevice;
+    ret = hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Passing different element size for hipMemcpy3DParms::srcArray"
+                   "and hipMemcpy3DParms::dstArray") {
+    myparams.srcArray = devArray1;
+    hipArray *devArray2;
+    HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc,
+              make_hipExtent(width+1, height+1, depth+1), hipArrayDefault));
+    myparams.dstArray = devArray2;
+    ret = hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams);
+    REQUIRE(hipErrorInvalidValue == ret);
+    hipFreeArray(devArray2);
+  }
+
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  hipFreeArray(devArray1);
+  free(hData);
+  free(hOutputData);
+}
+
+static void validateMemcpyNode3DArray(bool peerAccess = false) {
   constexpr int width{10}, height{10}, depth{10};
   hipArray *devArray1, *devArray2;
   hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
@@ -137,7 +274,7 @@ void validateMemcpyNode3DArray(bool peerAccess = false) {
   free(hOutputData);
 }
 
-void validateMemcpyNode2DArray(bool peerAccess = false) {
+static void validateMemcpyNode2DArray(bool peerAccess = false) {
   int harray2D[YSIZE][XSIZE]{};
   int harray2Dres[YSIZE][XSIZE]{};
   constexpr int width{XSIZE}, height{YSIZE};
@@ -236,7 +373,7 @@ void validateMemcpyNode2DArray(bool peerAccess = false) {
   hipFreeArray(devArray2);
 }
 
-void validateMemcpyNode1DArray(bool peerAccess = false) {
+static void validateMemcpyNode1DArray(bool peerAccess = false) {
   int harray1D[XSIZE]{};
   int harray1Dres[XSIZE]{};
   constexpr int width{XSIZE};
@@ -380,3 +517,4 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_PeerAccessFunctional") {
     validateMemcpyNode1DArray(true);
   }
 }
+

tests/catch/unit/graph/hipGraphEventRecordNodeGetEvent.cc

@@ -0,0 +1,130 @@
+/*
+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 :
+ 1) Validate that the event returned by hipGraphEventRecordNodeGetEvent matches
+with the event set in hipGraphAddEventRecordNode.
+ 2) Negative Scenarios
+    - Input node is a nullptr.
+    - Output event is a nullptr.
+    - Input node is an empty node.
+    - Input node is a memset node.
+    - Input node is an uninitialized node.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+/**
+ * Local Function to set and get event record node property.
+ */
+static void validateEventRecordNodeGetEvent(unsigned flag) {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipEvent_t event, event_out;
+  HIP_CHECK(hipEventCreateWithFlags(&event, flag));
+  hipGraphNode_t eventrec;
+  HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
+                                                            event));
+  HIP_CHECK(hipGraphEventRecordNodeGetEvent(eventrec, &event_out));
+  // validate set event and get event are same
+  REQUIRE(event == event_out);
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event));
+}
+
+/**
+ * Scenario: Validate scenario 1 for different event flags.
+ */
+TEST_CASE("Unit_hipGraphEventRecordNodeGetEvent_Functional") {
+  // Create event nodes with different flags and validate with
+  // hipGraphEventRecordNodeGetEvent.
+  SECTION("Flag = hipEventDefault") {
+    validateEventRecordNodeGetEvent(hipEventDefault);
+  }
+
+  SECTION("Flag = hipEventBlockingSync") {
+    validateEventRecordNodeGetEvent(hipEventBlockingSync);
+  }
+
+  SECTION("Flag = hipEventDisableTiming") {
+    validateEventRecordNodeGetEvent(hipEventDisableTiming);
+  }
+}
+
+/**
+ * Scenario 2: Negative tests.
+ */
+TEST_CASE("Unit_hipGraphEventRecordNodeGetEvent_Negative") {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipEvent_t event, event_out;
+  HIP_CHECK(hipEventCreate(&event));
+  hipGraphNode_t eventrec;
+  HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
+                                                            event));
+  SECTION("node = nullptr") {
+    REQUIRE(hipErrorInvalidValue == hipGraphEventRecordNodeGetEvent(nullptr,
+                                    &event_out));
+  }
+
+  SECTION("event_out = nullptr") {
+    REQUIRE(hipErrorInvalidValue == hipGraphEventRecordNodeGetEvent(eventrec,
+                                    nullptr));
+  }
+
+  SECTION("input node is empty node") {
+    hipGraphNode_t EmptyGraphNode;
+    HIP_CHECK(hipGraphAddEmptyNode(&EmptyGraphNode, graph, nullptr, 0));
+    REQUIRE(hipErrorInvalidValue ==
+            hipGraphEventRecordNodeGetEvent(EmptyGraphNode, &event_out));
+  }
+
+  SECTION("input node is memset node") {
+    constexpr size_t Nbytes = 1024;
+    char *A_d;
+    hipGraphNode_t memset_A;
+    hipMemsetParams memsetParams{};
+    HIP_CHECK(hipMalloc(&A_d, Nbytes));
+    memset(&memsetParams, 0, sizeof(memsetParams));
+    memsetParams.dst = reinterpret_cast<void*>(A_d);
+    memsetParams.value = 0;
+    memsetParams.pitch = 0;
+    memsetParams.elementSize = sizeof(char);
+    memsetParams.width = Nbytes;
+    memsetParams.height = 1;
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                    &memsetParams));
+    REQUIRE(hipErrorInvalidValue ==
+            hipGraphEventRecordNodeGetEvent(memset_A, &event_out));
+    HIP_CHECK(hipFree(A_d));
+  }
+
+  SECTION("input node is uninitialized node") {
+    hipGraphNode_t node_unit{};
+    REQUIRE(hipErrorInvalidValue ==
+            hipGraphEventRecordNodeGetEvent(node_unit, &event_out));
+  }
+
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event));
+}

tests/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc

@@ -0,0 +1,246 @@
+/*
+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 :
+ 1) Set a different type of event using hipGraphEventRecordNodeSetEvent and
+    validate using hipGraphEventRecordNodeGetEvent.
+ 2) Add different kinds of nodes to graph and add dependencies to nodes.
+    Create an event record node at the end with Default flag. Set a different
+    type of event using hipGraphEventRecordNodeSetEvent.  Instantiate and
+    Launch graph. Wait for the event to complete. Verify the results.
+ 3) Negative Scenarios
+    - Input node parameter is nullptr.
+    - Input event parameter is nullptr.
+    - Empty node is passed as input node.
+    - Memset node is passed as input node.
+    - Input node is an uninitialized node.
+    - Input event is an uninitialized event.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+/**
+ * Local Function: Set Get test
+ */
+static void validateEventRecordNodeSetEvent(unsigned flag) {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  // Create events
+  hipEvent_t event1, event2, event_out;
+  HIP_CHECK(hipEventCreate(&event1));
+  HIP_CHECK(hipEventCreateWithFlags(&event2, flag));
+  hipGraphNode_t eventrec;
+  HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
+                                                            event1));
+  // Set a different event
+  HIP_CHECK(hipGraphEventRecordNodeSetEvent(eventrec, event2));
+  HIP_CHECK(hipGraphEventRecordNodeGetEvent(eventrec, &event_out));
+  // validate set event and get event are same
+  REQUIRE(event2 == event_out);
+  // Free resources
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event1));
+  HIP_CHECK(hipEventDestroy(event2));
+}
+
+/**
+ * Local Function
+ */
+static void setEventWaitNode() {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  // Create events
+  hipEvent_t event1, event2;
+  HIP_CHECK(hipEventCreate(&event1));
+  HIP_CHECK(hipEventCreate(&event2));
+  hipGraphNode_t eventwait;
+  HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0,
+                                                           event1));
+  // Set a different event eventwait using hipGraphEventRecordNodeSetEvent
+  REQUIRE(hipErrorInvalidValue ==
+          hipGraphEventRecordNodeSetEvent(eventwait, event2));
+  // Free resources
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event1));
+  HIP_CHECK(hipEventDestroy(event2));
+}
+
+/**
+ * Scenario 2: Validate Change of event property in event record node.
+ */
+TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetEventProperty") {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  // Create events
+  hipEvent_t event1_start, event2_start, event1_end, event2_end;
+  HIP_CHECK(hipEventCreate(&event1_start));
+  HIP_CHECK(hipEventCreate(&event1_end));
+  HIP_CHECK(hipEventCreateWithFlags(&event2_start, hipEventDisableTiming));
+  HIP_CHECK(hipEventCreateWithFlags(&event2_end, hipEventDisableTiming));
+  // Create nodes
+  hipGraphNode_t event_start_rec, event_end_rec;
+  HIP_CHECK(hipGraphAddEventRecordNode(&event_start_rec, graph, nullptr, 0,
+                                                          event1_start));
+  HIP_CHECK(hipGraphAddEventRecordNode(&event_end_rec, graph, nullptr, 0,
+                                                          event1_end));
+  // Create memset node
+  constexpr size_t Nbytes = 1024;
+  char *A_d;
+  hipGraphNode_t memset_A;
+  hipMemsetParams memsetParams{};
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(A_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                  &memsetParams));
+  // Create dependencies
+  // event_start_rec --> memset_A --> event_end_rec
+  HIP_CHECK(hipGraphAddDependencies(graph, &event_start_rec, &memset_A, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &event_end_rec, 1));
+
+  // Instantiate and launch graph
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+  // Validate by measuring time difference between event_end_rec &
+  // event_start_rec
+  float t = 0.0f;
+  REQUIRE(hipSuccess == hipEventElapsedTime(&t, event1_start,
+                                            event1_end));
+  REQUIRE(t > 0.0f);
+  // Change the event property after instantiation
+  HIP_CHECK(hipGraphEventRecordNodeSetEvent(event_start_rec, event2_start));
+  HIP_CHECK(hipGraphEventRecordNodeSetEvent(event_end_rec, event2_end));
+  // Launch the graph again with the new settings
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+  // Validate hipEventElapsedTime() must return
+  // hipErrorInvalidHandle when events are created using
+  // hipEventDisableTiming flag.
+  t = 0.0f;
+  REQUIRE(hipErrorInvalidHandle ==
+          hipEventElapsedTime(&t, event2_start, event2_end));
+  // Free resources
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipEventDestroy(event1_start));
+  HIP_CHECK(hipEventDestroy(event2_start));
+  HIP_CHECK(hipEventDestroy(event1_end));
+  HIP_CHECK(hipEventDestroy(event2_end));
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+
+/**
+ * Scenario 1: Validate Set Get test for all Event flags
+ */
+TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetGet") {
+  SECTION("Flag = hipEventDefault") {
+    validateEventRecordNodeSetEvent(hipEventDefault);
+  }
+
+  SECTION("Flag = hipEventBlockingSync") {
+    validateEventRecordNodeSetEvent(hipEventBlockingSync);
+  }
+
+  SECTION("Flag = hipEventDisableTiming") {
+    validateEventRecordNodeSetEvent(hipEventDisableTiming);
+  }
+}
+
+/**
+ * Scenario 3: Negative Tests
+ */
+TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_Negative") {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipEvent_t event1, event2;
+  HIP_CHECK(hipEventCreate(&event1));
+  HIP_CHECK(hipEventCreate(&event2));
+  hipGraphNode_t eventrec;
+  HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
+                                                            event1));
+  SECTION("node = nullptr") {
+    REQUIRE(hipErrorInvalidValue == hipGraphEventRecordNodeSetEvent(nullptr,
+                                    event2));
+  }
+
+  SECTION("event_out = nullptr") {
+    REQUIRE(hipErrorInvalidValue == hipGraphEventRecordNodeSetEvent(eventrec,
+                                    nullptr));
+  }
+
+  SECTION("input node is empty node") {
+    hipGraphNode_t EmptyGraphNode;
+    HIP_CHECK(hipGraphAddEmptyNode(&EmptyGraphNode, graph, nullptr, 0));
+    REQUIRE(hipErrorInvalidValue ==
+            hipGraphEventRecordNodeSetEvent(EmptyGraphNode, event2));
+  }
+
+  SECTION("input node is memset node") {
+    constexpr size_t Nbytes = 1024;
+    char *A_d;
+    hipGraphNode_t memset_A;
+    hipMemsetParams memsetParams{};
+    HIP_CHECK(hipMalloc(&A_d, Nbytes));
+    memset(&memsetParams, 0, sizeof(memsetParams));
+    memsetParams.dst = reinterpret_cast<void*>(A_d);
+    memsetParams.value = 0;
+    memsetParams.pitch = 0;
+    memsetParams.elementSize = sizeof(char);
+    memsetParams.width = Nbytes;
+    memsetParams.height = 1;
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                    &memsetParams));
+    REQUIRE(hipErrorInvalidValue ==
+            hipGraphEventRecordNodeSetEvent(memset_A, event2));
+    HIP_CHECK(hipFree(A_d));
+  }
+
+  SECTION("input node is event wait node") {
+    setEventWaitNode();
+  }
+
+  SECTION("input node is uninitialized node") {
+    hipGraphNode_t node_uninit{};
+    REQUIRE(hipErrorInvalidValue ==
+            hipGraphEventRecordNodeSetEvent(node_uninit, event2));
+  }
+
+  SECTION("input event is uninitialized") {
+    hipEvent_t event_uninit{};
+    REQUIRE(hipErrorInvalidValue ==
+            hipGraphEventRecordNodeSetEvent(eventrec, event_uninit));
+  }
+
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event1));
+  HIP_CHECK(hipEventDestroy(event2));
+}

tests/catch/unit/graph/hipGraphEventWaitNodeGetEvent.cc

@@ -0,0 +1,130 @@
+/*
+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 :
+ 1) Validate that the event returned by hipGraphEventWaitNodeGetEvent matches
+with the event set in hipGraphAddEventWaitNode.
+ 2) Negative Scenarios
+    - Input node parameter is passed as nullptr.
+    - Output event parameter is passed as nullptr.
+    - Input node parameter is an empty node.
+    - Input node parameter is a memset node.
+    - Input node parameter is an uninitialized node.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+/**
+ * Local Function
+ */
+static void validateEventWaitNodeGetEvent(unsigned flag) {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipEvent_t event, event_out;
+  HIP_CHECK(hipEventCreateWithFlags(&event, flag));
+  hipGraphNode_t eventwait;
+  HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0,
+                                                          event));
+  HIP_CHECK(hipGraphEventWaitNodeGetEvent(eventwait, &event_out));
+  // validate set event and get event are same
+  REQUIRE(event == event_out);
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event));
+}
+
+/**
+ * Scenario 1
+ */
+TEST_CASE("Unit_hipGraphEventWaitNodeGetEvent_Functional") {
+  // Create event nodes with different flags and validate with
+  // hipGraphEventWaitNodeGetEvent.
+  SECTION("Flag = hipEventDefault") {
+    validateEventWaitNodeGetEvent(hipEventDefault);
+  }
+
+  SECTION("Flag = hipEventBlockingSync") {
+    validateEventWaitNodeGetEvent(hipEventBlockingSync);
+  }
+
+  SECTION("Flag = hipEventDisableTiming") {
+    validateEventWaitNodeGetEvent(hipEventDisableTiming);
+  }
+}
+
+/**
+ * Scenario 2
+ */
+TEST_CASE("Unit_hipGraphEventWaitNodeGetEvent_Negative") {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipEvent_t event, event_out;
+  HIP_CHECK(hipEventCreate(&event));
+  hipGraphNode_t eventwait;
+  HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0,
+                                                            event));
+  SECTION("node = nullptr") {
+    REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeGetEvent(nullptr,
+                                    &event_out));
+  }
+
+  SECTION("event_out = nullptr") {
+    REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeGetEvent(eventwait,
+                                    nullptr));
+  }
+
+  SECTION("input node is empty node") {
+    hipGraphNode_t EmptyGraphNode;
+    HIP_CHECK(hipGraphAddEmptyNode(&EmptyGraphNode, graph, nullptr, 0));
+    REQUIRE(hipErrorInvalidValue ==
+            hipGraphEventWaitNodeGetEvent(EmptyGraphNode, &event_out));
+  }
+
+  SECTION("input node is memset node") {
+    constexpr size_t Nbytes = 1024;
+    char *A_d;
+    hipGraphNode_t memset_A;
+    hipMemsetParams memsetParams{};
+    HIP_CHECK(hipMalloc(&A_d, Nbytes));
+    memset(&memsetParams, 0, sizeof(memsetParams));
+    memsetParams.dst = reinterpret_cast<void*>(A_d);
+    memsetParams.value = 0;
+    memsetParams.pitch = 0;
+    memsetParams.elementSize = sizeof(char);
+    memsetParams.width = Nbytes;
+    memsetParams.height = 1;
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                    &memsetParams));
+    REQUIRE(hipErrorInvalidValue ==
+            hipGraphEventWaitNodeGetEvent(memset_A, &event_out));
+    HIP_CHECK(hipFree(A_d));
+  }
+
+  SECTION("input node is uninitialized") {
+    hipGraphNode_t node_uninit{};
+    REQUIRE(hipErrorInvalidValue ==
+            hipGraphEventWaitNodeGetEvent(node_uninit, &event_out));
+  }
+
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event));
+}

tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParams.cc

@@ -0,0 +1,259 @@
+/*
+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 :
+Functional-
+1) Instantiate a graph with memcpy node, obtain executable graph and update the hipMemcpy3DParms node params with set. Make sure they are taking effect.
+Negative-
+1) Pass hGraphExec as nullptr and verify api returns error code.
+2) Pass node as nullptr and verify api returns error code.
+3) Pass pNodeParams as nullptr and verify api returns error code.
+4) Pass pNodeParams as empty structure object and verify api returns error code.
+5) API expects atleast one memcpy src pointer to be set. When hipMemcpy3DParms::srcArray and hipMemcpy3DParms::srcPtr.ptr both are nullptr, api expected to return error code.
+6) API expects atleast one memcpy dst pointer to be set. When hipMemcpy3DParms::dstArray and hipMemcpy3DParms::dstPtr.ptr both are nullptr, api expected to return error code.
+7) Passing different element size for hipMemcpy3DParms::srcArray and hipMemcpy3DParms::dstArray is expected to return error code.
+8) Pass node of different graph and verify api returns error code.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+
+/* Test verifies hipGraphExecMemcpyNodeSetParams API Negative scenarios.
+ */
+TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams_Negative") {
+  constexpr int width{10}, height{10}, depth{10};
+  hipArray *devArray, *devArray2;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparms;
+  hipError_t ret;
+  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));
+  HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc, make_hipExtent(width+1,
+                             height+1, depth+1), 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;
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms));
+
+  // Instantiate the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
+  SECTION("Pass hGraphExec as nullptr") {
+    ret = hipGraphExecMemcpyNodeSetParams(nullptr, memcpyNode, &myparms);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass node as nullptr") {
+    ret = hipGraphExecMemcpyNodeSetParams(graphExec, nullptr, &myparms);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass pNodeParams as nullptr") {
+    ret = hipGraphExecMemcpyNodeSetParams(graphExec, memcpyNode, nullptr);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass pNodeParams as empty structure object") {
+    hipMemcpy3DParms temp{};
+    ret = hipGraphExecMemcpyNodeSetParams(graphExec, memcpyNode, &temp);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("API expects atleast one memcpy src pointer to be set") {
+    hipMemcpy3DParms temp;
+    memset(&temp, 0x0, sizeof(hipMemcpy3DParms));
+    temp.srcPos = make_hipPos(0, 0, 0);
+    temp.dstPos = make_hipPos(0, 0, 0);
+    temp.extent = make_hipExtent(width , height, depth);
+    temp.dstArray = devArray;
+    temp.kind = hipMemcpyHostToDevice;
+    ret = hipGraphExecMemcpyNodeSetParams(graphExec, memcpyNode, &temp);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("API expects atleast one memcpy dst pointer to be set") {
+    hipMemcpy3DParms temp;
+    memset(&temp, 0x0, sizeof(hipMemcpy3DParms));
+    temp.srcPos = make_hipPos(0, 0, 0);
+    temp.dstPos = make_hipPos(0, 0, 0);
+    temp.extent = make_hipExtent(width , height, depth);
+    temp.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int),
+                                      width, height);
+    temp.kind = hipMemcpyHostToDevice;
+    ret = hipGraphExecMemcpyNodeSetParams(graphExec, memcpyNode, &temp);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Passing different element size for hipMemcpy3DParms::srcArray"
+                   "and hipMemcpy3DParms::dstArray") {
+    hipMemcpy3DParms temp;
+    memset(&temp, 0x0, sizeof(hipMemcpy3DParms));
+    temp.srcPos = make_hipPos(0, 0, 0);
+    temp.dstPos = make_hipPos(0, 0, 0);
+    temp.extent = make_hipExtent(width , height, depth);
+    temp.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int),
+                                      width, height);
+    temp.kind = hipMemcpyHostToDevice;
+    temp.srcArray = devArray;
+    temp.dstArray = devArray2;
+    ret = hipGraphExecMemcpyNodeSetParams(graphExec, memcpyNode, &temp);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Check with other graph node") {
+    hipGraph_t graph1;
+    hipGraphNode_t memcpyNode1;
+    HIP_CHECK(hipGraphCreate(&graph1, 0));
+    HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode1, graph1, NULL, 0, &myparms));
+    ret = hipGraphExecMemcpyNodeSetParams(graphExec, memcpyNode1, &myparms);
+    REQUIRE(hipErrorInvalidValue == ret);
+    HIP_CHECK(hipGraphDestroy(graph1));
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  hipFreeArray(devArray);
+  hipFreeArray(devArray2);
+  free(hData);
+}
+
+/* Test verifies hipGraphExecMemcpyNodeSetParams API Functional scenarios.
+ */
+TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams_Functional") {
+  constexpr int XSIZE = 1024;
+  int harray1D[XSIZE]{};
+  int harray1Dres[XSIZE]{};
+  constexpr int width{XSIZE};
+  hipArray *devArray1, *devArray2;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparams;
+  hipGraph_t graph;
+  hipGraphNode_t memcpyNode;
+  std::vector<hipGraphNode_t> dependencies;
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  // Initialize 1D object
+  for (int i = 0; i < XSIZE; i++) {
+    harray1D[i] = i + 1;
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
+                                                          0, 0, 0, formatKind);
+  // Allocate 1D device array by passing depth(0), height(0)
+  HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc,
+                       make_hipExtent(width, 0, 0), hipArrayDefault));
+  HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc,
+                       make_hipExtent(width, 0, 0), hipArrayDefault));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // Host to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, 1, 1);
+  myparams.srcPtr = make_hipPitchedPtr(harray1D, width * sizeof(int),
+                                       width, 1);
+  myparams.dstArray = devArray1;
+  myparams.kind = hipMemcpyHostToDevice;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams));
+  dependencies.push_back(memcpyNode);
+
+  // Device to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.srcArray = devArray1;
+  myparams.dstArray = devArray2;
+  myparams.extent = make_hipExtent(width, 1, 1);
+  myparams.kind = hipMemcpyDeviceToDevice;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                  dependencies.size(), &myparams));
+  dependencies.clear();
+  dependencies.push_back(memcpyNode);
+
+  // Device to host
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, 1, 1);
+  myparams.dstPtr = make_hipPitchedPtr(harray1Dres, width * sizeof(int),
+                                       width, 1);
+  myparams.srcArray = devArray2;
+  myparams.kind = hipMemcpyDeviceToHost;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                  dependencies.size(), &myparams));
+
+  // Instantiate the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+
+  int harray1Dupdate[XSIZE]{};
+  hipArray *devArray3;
+  HIP_CHECK(hipMalloc3DArray(&devArray3, &channelDesc,
+                       make_hipExtent(width, 0, 0), hipArrayDefault));
+
+  // D2H updated with different pointer harray1Dres -> harray1Dupdate
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, 1, 1);
+  myparams.dstPtr = make_hipPitchedPtr(harray1Dupdate, width * sizeof(int),
+                                      width, 1);
+  myparams.srcArray = devArray2;
+  myparams.kind = hipMemcpyDeviceToHost;
+
+  HIP_CHECK(hipGraphExecMemcpyNodeSetParams(graphExec, memcpyNode, &myparams));
+
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Validate result
+  for (int i = 0; i < XSIZE; i++) {
+    if (harray1D[i] != harray1Dupdate[i]) {
+      INFO("harray1D: " << harray1D[i] << " harray1Dupdate: " <<
+                      harray1Dupdate[i] << " mismatch at : " << i);
+      REQUIRE(false);
+    }
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  hipFreeArray(devArray1);
+  hipFreeArray(devArray2);
+}

tests/catch/unit/memory/CMakeLists.txt

@@ -75,6 +75,17 @@ set(TEST_SRC
     hipHostMalloc.cc
     hipMemcpy.cc
     hipMemcpyAsync.cc
+    hipMallocPitch.cc
+    hipMallocArray.cc
+    hipMalloc3D.cc
+    hipMalloc3DArray.cc
+    hipArrayCreate.cc
+    hipDrvMemcpy3D.cc
+    hipDrvMemcpy3DAsync.cc
+    hipPointerGetAttribute.cc
+    hipDrvPtrGetAttributes.cc
+    hipMallocMngdMultiThread.cc
+    hipMemPrefetchAsync.cc
 )
 else()
 set(TEST_SRC
@@ -129,6 +140,16 @@ set(TEST_SRC
     hipHostMalloc.cc
     hipMemcpy.cc
     hipMemcpyAsync.cc
+    hipMallocPitch.cc
+    hipMallocArray.cc
+    hipMalloc3D.cc
+    hipMalloc3DArray.cc
+    hipArrayCreate.cc
+    hipDrvMemcpy3D.cc
+    hipDrvMemcpy3DAsync.cc
+    hipPointerGetAttribute.cc
+    hipDrvPtrGetAttributes.cc
+    hipMemPrefetchAsync.cc
 )
 endif()
 

tests/catch/unit/memory/hipArrayCreate.cc

@@ -0,0 +1,148 @@
+/*
+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.
+*/
+
+/*
+hipArrayCreate API test scenarios
+1. Negative Scenarios
+2. Allocating Small and big chunk data
+3. Multithreaded scenario
+*/
+
+#include <hip_test_common.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 auto ARRAY_LOOP{100};
+
+/*
+ * This API verifies  memory allocations for small and
+ * bigger chunks of data.
+ * Two scenarios are verified in this API
+ * 1. SmallArray: Allocates NUM_W*NUM_H in a loop and
+ *    releases the memory and verifies the meminfo.
+ * 2. BigArray: Allocates BIGNUM_W*BIGNUM_H in a loop and
+ *    releases the memory and verifies the meminfo
+ *
+ * In both cases, the memory info before allocation and
+ * after releasing the memory should be the same.
+ *
+ */
+
+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);
+    }
+  }
+}
+
+/*Thread function*/
+static void ArrayCreateThreadFunc(int gpu) {
+  ArrayCreate_DiffSizes(gpu);
+}
+
+/* This testcase verifies hipArrayCreate API for small and big chunks data*/
+TEST_CASE("Unit_hipArrayCreate_DiffSizes") {
+  ArrayCreate_DiffSizes(0);
+}
+
+
+/* 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
+and verifies the hipArrayCreate API with small and big chunks data
+*/
+TEST_CASE("Unit_hipArrayCreate_MultiThread") {
+  std::vector<std::thread> threadlist;
+  int devCnt = 0;
+
+  devCnt = HipTest::getDeviceCount();
+
+  size_t tot, avail, ptot, pavail;
+  HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
+  for (int i = 0; i < devCnt; i++) {
+    threadlist.push_back(std::thread(ArrayCreateThreadFunc, i));
+  }
+
+  for (auto &t : threadlist) {
+    t.join();
+  }
+  HIP_CHECK(hipMemGetInfo(&avail, &tot));
+
+  if (pavail != avail) {
+    WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
+    REQUIRE(false);
+  }
+}
+

tests/catch/unit/memory/hipDrvMemcpy3D.cc

@@ -0,0 +1,573 @@
+/*
+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.
+*/
+/*
+ * Test Scenarios
+ * 1. Verifying hipDrvMemcpy3D API for H2A,A2A,A2H scenarios
+ * 2. Verifying hipDrvMemcpy3D API for H2D,D2D,D2H scenarios
+ * 3. Verifying Negative Scenarios
+ * 4. Verifying Extent validation scenarios by passing 0
+ * 5. Verifying hipDrvMemcpy3D API by allocating Memory in
+ *    one GPU and trigger hipDrvMemcpy3D from peer GPU for
+ *    H2D,D2D,D2H scenarios
+ * 6. Verifying hipDrvMemcpy3D API by allocating Memory in
+ *    one GPU and trigger hipDrvMemcpy3D from peer GPU for
+ *    H2A,A2A,A2H scenarios
+ *
+ *    Scenarios 3 is temporarily suspended on AMD
+ *    Scenario 5&6 are not supported in CUDA platform
+ */
+
+#include "hip_test_common.hh"
+#include "hip_test_checkers.hh"
+
+template<typename T>
+class DrvMemcpy3D {
+  int width, height, depth;
+  unsigned int size;
+  hipArray_Format formatKind;
+  hiparray arr, arr1;
+  size_t pitch_D, pitch_E;
+  HIP_MEMCPY3D myparms;
+  hipDeviceptr_t D_m, E_m;
+  T* hData{nullptr};
+ public:
+  DrvMemcpy3D(int l_width, int l_height, int l_depth,
+      hipArray_Format l_format);
+  DrvMemcpy3D() = delete;
+  void AllocateMemory();
+  void SetDefaultData();
+  void HostArray_DrvMemcpy3D(bool device_context_change = false);
+  void HostDevice_DrvMemcpy3D(bool device_context_change = false);
+  void Extent_Validation();
+  void NegativeTests();
+  void DeAllocateMemory();
+};
+
+/* Intializes class variables */
+template <typename T>
+DrvMemcpy3D<T>::DrvMemcpy3D(int l_width, int l_height, int l_depth,
+    hipArray_Format l_format) {
+  width = l_width;
+  height = l_height;
+  depth = l_depth;
+  formatKind = l_format;
+}
+
+/* Allocating Memory */
+template <typename T>
+void DrvMemcpy3D<T>::AllocateMemory() {
+  size = width * height * depth * sizeof(T);
+  hData = reinterpret_cast<T*>(malloc(size));
+  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;
+      }
+    }
+  }
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&D_m),
+                          &pitch_D, width*sizeof(T), height));
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&E_m),
+                          &pitch_E, width*sizeof(T), height));
+  HIP_ARRAY3D_DESCRIPTOR *desc;
+  desc = reinterpret_cast<HIP_ARRAY3D_DESCRIPTOR*>
+         (malloc(sizeof(HIP_ARRAY3D_DESCRIPTOR)));
+  desc->Format = formatKind;
+  desc->NumChannels = 1;
+  desc->Width = width;
+  desc->Height = height;
+  desc->Depth = depth;
+  desc->Flags = hipArrayDefault;
+  HIP_CHECK(hipArray3DCreate(&arr, desc));
+  HIP_CHECK(hipArray3DCreate(&arr1, desc));
+}
+
+/* Setting the default data */
+template <typename T>
+void DrvMemcpy3D<T>::SetDefaultData() {
+  memset(&myparms, 0x0, sizeof(HIP_MEMCPY3D));
+  myparms.srcXInBytes = 0;
+  myparms.srcY = 0;
+  myparms.srcZ = 0;
+  myparms.srcLOD = 0;
+  myparms.dstXInBytes = 0;
+  myparms.dstY = 0;
+  myparms.dstZ = 0;
+  myparms.dstLOD = 0;
+  myparms.WidthInBytes = width*sizeof(T);
+  myparms.Height = height;
+  myparms.Depth = depth;
+}
+
+/*
+This function verifies the negative scenarios of
+hipDrvMemcpy3D API
+*/
+template <typename T>
+void DrvMemcpy3D<T>::NegativeTests() {
+  HIP_CHECK(hipSetDevice(0));
+  AllocateMemory();
+  SetDefaultData();
+  int deviceId;
+  HIP_CHECK(hipGetDevice(&deviceId));
+  unsigned int MaxPitch;
+  HIP_CHECK(hipDeviceGetAttribute(reinterpret_cast<int *>(&MaxPitch),
+                                  hipDeviceAttributeMaxPitch, deviceId));
+  myparms.srcHost = hData;
+  myparms.dstArray = arr;
+  myparms.srcPitch = width * sizeof(T);
+  myparms.srcHeight = height;
+#if HT_NVIDIA
+  myparms.srcMemoryType = CU_MEMORYTYPE_HOST;
+  myparms.dstMemoryType = CU_MEMORYTYPE_ARRAY;
+#else
+  myparms.srcMemoryType = hipMemoryTypeHost;
+  myparms.dstMemoryType = hipMemoryTypeArray;
+#endif
+
+  SECTION("Passing nullptr to Source Host") {
+    myparms.srcHost = nullptr;
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("Passing both dst host and device") {
+    myparms.dstHost = hData;
+    myparms.dstArray = nullptr;
+    myparms.dstDevice = D_m;
+    myparms.WidthInBytes = pitch_D;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("Passing max value to WidthInBytes") {
+    myparms.WidthInBytes = std::numeric_limits<int>::max();
+    myparms.Height = std::numeric_limits<int>::max();
+    myparms.Depth = std::numeric_limits<int>::max();
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("Passing width > max width size") {
+    myparms.WidthInBytes = width*sizeof(T) + 1;
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("Passing height > max height size") {
+    myparms.Height = height + 1;
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("Passing depth > max depth size") {
+    myparms.Depth = depth + 1;
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("widthinbytes + srcXinBytes is out of bound") {
+    myparms.srcXInBytes = 1;
+    myparms.dstArray = nullptr;
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("widthinbytes + dstXinBytes is out of bound") {
+    myparms.dstXInBytes = pitch_D;
+    myparms.dstArray = nullptr;
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("srcY + height is out of bound") {
+    myparms.srcY = 1;
+    myparms.dstArray = nullptr;
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("dstY + height out of bounds") {
+    myparms.dstY = 1;
+    myparms.dstArray = nullptr;
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("src pitch greater than Max allowed pitch") {
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+    myparms.dstMemoryType = CU_MEMORYTYPE_HOST;
+#else
+    myparms.srcMemoryType = hipMemoryTypeDevice;
+    myparms.dstMemoryType = hipMemoryTypeHost;
+#endif
+    myparms.srcDevice = D_m;
+    myparms.srcHost = nullptr;
+    myparms.srcPitch = MaxPitch;
+    myparms.srcHeight = height;
+    myparms.dstHost = hData;
+    myparms.dstArray = nullptr;
+    myparms.dstPitch = width*sizeof(T);
+    myparms.dstHeight = height;
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("dst pitch greater than Max allowed pitch") {
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstArray = nullptr;
+    myparms.dstPitch = MaxPitch+1;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("Nullptr to src/dst device") {
+    myparms.dstDevice = hipDeviceptr_t(nullptr);
+    myparms.dstArray = nullptr;
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("Nullptr to src/dst array") {
+    myparms.dstArray = nullptr;
+    REQUIRE(hipDrvMemcpy3D(&myparms) != hipSuccess);
+  }
+
+  SECTION("Nullptr to hipDrvMemcpy3D") {
+    REQUIRE(hipDrvMemcpy3D(nullptr) != hipSuccess);
+  }
+
+  DeAllocateMemory();
+}
+/*
+This function verifies the Extent validation scenarios of
+hipDrvMemcpy3D API
+*/
+template <typename T>
+void DrvMemcpy3D<T>::Extent_Validation() {
+  HIP_CHECK(hipSetDevice(0));
+  // Allocating the memory
+  AllocateMemory();
+
+  // Setting default data
+  SetDefaultData();
+#if HT_NVIDIA
+  myparms.srcMemoryType = CU_MEMORYTYPE_HOST;
+  myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+  myparms.srcMemoryType = hipMemoryTypeHost;
+  myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+  myparms.srcHost = hData;
+  myparms.srcPitch = width * sizeof(T);
+  myparms.srcHeight = height;
+  myparms.dstDevice = D_m;
+  myparms.dstPitch = pitch_D;
+  myparms.dstHeight = height;
+
+  SECTION("WidthInBytes is 0") {
+    myparms.WidthInBytes = 0;
+    HIP_CHECK(hipDrvMemcpy3D(&myparms));
+  }
+
+  SECTION("Height is 0") {
+    myparms.Height = 0;
+    HIP_CHECK(hipDrvMemcpy3D(&myparms));
+  }
+
+  SECTION("Depth is 0") {
+    myparms.Depth = 0;
+    HIP_CHECK(hipDrvMemcpy3D(&myparms));
+  }
+
+  DeAllocateMemory();
+}
+/*
+This Function verifies following functionalities of hipDrvMemcpy3D API
+1. Host to Device copy
+2. Device to Device
+3. Device to Host
+In the end  validates the results.
+
+This functionality is verified in 2 scenarios
+1. Basic scenario on same GPU device
+2. Device context change scenario where memory is allocated in 1 GPU
+   and hipDrvMemcpy3D API is trigerred from another GPU
+*/
+template <typename T>
+void DrvMemcpy3D<T>::HostDevice_DrvMemcpy3D(bool device_context_change) {
+  HIP_CHECK(hipSetDevice(0));
+  bool skip_test = false;
+  int peerAccess = 0;
+  AllocateMemory();
+  if (device_context_change) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 0, 1));
+    if (!peerAccess) {
+      WARN("skipped the testcase as no peer access");
+      skip_test = true;
+    } else {
+      HIP_CHECK(hipSetDevice(1));
+    }
+  }
+  if (!skip_test) {
+    SetDefaultData();
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_HOST;
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.srcMemoryType = hipMemoryTypeHost;
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    myparms.srcHost = hData;
+    myparms.srcPitch = width * sizeof(T);
+    myparms.srcHeight = height;
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+    HIP_CHECK(hipDrvMemcpy3D(&myparms));
+
+    // Device to Device
+    SetDefaultData();
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.srcMemoryType = hipMemoryTypeDevice;
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    myparms.srcDevice = hipDeviceptr_t(D_m);
+    myparms.srcPitch = pitch_D;
+    myparms.srcHeight = height;
+    myparms.dstDevice = hipDeviceptr_t(E_m);
+    myparms.dstPitch = pitch_E;
+    myparms.dstHeight = height;
+    HIP_CHECK(hipDrvMemcpy3D(&myparms));
+    T *hOutputData = reinterpret_cast<T*>(malloc(size));
+    memset(hOutputData, 0,  size);
+
+    // Device to host
+    SetDefaultData();
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+    myparms.dstMemoryType = CU_MEMORYTYPE_HOST;
+#else
+    myparms.srcMemoryType = hipMemoryTypeDevice;
+    myparms.dstMemoryType = hipMemoryTypeHost;
+#endif
+    myparms.srcDevice = hipDeviceptr_t(E_m);
+    myparms.srcPitch = pitch_E;
+    myparms.srcHeight = height;
+    myparms.dstHost = hOutputData;
+    myparms.dstPitch = width * sizeof(T);
+    myparms.dstHeight = height;
+    HIP_CHECK(hipDrvMemcpy3D(&myparms));
+
+    HipTest::checkArray(hData, hOutputData, width, height, depth);
+    free(hOutputData);
+  }
+  DeAllocateMemory();
+}
+
+/*
+This Function verifies following functionalities of hipDrvMemcpy3D API
+1. Host to Array copy
+2. Array to Array
+3. Array to Host
+In the end  validates the results.
+
+This functionality is verified in 2 scenarios
+1. Basic scenario on same GPU device
+2. Device context change scenario where memory is allocated in 1 GPU
+   and hipDrvMemcpy3D API is trigerred from another GPU
+*/
+template <typename T>
+void DrvMemcpy3D<T>::HostArray_DrvMemcpy3D(bool device_context_change) {
+  HIP_CHECK(hipSetDevice(0));
+  bool skip_test = false;
+  int peerAccess = 0;
+  AllocateMemory();
+  if (device_context_change) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 0, 1));
+    if (!peerAccess) {
+      WARN("skipped the testcase as no peer access");
+      skip_test = true;
+    } else {
+      HIP_CHECK(hipSetDevice(1));
+    }
+  }
+  if (!skip_test) {
+    SetDefaultData();
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_HOST;
+    myparms.dstMemoryType = CU_MEMORYTYPE_ARRAY;
+#else
+    myparms.srcMemoryType = hipMemoryTypeHost;
+    myparms.dstMemoryType = hipMemoryTypeArray;
+#endif
+    myparms.srcHost = hData;
+    myparms.srcPitch = width * sizeof(T);
+    myparms.srcHeight = height;
+    myparms.dstArray = arr;
+    HIP_CHECK(hipDrvMemcpy3D(&myparms));
+    // Array to Array
+    SetDefaultData();
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_ARRAY;
+    myparms.dstMemoryType = CU_MEMORYTYPE_ARRAY;
+#else
+    myparms.srcMemoryType = hipMemoryTypeArray;
+    myparms.dstMemoryType = hipMemoryTypeArray;
+#endif
+    myparms.srcArray = arr;
+    myparms.dstArray = arr1;
+    HIP_CHECK(hipDrvMemcpy3D(&myparms));
+    T *hOutputData = reinterpret_cast<T*>(malloc(size));
+    memset(hOutputData, 0,  size);
+    SetDefaultData();
+    // Device to host
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_ARRAY;
+    myparms.dstMemoryType = CU_MEMORYTYPE_HOST;
+#else
+    myparms.srcMemoryType = hipMemoryTypeArray;
+    myparms.dstMemoryType = hipMemoryTypeHost;
+#endif
+    myparms.srcArray = arr1;
+    myparms.dstHost = hOutputData;
+    myparms.dstPitch = width * sizeof(T);
+    myparms.dstHeight = height;
+    HIP_CHECK(hipDrvMemcpy3D(&myparms));
+
+    HipTest::checkArray(hData, hOutputData, width, height, depth);
+    free(hOutputData);
+  }
+  DeAllocateMemory();
+}
+/* DeAllocating the memory */
+template <typename T>
+void DrvMemcpy3D<T>::DeAllocateMemory() {
+  hipArrayDestroy(arr);
+  hipArrayDestroy(arr1);
+  free(hData);
+}
+
+/* Verifying hipDrvMemcpy3D API Host to Array for different datatypes */
+TEMPLATE_TEST_CASE("Unit_hipDrvMemcpy3D_MultipleDataTypes", "",
+    uint8_t, int, float) {
+  for (int i = 1; i < 25; i++) {
+    if (std::is_same<TestType, float>::value) {
+      DrvMemcpy3D<TestType> memcpy3d_float(i, i, i, HIP_AD_FORMAT_FLOAT);
+      memcpy3d_float.HostArray_DrvMemcpy3D();
+    } else if (std::is_same<TestType, uint8_t>::value)  {
+      DrvMemcpy3D<TestType> memcpy3d_intx(i, i, i, HIP_AD_FORMAT_UNSIGNED_INT8);
+      memcpy3d_intx.HostArray_DrvMemcpy3D();
+    } else if (std::is_same<TestType, int>::value)  {
+      DrvMemcpy3D<TestType> memcpy3d_inty(i, i, i, HIP_AD_FORMAT_SIGNED_INT32);
+      memcpy3d_inty.HostArray_DrvMemcpy3D();
+    }
+  }
+}
+
+/* This testcase verifies H2D copy of hipDrvMemcpy3D API */
+TEST_CASE("Unit_hipDrvMemcpy3D_HosttoDevice") {
+  DrvMemcpy3D<float> memcpy3d_D2H_float(10, 10, 1, HIP_AD_FORMAT_FLOAT);
+  memcpy3d_D2H_float.HostDevice_DrvMemcpy3D();
+}
+
+/* This testcase verifies negative scenarios of hipDrvMemcpy3D API */
+#if HT_NVIDIA
+TEST_CASE("Unit_hipDrvMemcpy3D_Negative") {
+  DrvMemcpy3D<float> memcpy3d(10, 10, 1, HIP_AD_FORMAT_FLOAT);
+  memcpy3d.NegativeTests();
+}
+#endif
+
+/* This testcase verifies extent validation scenarios of hipDrvMemcpy3D API */
+TEST_CASE("Unit_hipDrvMemcpy3D_ExtentValidation") {
+  DrvMemcpy3D<float> memcpy3d(10, 10, 1, HIP_AD_FORMAT_FLOAT);
+  memcpy3d.Extent_Validation();
+}
+
+#if HT_AMD
+/* This testcase verifies H2D copy in device context
+change scenario for hipDrvMemcpy3D API */
+TEST_CASE("Unit_hipDrvMemcpy3D_H2DDeviceContextChange") {
+  int numDevices = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    DrvMemcpy3D<float> memcpy3d(10, 10, 1, HIP_AD_FORMAT_FLOAT);
+    memcpy3d.HostDevice_DrvMemcpy3D(true);
+  } else {
+    SUCCEED("skipped testcase as Device count is < 2");
+  }
+}
+
+
+/* This testcase verifies Host to Array copy in device context
+change scenario for hipDrvMemcpy3D API */
+TEST_CASE("Unit_hipDrvMemcpy3D_Host2ArrayDeviceContextChange") {
+  int numDevices = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    DrvMemcpy3D<float> memcpy3d(10, 10, 1, HIP_AD_FORMAT_FLOAT);
+    memcpy3d.HostArray_DrvMemcpy3D(true);
+  } else {
+    SUCCEED("skipped testcase as Device count is < 2");
+  }
+}
+#endif

tests/catch/unit/memory/hipDrvMemcpy3DAsync.cc

@@ -0,0 +1,594 @@
+/*
+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.
+*/
+/*
+ * Test Scenarios
+ * 1. Verifying hipDrvMemcpy3DAsync API for H2A,A2A,A2H scenarios
+ * 2. Verifying hipDrvMemcpy3DAsync API for H2D,D2D,D2H scenarios
+ * 3. Verifying Negative Scenarios
+ * 4. Verifying Extent validation scenarios by passing 0
+ * 5. Verifying hipDrvMemcpy3DAsync API by allocating Memory in
+ *    one GPU and trigger hipDrvMemcpy3DAsync from peer GPU for
+ *    H2D,D2D,D2H scenarios
+ * 6. Verifying hipDrvMemcpy3DAsync API by allocating Memory in
+ *    one GPU and trigger hipDrvMemcpy3DAsync from peer GPU for
+ *    H2A,A2A,A2H scenarios
+ *
+ *    Scenarios 3 is temporarily excluded in AMD platform
+ *    Scenario 5&6 are excluded in CUDA platform
+ */
+
+#include "hip_test_common.hh"
+#include "hip_test_checkers.hh"
+
+template<typename T>
+class DrvMemcpy3DAsync {
+  int width, height, depth;
+  unsigned int size;
+  hipArray_Format formatKind;
+  hiparray arr, arr1;
+  hipStream_t stream;
+  size_t pitch_D, pitch_E;
+  HIP_MEMCPY3D myparms;
+  hipDeviceptr_t D_m, E_m;
+  T* hData{nullptr};
+ public:
+  DrvMemcpy3DAsync(int l_width, int l_height, int l_depth,
+      hipArray_Format l_format);
+  DrvMemcpy3DAsync() = delete;
+  void AllocateMemory();
+  void SetDefaultData();
+  void HostArray_DrvMemcpy3DAsync(bool device_context_change = false);
+  void HostDevice_DrvMemcpy3DAsync(bool device_context_change = false);
+  void Extent_Validation();
+  void NegativeTests();
+  void DeAllocateMemory();
+};
+
+/* Intializes class variables */
+template <typename T>
+DrvMemcpy3DAsync<T>::DrvMemcpy3DAsync(int l_width, int l_height, int l_depth,
+    hipArray_Format l_format) {
+  width = l_width;
+  height = l_height;
+  depth = l_depth;
+  formatKind = l_format;
+}
+
+/* Allocating Memory */
+template <typename T>
+void DrvMemcpy3DAsync<T>::AllocateMemory() {
+  size = width * height * depth * sizeof(T);
+  hData = reinterpret_cast<T*>(malloc(size));
+  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;
+      }
+    }
+  }
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&D_m),
+                          &pitch_D, width*sizeof(T), height));
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&E_m),
+                          &pitch_E, width*sizeof(T), height));
+  HIP_ARRAY3D_DESCRIPTOR *desc;
+  desc = reinterpret_cast<HIP_ARRAY3D_DESCRIPTOR*>
+         (malloc(sizeof(HIP_ARRAY3D_DESCRIPTOR)));
+  desc->Format = formatKind;
+  desc->NumChannels = 1;
+  desc->Width = width;
+  desc->Height = height;
+  desc->Depth = depth;
+  desc->Flags = hipArrayDefault;
+  HIP_CHECK(hipArray3DCreate(&arr, desc));
+  HIP_CHECK(hipArray3DCreate(&arr1, desc));
+}
+
+/* Setting the default data */
+template <typename T>
+void DrvMemcpy3DAsync<T>::SetDefaultData() {
+  memset(&myparms, 0x0, sizeof(HIP_MEMCPY3D));
+  myparms.srcXInBytes = 0;
+  myparms.srcY = 0;
+  myparms.srcZ = 0;
+  myparms.srcLOD = 0;
+  myparms.dstXInBytes = 0;
+  myparms.dstY = 0;
+  myparms.dstZ = 0;
+  myparms.dstLOD = 0;
+  myparms.WidthInBytes = width*sizeof(T);
+  myparms.Height = height;
+  myparms.Depth = depth;
+}
+
+/*
+This function verifies the negative scenarios of
+hipDrvMemcpy3DAsync API
+*/
+template <typename T>
+void DrvMemcpy3DAsync<T>::NegativeTests() {
+  HIP_CHECK(hipSetDevice(0));
+  AllocateMemory();
+  SetDefaultData();
+  int deviceId;
+  HIP_CHECK(hipGetDevice(&deviceId));
+  unsigned int MaxPitch;
+  HIP_CHECK(hipDeviceGetAttribute(reinterpret_cast<int *>(&MaxPitch),
+                                  hipDeviceAttributeMaxPitch, deviceId));
+  myparms.srcHost = hData;
+  myparms.dstArray = arr;
+  myparms.srcPitch = width * sizeof(T);
+  myparms.srcHeight = height;
+#if HT_NVIDIA
+  myparms.srcMemoryType = CU_MEMORYTYPE_HOST;
+  myparms.dstMemoryType = CU_MEMORYTYPE_ARRAY;
+#else
+  myparms.srcMemoryType = hipMemoryTypeHost;
+  myparms.dstMemoryType = hipMemoryTypeArray;
+#endif
+
+  SECTION("Passing nullptr to Source Host") {
+    myparms.srcHost = nullptr;
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("Passing both dst host and device") {
+    myparms.dstHost = hData;
+    myparms.dstArray = nullptr;
+    myparms.dstDevice = D_m;
+    myparms.WidthInBytes = pitch_D;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("Passing max value to WidthInBytes") {
+    myparms.WidthInBytes = std::numeric_limits<int>::max();
+    myparms.Height = std::numeric_limits<int>::max();
+    myparms.Depth = std::numeric_limits<int>::max();
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("Passing width > max width size") {
+    myparms.WidthInBytes = width*sizeof(T) + 1;
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("Passing height > max height size") {
+    myparms.Height = height + 1;
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("Passing depth > max depth size") {
+    myparms.Depth = depth + 1;
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("widthinbytes + srcXinBytes is out of bound") {
+    myparms.srcXInBytes = 1;
+    myparms.dstArray = nullptr;
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("widthinbytes + dstXinBytes is out of bound") {
+    myparms.dstXInBytes = pitch_D;
+    myparms.dstArray = nullptr;
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("srcY + height is out of bound") {
+    myparms.srcY = 1;
+    myparms.dstArray = nullptr;
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("dstY + height out of bounds") {
+    myparms.dstY = 1;
+    myparms.dstArray = nullptr;
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("src pitch greater than Max allowed pitch") {
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+    myparms.dstMemoryType = CU_MEMORYTYPE_HOST;
+#else
+    myparms.srcMemoryType = hipMemoryTypeDevice;
+    myparms.dstMemoryType = hipMemoryTypeHost;
+#endif
+    myparms.srcDevice = D_m;
+    myparms.srcHost = nullptr;
+    myparms.srcPitch = MaxPitch;
+    myparms.srcHeight = height;
+    myparms.dstHost = hData;
+    myparms.dstArray = nullptr;
+    myparms.dstPitch = width*sizeof(T);
+    myparms.dstHeight = height;
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("dst pitch greater than Max allowed pitch") {
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstArray = nullptr;
+    myparms.dstPitch = MaxPitch+1;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("Nullptr to src/dst device") {
+    myparms.dstDevice = hipDeviceptr_t(nullptr);
+    myparms.dstArray = nullptr;
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+#if HT_NVIDIA
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("Nullptr to src/dst array") {
+    myparms.dstArray = nullptr;
+    REQUIRE(hipDrvMemcpy3DAsync(&myparms, stream) != hipSuccess);
+  }
+
+  SECTION("Nullptr to hipDrvMemcpy3DAsync") {
+    REQUIRE(hipDrvMemcpy3DAsync(nullptr, stream) != hipSuccess);
+  }
+
+  DeAllocateMemory();
+}
+/*
+This function verifies the Extent validation scenarios of
+hipDrvMemcpy3DAsync API
+*/
+template <typename T>
+void DrvMemcpy3DAsync<T>::Extent_Validation() {
+  HIP_CHECK(hipSetDevice(0));
+  // Allocating the memory
+  AllocateMemory();
+
+  // Setting default data
+  SetDefaultData();
+#if HT_NVIDIA
+  myparms.srcMemoryType = CU_MEMORYTYPE_HOST;
+  myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+  myparms.srcMemoryType = hipMemoryTypeHost;
+  myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+  myparms.srcHost = hData;
+  myparms.srcPitch = width * sizeof(T);
+  myparms.srcHeight = height;
+  myparms.dstDevice = D_m;
+  myparms.dstPitch = pitch_D;
+  myparms.dstHeight = height;
+
+  SECTION("WidthInBytes is 0") {
+    myparms.WidthInBytes = 0;
+    HIP_CHECK(hipDrvMemcpy3DAsync(&myparms, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+  }
+
+  SECTION("Height is 0") {
+    myparms.Height = 0;
+    HIP_CHECK(hipDrvMemcpy3DAsync(&myparms, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+  }
+
+  SECTION("Depth is 0") {
+    myparms.Depth = 0;
+    HIP_CHECK(hipDrvMemcpy3DAsync(&myparms, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+  }
+
+  DeAllocateMemory();
+}
+/*
+This Function verifies following functionalities of hipDrvMemcpy3DAsync API
+1. Host to Device copy
+2. Device to Device
+3. Device to Host
+In the end  validates the results.
+
+This functionality is verified in 2 scenarios
+1. Basic scenario on same GPU device
+2. Device context change scenario where memory is allocated in 1 GPU
+   and hipDrvMemcpy3DAsync API is trigerred from another GPU
+*/
+template <typename T>
+void DrvMemcpy3DAsync<T>::HostDevice_DrvMemcpy3DAsync
+     (bool device_context_change) {
+  HIP_CHECK(hipSetDevice(0));
+  bool skip_test = false;
+  int peerAccess = 0;
+  AllocateMemory();
+  if (device_context_change) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 0, 1));
+    if (!peerAccess) {
+      WARN("skipped the testcase as no peer access");
+      skip_test = true;
+    } else {
+      HIP_CHECK(hipSetDevice(1));
+    }
+  }
+  if (!skip_test) {
+    SetDefaultData();
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_HOST;
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.srcMemoryType = hipMemoryTypeHost;
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    myparms.srcHost = hData;
+    myparms.srcPitch = width * sizeof(T);
+    myparms.srcHeight = height;
+    myparms.dstDevice = hipDeviceptr_t(D_m);
+    myparms.dstPitch = pitch_D;
+    myparms.dstHeight = height;
+    HIP_CHECK(hipDrvMemcpy3DAsync(&myparms, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+
+    // Device to Device
+    SetDefaultData();
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+    myparms.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+#else
+    myparms.srcMemoryType = hipMemoryTypeDevice;
+    myparms.dstMemoryType = hipMemoryTypeDevice;
+#endif
+    myparms.srcDevice = hipDeviceptr_t(D_m);
+    myparms.srcPitch = pitch_D;
+    myparms.srcHeight = height;
+    myparms.dstDevice = hipDeviceptr_t(E_m);
+    myparms.dstPitch = pitch_E;
+    myparms.dstHeight = height;
+    HIP_CHECK(hipDrvMemcpy3DAsync(&myparms, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+    T *hOutputData = reinterpret_cast<T*>(malloc(size));
+    memset(hOutputData, 0,  size);
+
+    // Device to host
+    SetDefaultData();
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+    myparms.dstMemoryType = CU_MEMORYTYPE_HOST;
+#else
+    myparms.srcMemoryType = hipMemoryTypeDevice;
+    myparms.dstMemoryType = hipMemoryTypeHost;
+#endif
+    myparms.srcDevice = hipDeviceptr_t(E_m);
+    myparms.srcPitch = pitch_E;
+    myparms.srcHeight = height;
+    myparms.dstHost = hOutputData;
+    myparms.dstPitch = width * sizeof(T);
+    myparms.dstHeight = height;
+    HIP_CHECK(hipDrvMemcpy3DAsync(&myparms, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+
+    HipTest::checkArray(hData, hOutputData, width, height, depth);
+    free(hOutputData);
+  }
+  DeAllocateMemory();
+}
+
+/*
+This Function verifies following functionalities of hipDrvMemcpy3DAsync API
+1. Host to Array copy
+2. Array to Array
+3. Array to Host
+In the end  validates the results.
+
+This functionality is verified in 2 scenarios
+1. Basic scenario on same GPU device
+2. Device context change scenario where memory is allocated in 1 GPU
+   and hipDrvMemcpy3DAsync API is trigerred from another GPU
+*/
+template <typename T>
+void DrvMemcpy3DAsync<T>::HostArray_DrvMemcpy3DAsync
+                          (bool device_context_change) {
+  HIP_CHECK(hipSetDevice(0));
+  bool skip_test = false;
+  int peerAccess = 0;
+  AllocateMemory();
+  if (device_context_change) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 0, 1));
+    if (!peerAccess) {
+      WARN("skipped the testcase as no peer access");
+      skip_test = true;
+    } else {
+      HIP_CHECK(hipSetDevice(1));
+    }
+  }
+  if (!skip_test) {
+    SetDefaultData();
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_HOST;
+    myparms.dstMemoryType = CU_MEMORYTYPE_ARRAY;
+#else
+    myparms.srcMemoryType = hipMemoryTypeHost;
+    myparms.dstMemoryType = hipMemoryTypeArray;
+#endif
+    myparms.srcHost = hData;
+    myparms.srcPitch = width * sizeof(T);
+    myparms.srcHeight = height;
+    myparms.dstArray = arr;
+    HIP_CHECK(hipDrvMemcpy3DAsync(&myparms, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+    // Array to Array
+    SetDefaultData();
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_ARRAY;
+    myparms.dstMemoryType = CU_MEMORYTYPE_ARRAY;
+#else
+    myparms.srcMemoryType = hipMemoryTypeArray;
+    myparms.dstMemoryType = hipMemoryTypeArray;
+#endif
+    myparms.srcArray = arr;
+    myparms.dstArray = arr1;
+    HIP_CHECK(hipDrvMemcpy3DAsync(&myparms, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+    T *hOutputData = reinterpret_cast<T*>(malloc(size));
+    memset(hOutputData, 0,  size);
+    SetDefaultData();
+    // Device to host
+#if HT_NVIDIA
+    myparms.srcMemoryType = CU_MEMORYTYPE_ARRAY;
+    myparms.dstMemoryType = CU_MEMORYTYPE_HOST;
+#else
+    myparms.srcMemoryType = hipMemoryTypeArray;
+    myparms.dstMemoryType = hipMemoryTypeHost;
+#endif
+    myparms.srcArray = arr1;
+    myparms.dstHost = hOutputData;
+    myparms.dstPitch = width * sizeof(T);
+    myparms.dstHeight = height;
+    HIP_CHECK(hipDrvMemcpy3DAsync(&myparms, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+
+    HipTest::checkArray(hData, hOutputData, width, height, depth);
+    free(hOutputData);
+  }
+  DeAllocateMemory();
+}
+
+/* DeAllocating the memory */
+template <typename T>
+void DrvMemcpy3DAsync<T>::DeAllocateMemory() {
+  HIP_CHECK(hipArrayDestroy(arr));
+  HIP_CHECK(hipArrayDestroy(arr1));
+  HIP_CHECK(hipStreamDestroy(stream));
+  free(hData);
+}
+
+/* Verifying hipDrvMemcpy3DAsync API Host to Array for different datatypes */
+TEMPLATE_TEST_CASE("Unit_hipDrvMemcpy3DAsync_MultipleDataTypes", "",
+                   uint8_t, int, float) {
+  for (int i = 1; i < 25; i++) {
+    if (std::is_same<TestType, float>::value) {
+      DrvMemcpy3DAsync<TestType> memcpy3d_float(i, i, i,
+          HIP_AD_FORMAT_FLOAT);
+      memcpy3d_float.HostArray_DrvMemcpy3DAsync();
+    } else if (std::is_same<TestType, uint8_t>::value)  {
+      DrvMemcpy3DAsync<TestType> memcpy3d_intx(i, i, i,
+          HIP_AD_FORMAT_UNSIGNED_INT8);
+      memcpy3d_intx.HostArray_DrvMemcpy3DAsync();
+    } else if (std::is_same<TestType, int>::value)  {
+      DrvMemcpy3DAsync<TestType> memcpy3d_inty(i, i, i,
+          HIP_AD_FORMAT_SIGNED_INT32);
+      memcpy3d_inty.HostArray_DrvMemcpy3DAsync();
+    }
+  }
+}
+
+/* This testcase verifies H2D copy of hipDrvMemcpy3DAsync API */
+TEST_CASE("Unit_hipDrvMemcpy3DAsync_HosttoDevice") {
+  DrvMemcpy3DAsync<float> memcpy3d_D2H_float(10, 10, 1, HIP_AD_FORMAT_FLOAT);
+  memcpy3d_D2H_float.HostDevice_DrvMemcpy3DAsync();
+}
+
+/* This testcase verifies negative scenarios of hipDrvMemcpy3DAsync API */
+#if HT_NVIDIA
+TEST_CASE("Unit_hipDrvMemcpy3DAsync_Negative") {
+  DrvMemcpy3DAsync<float> memcpy3d(10, 10, 1, HIP_AD_FORMAT_FLOAT);
+  memcpy3d.NegativeTests();
+}
+#endif
+
+/* This testcase verifies extent validation scenarios of
+   hipDrvMemcpy3DAsync API */
+TEST_CASE("Unit_hipDrvMemcpy3DAsync_ExtentValidation") {
+  DrvMemcpy3DAsync<float> memcpy3d(10, 10, 1, HIP_AD_FORMAT_FLOAT);
+  memcpy3d.Extent_Validation();
+}
+
+/* This testcase verifies H2D copy in device context
+change scenario for hipDrvMemcpy3DAsync API */
+#if HT_AMD
+TEST_CASE("Unit_hipDrvMemcpy3DAsync_H2DDeviceContextChange") {
+  int numDevices = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    DrvMemcpy3DAsync<float> memcpy3d(10, 10, 1, HIP_AD_FORMAT_FLOAT);
+    memcpy3d.HostDevice_DrvMemcpy3DAsync(true);
+  } else {
+    SUCCEED("skipped testcase as Device count is < 2");
+  }
+}
+
+
+/* This testcase verifies Host to Array copy in device context
+change scenario for hipDrvMemcpy3DAsync API */
+TEST_CASE("Unit_hipDrvMemcpy3DAsync_Host2ArrayDeviceContextChange") {
+  int numDevices = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    DrvMemcpy3DAsync<float> memcpy3d(10, 10, 10, HIP_AD_FORMAT_FLOAT);
+    memcpy3d.HostArray_DrvMemcpy3DAsync(true);
+  } else {
+    SUCCEED("skipped testcase as Device count is < 2");
+  }
+}
+#endif
+
+

tests/catch/unit/memory/hipDrvPtrGetAttributes.cc

@@ -0,0 +1,176 @@
+/*
+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.
+*/
+
+/*
+ *  Tests for hipDrvPointerGetAttributes API
+ Functional Tests:
+   1. Pass multiple device related attributes in the attributes of hipDrvPointerGetAttributes API
+      for the device pointer and check the behaviour
+   2. Pass device and host attributes in the attributes of hipDrvPointerGetAttributes API and validate the behaviour
+   3. Pass invalid pointer to hipDrvPointerGetAttributes API and validate the behaviour.
+
+ Negative Tests:
+  1. Pass invalid numAttributes
+  2. Pass nullptr to attributes
+  3. Pass nullptr to data
+  4. Pass nullptr to device pointer
+*/
+#include <hip_test_common.hh>
+
+static size_t Nbytes = 0;
+constexpr size_t N {1000000};
+
+/* This testcase verifies Negative Scenarios of
+ * hipDrvPointerGetAttributes API */
+TEST_CASE("Unit_hipDrvPtrGetAttributes_Negative") {
+  HIP_CHECK(hipSetDevice(0));
+  Nbytes = N * sizeof(int);
+  int deviceId;
+  int numDevices = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  int* A_d;
+  int* A_Pinned_h;
+
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_Pinned_h), Nbytes,
+                         hipHostMallocDefault));
+  HIP_CHECK(hipGetDevice(&deviceId));
+  unsigned int device_ordinal;
+  int *dev_ptr;
+  void *data[2];
+  data[0] = &dev_ptr;
+  data[1] = &device_ordinal;
+
+  hipPointer_attribute attributes[] = {HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,
+                                       HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL};
+
+  SECTION("Passing nullptr to attributes") {
+    REQUIRE(hipDrvPointerGetAttributes(2, nullptr, data,
+            reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to data") {
+    REQUIRE(hipDrvPointerGetAttributes(2, attributes, nullptr,
+              reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorInvalidValue);
+  }
+
+#if HT_AMD
+  SECTION("Passing nullptr to device Pointer") {
+    hipDeviceptr_t ptr = 0;
+    REQUIRE(hipDrvPointerGetAttributes(2, attributes, data,
+            ptr) == hipErrorInvalidValue);
+  }
+#endif
+#if HT_NVIDIA
+  SECTION("Passing invalid dependencies") {
+    hipPointer_attribute attributes1[] = {HIP_POINTER_ATTRIBUTE_DEVICE_POINTER};
+    REQUIRE(hipDrvPointerGetAttributes(2, attributes1, data,
+            reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorInvalidValue);
+  }
+#endif
+}
+
+// Testcase verifies functional scenarios of hipDrvPointerGetAttributes API
+TEST_CASE("Unit_hipDrvPtrGetAttributes_Functional") {
+  HIP_CHECK(hipSetDevice(0));
+  Nbytes = N * sizeof(int);
+  int deviceId;
+  int numDevices = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  int* A_d;
+  int* A_Pinned_h;
+
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_Pinned_h), Nbytes,
+                         hipHostMallocDefault));
+  HIP_CHECK(hipGetDevice(&deviceId));
+
+  SECTION("Passing device attributes to device pointer") {
+    unsigned int memory_type;
+    int device_ordinal;
+    int *dev{nullptr};
+    int *dev_ptr{nullptr};
+    int *dev_ptr1{nullptr};
+    unsigned int range_size;
+    int *start_addr{nullptr};
+    void *data[5];
+    data[0] = (&memory_type);
+    data[1] = (&device_ordinal);
+    data[2] = (&dev_ptr);
+    data[3] = (&range_size);
+    data[4] = (&start_addr);
+
+    // Device memory
+    hipPointer_attribute attributes[] = {HIP_POINTER_ATTRIBUTE_MEMORY_TYPE,
+                                         HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL,
+                                         HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,
+                                         HIP_POINTER_ATTRIBUTE_RANGE_SIZE,
+                                        HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR};
+    HIP_CHECK(hipPointerGetAttribute(&dev_ptr1,
+              HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,
+              reinterpret_cast<hipDeviceptr_t>(A_d + 100)));
+    HIP_CHECK(hipPointerGetAttribute(&dev,
+	      HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,
+              reinterpret_cast<hipDeviceptr_t>(A_d)));
+    HIP_CHECK(hipDrvPointerGetAttributes(5, attributes, data,
+              reinterpret_cast<hipDeviceptr_t>(A_d + 100)));
+
+    REQUIRE(dev_ptr == dev_ptr1);
+#if HT_NVIDIA
+    REQUIRE(memory_type == CU_MEMORYTYPE_DEVICE);
+#else
+    REQUIRE(memory_type == hipMemoryTypeDevice);
+#endif
+    REQUIRE(device_ordinal == deviceId);
+    REQUIRE(range_size == Nbytes);
+    REQUIRE(start_addr == dev);
+  }
+
+  SECTION("Passing device and host attributes to device pointer") {
+    int device_ordinal;
+    int *host_ptr;
+    void *data[2];
+    data[0] = (&host_ptr);
+    data[1] = (&device_ordinal);
+
+    hipPointer_attribute attributes[] = {HIP_POINTER_ATTRIBUTE_HOST_POINTER,
+                                         HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL};
+    HIP_CHECK(hipDrvPointerGetAttributes(2, attributes, data,
+              reinterpret_cast<hipDeviceptr_t>(A_d)));
+    REQUIRE(host_ptr == nullptr);
+    REQUIRE(device_ordinal == deviceId);
+  }
+
+  SECTION("Passing host related attributes to host pointer") {
+    int device_ordinal;
+    void *data[2];
+    int *host_ptr;
+    data[0] = (&host_ptr);
+    data[1] = (&device_ordinal);
+    hipPointer_attribute attributes[] = {HIP_POINTER_ATTRIBUTE_HOST_POINTER,
+                                         HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL};
+    HIP_CHECK(hipDrvPointerGetAttributes(2, attributes, data,
+              reinterpret_cast<hipDeviceptr_t>(A_Pinned_h)));
+    REQUIRE(host_ptr == A_Pinned_h);
+    REQUIRE(device_ordinal == deviceId);
+  }
+}

tests/catch/unit/memory/hipMalloc3D.cc

@@ -0,0 +1,133 @@
+/*
+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.
+*/
+
+/*
+hipMalloc3D API test scenarios
+1. Basic Functionality
+2. Negative Scenarios
+3. Allocating Small and big chunk data
+4. Multithreaded scenario
+*/
+
+#include <hip_test_common.hh>
+static constexpr auto SMALL_SIZE{4};
+static constexpr auto CHUNK_LOOP{100};
+static constexpr auto BIG_SIZE{100};
+/*
+This API verifies hipMalloc3D API by allocating memory in smaller chunks for
+CHUNK_LOOP iterations and checks for the memory leaks by get the memory
+info before and after the hipMalloc3D API and the difference should
+match with the allocated memory
+*/
+static void MemoryAlloc3DDiffSizes(int gpu) {
+  HIPCHECK(hipSetDevice(gpu));
+  std::vector<size_t> array_size;
+  array_size.push_back(SMALL_SIZE);
+  array_size.push_back(BIG_SIZE);
+  for (auto &sizes : array_size) {
+    size_t width = sizes * sizeof(float);
+    size_t height{sizes}, depth{sizes};
+    hipPitchedPtr devPitchedPtr[CHUNK_LOOP];
+    hipExtent extent = make_hipExtent(width, height, depth);
+    size_t tot, avail, ptot, pavail;
+    HIPCHECK(hipMemGetInfo(&pavail, &ptot));
+    for (int i = 0; i < CHUNK_LOOP; i++) {
+      HIPCHECK(hipMalloc3D(&devPitchedPtr[i], extent));
+    }
+    for (int i = 0; i < CHUNK_LOOP; i++) {
+      HIPCHECK(hipFree(devPitchedPtr[i].ptr));
+    }
+    HIPCHECK(hipMemGetInfo(&avail, &tot));
+    if ((pavail != avail)) {
+      HIPASSERT(false);
+    }
+  }
+}
+
+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
+ */
+TEST_CASE("Unit_hipMalloc3D_Basic") {
+  size_t width = SMALL_SIZE * sizeof(char);
+  size_t height{SMALL_SIZE}, depth{SMALL_SIZE};
+  hipPitchedPtr devPitchedPtr;
+  hipExtent extent = make_hipExtent(width, height, depth);
+  REQUIRE(hipMalloc3D(&devPitchedPtr, extent) == hipSuccess);
+}
+
+/*
+This testcase verifies the hipMalloc3D API by allocating
+smaller and big chunk data.
+*/
+TEST_CASE("Unit_hipMalloc3D_SmallandBigChunks") {
+  MemoryAlloc3DDiffSizes(0);
+}
+
+/*
+This testcase verifies the hipMalloc3D API in multithreaded
+scenario by launching threads in parallel on multiple GPUs
+and verifies the hipMalloc3D API with small and big chunks data
+*/
+TEST_CASE("Unit_hipMalloc3D_MultiThread") {
+  std::vector<std::thread> threadlist;
+  int devCnt = 0;
+
+  devCnt = HipTest::getDeviceCount();
+
+  size_t tot, avail, ptot, pavail;
+  HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
+  for (int i = 0; i < devCnt; i++) {
+    threadlist.push_back(std::thread(Malloc3DThreadFunc, i));
+  }
+
+  for (auto &t : threadlist) {
+    t.join();
+  }
+  HIP_CHECK(hipMemGetInfo(&avail, &tot));
+
+  if (pavail != avail) {
+    WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
+    REQUIRE(false);
+  }
+}

tests/catch/unit/memory/hipMalloc3DArray.cc

@@ -0,0 +1,195 @@
+/*
+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.
+*/
+
+/*
+hipMalloc3DArray API test scenarios
+1. Basic Functionality
+2. Negative Scenarios
+3. Allocating Small and big chunk data
+4. Multithreaded scenario
+*/
+
+
+
+#include <hip_test_common.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.
+ * Two scenarios are verified in this API
+ * 1. SmallArray: Allocates ARRAY_SIZE in a loop and
+ *    releases the memory and verifies the meminfo.
+ * 2. BigArray: Allocates BIG_ARRAY_SIZE in a loop and
+ *    releases the memory and verifies the meminfo
+ *
+ * In both cases, the memory info before allocation and
+ * after releasing the memory should be the same
+ *
+ */
+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);
+    }
+  }
+}
+
+/* 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);
+  }
+
+  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));
+}
+
+TEST_CASE("Unit_hipMalloc3DArray_DiffSizes") {
+  Malloc3DArray_DiffSizes(0);
+}
+/*
+This testcase verifies the hipMalloc3DArray API in multithreaded
+scenario by launching threads in parallel on multiple GPUs
+and verifies the hipMalloc3DArray API with small and big chunks data
+*/
+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));
+  for (int i = 0; i < devCnt; i++) {
+    threadlist.push_back(std::thread(Malloc3DArrayThreadFunc, i));
+  }
+
+  for (auto &t : threadlist) {
+    t.join();
+  }
+  HIP_CHECK(hipMemGetInfo(&avail, &tot));
+
+  if (pavail != avail) {
+    WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
+    REQUIRE(false);
+  }
+}
+

tests/catch/unit/memory/hipMallocArray.cc

@@ -0,0 +1,172 @@
+/*
+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.
+*/
+
+/*
+hipMallocArray API test scenarios
+1. Basic Functionality
+2. Negative Scenarios
+3. Allocating Small and big chunk data
+4. Multithreaded scenario
+*/
+
+
+#include <hip_test_common.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};
+
+/*
+ * This API verifies  memory allocations for small and
+ * bigger chunks of data.
+ * Two scenarios are verified in this API
+ * 1. NUM_W(small Data): Allocates NUM_W*NUM_H in a loop and
+ *    releases the memory and verifies the meminfo.
+ * 2. BIGNUM_W(big data): Allocates BIGNUM_W*BIGNUM_H in a loop and
+ *    releases the memory and verifies the meminfo
+ *
+ * In both cases, the memory info before allocation and
+ * after releasing the memory should be the same
+ *
+ */
+
+static void MallocArray_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) {
+    hipArray* A_d[ARRAY_LOOP];
+    size_t tot, avail, ptot, pavail;
+    hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
+    HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
+    for (int i = 0; i < ARRAY_LOOP; i++) {
+      if (size == NUM_W) {
+        HIP_CHECK(hipMallocArray(&A_d[i], &desc,
+              NUM_W, NUM_H,
+              hipArrayDefault));
+      } else {
+        HIP_CHECK(hipMallocArray(&A_d[i], &desc,
+              BIGNUM_W, BIGNUM_H,
+              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);
+    }
+  }
+}
+
+static void MallocArrayThreadFunc(int gpu) {
+  MallocArray_DiffSizes(gpu);
+}
+
+/*
+ * This testcase verifies the negative scenarios of
+ * hipMallocArray API
+ */
+TEST_CASE("Unit_hipMallocArray_Negative") {
+  hipArray* A_d;
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
+#if HT_NVIDIA
+  SECTION("NullPointer to Array") {
+    REQUIRE(hipMallocArray(nullptr, &desc,
+            NUM_W, NUM_H, hipArrayDefault) != hipSuccess);
+  }
+
+  SECTION("NullPointer to Channel Descriptor") {
+    REQUIRE(hipMallocArray(&A_d, nullptr,
+            NUM_W, NUM_H, hipArrayDefault) != hipSuccess);
+  }
+#endif
+  SECTION("Width 0 in hipMallocArray") {
+    REQUIRE(hipMallocArray(&A_d, &desc,
+            0, NUM_H, hipArrayDefault) != hipSuccess);
+  }
+
+  SECTION("Height 0 in hipMallocArray") {
+    REQUIRE(hipMallocArray(&A_d, &desc,
+            NUM_W, 0, hipArrayDefault) == hipSuccess);
+  }
+
+  SECTION("Invalid Flag") {
+    REQUIRE(hipMallocArray(&A_d, &desc,
+            NUM_W, NUM_H, 100) != hipSuccess);
+  }
+
+  SECTION("Max int values") {
+    REQUIRE(hipMallocArray(&A_d, &desc,
+            std::numeric_limits<int>::max(),
+            std::numeric_limits<int>::max(),
+            hipArrayDefault) != hipSuccess);
+  }
+}
+/*
+ * This testcase verifies the basic scenario of
+ * hipMallocArray API for different datatypes
+ * of size 10
+ */
+TEMPLATE_TEST_CASE("Unit_hipMallocArray_Basic",
+                   "", int, unsigned int, float) {
+  hipArray* A_d;
+  hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
+  REQUIRE(hipMallocArray(&A_d, &desc,
+                         NUM_W, NUM_H,
+                         hipArrayDefault) == hipSuccess);
+  HIP_CHECK(hipFreeArray(A_d));
+}
+
+
+TEST_CASE("Unit_hipMallocArray_DiffSizes") {
+  MallocArray_DiffSizes(0);
+}
+
+
+/*
+This testcase verifies the hipMallocArray API in multithreaded
+scenario by launching threads in parallel on multiple GPUs
+and verifies the hipMallocArray API with small and big chunks data
+*/
+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));
+  for (int i = 0; i < devCnt; i++) {
+    threadlist.push_back(std::thread(MallocArrayThreadFunc, i));
+  }
+
+  for (auto &t : threadlist) {
+    t.join();
+  }
+  HIP_CHECK(hipMemGetInfo(&avail, &tot));
+
+  if (pavail != avail) {
+    WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
+    REQUIRE(false);
+  }
+}
+

tests/catch/unit/memory/hipMallocMngdMultiThread.cc

@@ -0,0 +1,546 @@
+/*
+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.
+*/
+#include <hip_test_common.hh>
+#include <atomic>
+
+
+// Kernel functions
+__global__ void HmmMultiThread(int n, float *x, float *y) {
+  int index = blockIdx.x * blockDim.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  for (int i = index; i < n; i += stride)
+    y[i] = x[i] * x[i];
+}
+
+__global__ void KrnlWth2MemTypes(int *Hmm, int *Dptr, size_t n) {
+  size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+  for (size_t i = index; i < n; i++) {
+    Hmm[i] = Dptr[i] + 10;
+  }
+}
+
+__global__ void KernelMul_MngdMem123(int *Hmm, int *Dptr, size_t n) {
+  size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+  size_t stride = blockDim.x * gridDim.x;
+  for (size_t i = index; i < n; i += stride) {
+    Hmm[i] = Dptr[i] * 10;
+  }
+}
+
+
+
+// The following variable is used to determine the failure of test case
+static bool  IfTestPassed = true;
+
+static void LaunchKrnl(int *Hmm1, size_t NumElms, int InitVal, int GpuOrdnl,
+                       int AdviseFlg) {
+  int *Hmm2 = NULL;
+  hipStream_t strm;
+  HIPCHECK(hipSetDevice(GpuOrdnl));
+  HIPCHECK(hipStreamCreate(&strm));
+  if (AdviseFlg == 0) {
+    HIPCHECK(hipMemAdvise(Hmm1 , NumElms * sizeof(int),
+                           hipMemAdviseSetReadMostly, GpuOrdnl));
+  } else if (AdviseFlg == 1) {
+    HIPCHECK(hipMemAdvise(Hmm1 , NumElms * sizeof(int),
+                           hipMemAdviseSetPreferredLocation, GpuOrdnl));
+  } else if (AdviseFlg == 2) {
+    HIPCHECK(hipMemAdvise(Hmm1 , NumElms * sizeof(int),
+                           hipMemAdviseSetAccessedBy, GpuOrdnl));
+  } else if (AdviseFlg == 3) {
+    HIPCHECK(hipMemPrefetchAsync(Hmm1, NumElms * sizeof(int), GpuOrdnl, strm));
+    HIPCHECK(hipStreamSynchronize(strm));
+  }
+  HIPCHECK(hipMallocManaged(&Hmm2, (sizeof(int) * NumElms)));
+  for (int i = 0; i < 2; ++i) {
+    KrnlWth2MemTypes<<<((NumElms + 63)/64), 64, 0, strm>>>(Hmm2, Hmm1, NumElms);
+    HIPCHECK(hipStreamSynchronize(strm));
+  }
+  // Verifying the result
+  int DataMismatch = 0;
+  for (size_t i = 0; i < NumElms; ++i) {
+    if (Hmm2[i] != (InitVal + 10)) {
+      DataMismatch++;
+    }
+  }
+  if (DataMismatch != 0) {
+    WARN("Data Mismatch observed at line: " << __LINE__);
+    IfTestPassed = false;
+  }
+}
+
+static void LaunchKrnl2(int *Hmm, size_t NumElms, int InitVal, int HmmMem) {
+  int *ptr = nullptr, blockSize = 64, *HstPtr = nullptr;
+  hipStream_t strm;
+  HIPCHECK(hipStreamCreate(&strm));
+  if (HmmMem == 0) {
+    HstPtr = reinterpret_cast<int*>(new int[NumElms]);
+    HIPCHECK(hipMalloc(&ptr, (sizeof(int) * NumElms)));
+  } else {
+    HIPCHECK(hipMallocManaged(&ptr, (sizeof(int) * NumElms)));
+  }
+  dim3 dimBlock(blockSize, 1, 1);
+  dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
+  for (int i = 0; i < 2; ++i) {
+    KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(ptr, Hmm, NumElms);
+  }
+  HIPCHECK(hipStreamSynchronize(strm));
+  // Verifying the result
+  int DataMismatch = 0;
+  if (HmmMem == 0) {
+    HIPCHECK(hipMemcpy(HstPtr, ptr, (sizeof(int) * NumElms),
+                        hipMemcpyDeviceToHost));
+    for (size_t i = 0; i < NumElms; ++i) {
+      if (HstPtr[i] != (InitVal + 10)) {
+        DataMismatch++;
+      }
+    }
+  } else {
+    for (size_t i = 0; i < NumElms; ++i) {
+      if (ptr[i] != (InitVal + 10)) {
+        DataMismatch++;
+      }
+    }
+  }
+  if (DataMismatch != 0) {
+    INFO("Data Mismatch observed at line: " << __LINE__);
+    REQUIRE(false);
+  }
+}
+
+static void LaunchKrnl3(int *Dptr, size_t NumElms, int InitVal) {
+  int *Hmm = NULL, blockSize = 64;
+  hipStream_t strm;
+  HIPCHECK(hipStreamCreate(&strm));
+  HIPCHECK(hipMallocManaged(&Hmm, (sizeof(int) * NumElms)));
+  dim3 dimBlock(blockSize, 1, 1);
+  dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
+  for (int i = 0; i < 2; ++i) {
+    KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
+  }
+  HIPCHECK(hipStreamSynchronize(strm));
+  // Verifying the result
+  int DataMismatch = 0;
+  for (size_t i = 0; i < NumElms; ++i) {
+    if (Hmm[i] != (InitVal + 10)) {
+      DataMismatch++;
+    }
+  }
+  if (DataMismatch != 0) {
+    INFO("Data Mismatch observed at line: " << __LINE__);
+    REQUIRE(false);
+  }
+}
+
+
+static void LaunchKrnl5(int *Hmm1, size_t NumElms, int InitVal,
+                        int KerneltoLaunch) {
+  int *Hmm2 = NULL, blockSize = 64;
+  hipStream_t strm;
+  HIPCHECK(hipStreamCreate(&strm));
+  HIPCHECK(hipMallocManaged(&Hmm2, (sizeof(int) * NumElms)));
+  dim3 dimBlock(blockSize, 1, 1);
+  dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
+  for (int i = 0; i < 2; ++i) {
+    if (KerneltoLaunch == 0) {
+      KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm2, Hmm1, NumElms);
+    } else {
+      KernelMul_MngdMem123<<<dimGrid, dimBlock, 0, strm>>>(Hmm2, Hmm1, NumElms);
+    }
+  }
+  HIPCHECK(hipStreamSynchronize(strm));
+  // Verifying the result
+  int DataMismatch = 0;
+  if (KerneltoLaunch == 0) {
+    for (size_t i = 0; i < NumElms; ++i) {
+      if (Hmm2[i] != (InitVal + 10)) {
+        DataMismatch++;
+      }
+    }
+  } else {
+    for (size_t i = 0; i < NumElms; ++i) {
+      if (Hmm2[i] != (InitVal * 10)) {
+        DataMismatch++;
+      }
+    }
+  }
+  if (DataMismatch != 0) {
+    INFO("Data Mismatch observed at line: " << __LINE__);
+    REQUIRE(false);
+  }
+}
+
+
+static void TestFlagParamGlobal(int dev) {
+  std::atomic<int> DataMismatch{0};
+  int NUM_ELMS = 4096, ITERATIONS = 10;
+  float *HmmAG = NULL, INIT_VAL = 2.5;
+  float *Ad = NULL, *Ah = NULL;
+  Ah = new float[NUM_ELMS];
+  hipStream_t strm;
+  HIPCHECK(hipSetDevice(dev));
+  HIPCHECK(hipStreamCreate(&strm));
+  // Testing hipMemAttachGlobal Flag
+  HIPCHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float),
+                            hipMemAttachGlobal));
+
+  // Initializing HmmAG memory
+  for (int i = 0; i < NUM_ELMS; i++) {
+    HmmAG[i] = INIT_VAL;
+    Ah[i] = 0;
+  }
+
+  int blockSize = 256;
+  int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize;
+  dim3 dimGrid(numBlocks, 1, 1);
+  dim3 dimBlock(blockSize, 1, 1);
+  HIPCHECK(hipSetDevice(dev));
+  HIPCHECK(hipMalloc(&Ad, NUM_ELMS * sizeof(float)));
+  HIPCHECK(hipMemset(Ad, 0, NUM_ELMS * sizeof(float)));
+  for (int i = 0; i < ITERATIONS; ++i) {
+    HmmMultiThread<<<dimGrid, dimBlock, 0, strm>>>(NUM_ELMS, HmmAG, Ad);
+    HIPCHECK(hipStreamSynchronize(strm));
+  }
+  HIPCHECK(hipMemcpy(Ah, Ad, NUM_ELMS * sizeof(float), hipMemcpyDeviceToHost));
+  for (int j = 0; j < NUM_ELMS; ++j) {
+    if (Ah[j] != (INIT_VAL * INIT_VAL)) {
+      DataMismatch++;
+      break;
+    }
+  }
+  if (DataMismatch != 0) {
+    INFO("Data Mismatch observed when kernel launched on device: " << dev);
+    IfTestPassed = false;
+  }
+
+  HIPCHECK(hipFree(Ad));
+  delete[] Ah;
+  HIPCHECK(hipFree(HmmAG));
+  HIPCHECK(hipStreamDestroy(strm));
+}
+
+
+static void TestFlagParamHost(int dev) {
+  std::atomic<int> DataMismatch{0};
+  float *HmmAH1 = nullptr, *HmmAH2 = nullptr, INIT_VAL = 2.5;
+  int NUM_ELMS = 4096, ITERATIONS = 10;
+  hipStream_t strm;
+  HIPCHECK(hipSetDevice(dev));
+  HIPCHECK(hipStreamCreate(&strm));
+  HIPCHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float),
+                            hipMemAttachHost));
+  HIPCHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float),
+                            hipMemAttachHost));
+  // Initializing HmmAH memory
+  for (int i = 0; i < NUM_ELMS; i++) {
+    HmmAH1[i] = INIT_VAL;
+    HmmAH2[i] = 0;
+  }
+  int blockSize = 256;
+  int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize;
+  dim3 dimGrid(numBlocks, 1, 1);
+  dim3 dimBlock(blockSize, 1, 1);
+  for (int i = 0; i < ITERATIONS; ++i) {
+    HmmMultiThread<<<dimGrid, dimBlock, 0, strm>>>(NUM_ELMS, HmmAH1, HmmAH2);
+    HIPCHECK(hipStreamSynchronize(strm));
+  }
+  for (int j = 0; j < NUM_ELMS; ++j) {
+    if (HmmAH2[j] != (INIT_VAL * INIT_VAL)) {
+      IfTestPassed = false;
+      DataMismatch++;
+      break;
+    }
+  }
+  if (DataMismatch != 0) {
+    INFO("Data Mismatch observed when kernel launched on device: " << dev);
+    IfTestPassed = false;
+  }
+  HIPCHECK(hipFree(HmmAH1));
+  HIPCHECK(hipFree(HmmAH2));
+  HIPCHECK(hipStreamDestroy(strm));
+}
+
+static void AllocateHmmMemory(int flag, int device) {
+  int ITERATIONS = 10;
+  void *HmmAG = NULL, *HmmAH = NULL;
+  HIPCHECK(hipSetDevice(device));
+  for (int i = 0; i < ITERATIONS; ++i) {
+    if (!flag) {
+      HIPCHECK(hipMallocManaged(&HmmAG, (2 * 4096), hipMemAttachGlobal));
+      HIPCHECK(hipFree(HmmAG));
+    } else {
+      HIPCHECK(hipMallocManaged(&HmmAH, (2 * 4096), hipMemAttachHost));
+      HIPCHECK(hipFree(HmmAH));
+    }
+  }
+}
+
+
+static int HmmAttrPrint() {
+  int managed = 0;
+  INFO("The following are the attribute values related to HMM for"
+         " device 0:\n");
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+              hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
+  INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                 hipDeviceAttributeConcurrentManagedAccess, 0));
+  INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+              hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
+  INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
+         << managed);
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+  return managed;
+}
+
+TEST_CASE("Unit_hipMallocManaged_MultiThread") {
+  IfTestPassed = true;
+  int NumDevs = 0, managed = 0, ATTACH_GLOBAL = 0, ATTACH_HOST = 1;
+  int ITERATIONS = 10;
+  managed = HmmAttrPrint();
+  if (managed) {
+    HIP_CHECK(hipGetDeviceCount(&NumDevs));
+    std::vector<std::thread> T1;
+    std::vector<std::thread> T2;
+    for (int i = 0; i < NumDevs; ++i) {
+      for (int j = 0; j < ITERATIONS; ++j) {
+        T1.push_back(std::thread(TestFlagParamGlobal, i));
+        T2.push_back(std::thread(AllocateHmmMemory, ATTACH_GLOBAL, i));
+      }
+      for (auto &t1 : T1) {
+        if (t1.joinable()) {
+          t1.join();
+        }
+      }
+      for (auto &t2 : T2) {
+        if (t2.joinable()) {
+          t2.join();
+        }
+      }
+    }
+    T1.clear();
+    T2.clear();
+    for (int i = 0; i < NumDevs; ++i) {
+      for (int j = 0; j < ITERATIONS; ++j) {
+        T1.push_back(std::thread(TestFlagParamHost, i));
+        T2.push_back(std::thread(AllocateHmmMemory, ATTACH_HOST, i));
+      }
+      for (auto &t1 : T1) {
+        if (t1.joinable()) {
+          t1.join();
+        }
+      }
+      for (auto &t2 : T2) {
+        if (t2.joinable()) {
+          t2.join();
+        }
+      }
+    }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory"
+            "attribute. Hence skipping the testing with Pass result.\n");
+  }
+  REQUIRE(IfTestPassed);
+}
+
+// The following test checks what happens when same Hmm memory is used to
+// launch multiple threads over multiple gpus
+TEST_CASE("Unit_hipMallocManaged_MGpuMThread") {
+  IfTestPassed = true;
+  int Ngpus = 0;
+  HIP_CHECK(hipGetDeviceCount(&Ngpus));
+  if (Ngpus < 2) {
+    WARN("This test needs atleast 2 or more gpus, but the system");
+    WARN(" has only " << Ngpus);
+    WARN(" gpus. Hence skipping the test.");
+    SUCCEED("\n");
+  }
+  int managed = HmmAttrPrint();
+  if (managed == 1) {
+    int InitVal = 123, *Hmm1 = NULL, NumElms = 4096*4;
+    HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int))));
+    for (int i = 0; i < NumElms; ++i) {
+      Hmm1[i] = InitVal;
+    }
+
+    std::vector<std::thread> Thrds;
+    // AdviseFlg=0 for ReadMostly to be applied
+    // AdviseFlg=1 for PreferredLocation to be applied
+    // AdviseFlg=2 for AccessedBy to be applied
+    // AdviseFlg=3 to prefetch the memory to particular gpu
+    for (int AdviseFlg = 0; AdviseFlg < 4; ++AdviseFlg) {
+      for (int i = 0; i < Ngpus; ++i) {
+        Thrds.push_back(std::thread(LaunchKrnl, Hmm1, NumElms, InitVal, i,
+                                    AdviseFlg));
+      }
+      for (auto &thr : Thrds) {
+        if (thr.joinable()) {
+          thr.join();
+        }
+      }
+    }
+    REQUIRE(IfTestPassed);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+
+// The following test checks what happens when multiple kernels are launched
+// with same Hmm memory
+TEST_CASE("Unit_hipMallocManaged_MultiKrnlComnHmm") {
+  IfTestPassed = true;
+  int managed = HmmAttrPrint();
+  if (managed == 1) {
+    int InitVal = 123, *Hmm = NULL, NumElms = 1024*4, TotThrds = 2;
+    int HmmMem2 = 0, *HstPtr = nullptr;  //  to indicate the thread that
+    //  hipMalloc() memory has to be used
+    HstPtr = reinterpret_cast<int*>(new int[NumElms]);
+    HIP_CHECK(hipMalloc(&Hmm, (NumElms * sizeof(int))));
+    for (int i = 0; i < NumElms; ++i) {
+      HstPtr[i] = InitVal;
+    }
+    HIP_CHECK(hipMemcpy(Hmm, HstPtr, (NumElms * sizeof(int)),
+                        hipMemcpyHostToDevice));
+    std::vector<std::thread> Thrds;
+    for (int i = 0; i < TotThrds; ++i) {
+      Thrds.push_back(std::thread(LaunchKrnl2, Hmm, NumElms, InitVal, HmmMem2));
+    }
+
+    for (auto &thr : Thrds) {
+      if (thr.joinable()) {
+        thr.join();
+      }
+    }
+    delete[] HstPtr;
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+
+// The following test checks what happens when multiple kernels are launched
+// with same hipMalloc() memory
+TEST_CASE("Unit_hipMallocManaged_MultiKrnlComnMalloc") {
+  IfTestPassed = true;
+  int managed = HmmAttrPrint();
+  if (managed) {
+    int InitVal = 123, *Dptr = NULL, NumElms = 4096*8, TotThrds = 2;
+    int *HstPtr = reinterpret_cast<int*>(new int[NumElms]);
+    HIP_CHECK(hipMalloc(&Dptr, (NumElms * sizeof(int))));
+    for (int i = 0; i < NumElms; ++i) {
+      HstPtr[i] = InitVal;
+    }
+    HIP_CHECK(hipMemcpy(Dptr, HstPtr, (NumElms * sizeof(int)),
+                        hipMemcpyHostToDevice));
+    std::vector<std::thread> Thrds;
+    for (int i = 0; i < TotThrds; ++i) {
+      Thrds.push_back(std::thread(LaunchKrnl3, Dptr, NumElms, InitVal));
+    }
+
+    for (auto &thr : Thrds) {
+      if (thr.joinable()) {
+        thr.join();
+      }
+    }
+    delete[] HstPtr;
+    HIP_CHECK(hipFree(Dptr));
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+//  The following section tests the scenario wherein multiple threads use their
+//  own stream to launch kernel on common Hmm memory
+TEST_CASE("Unit_hipMallocManaged_MultiThrdMultiStrm") {
+  IfTestPassed = true;
+  int managed = HmmAttrPrint();
+  if (managed == 1) {
+    int NumElms = 4096*4;
+    int *Hmm1 = NULL, TotlThrds = 4, InitVal = 123;
+    int HmmMem = 1;  //  to indicate the thread that Hmm memory need to be
+    //  used inside it
+    HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int))));
+    for (int i = 0; i < NumElms; ++i) {
+      Hmm1[i] = InitVal;
+    }
+    std::vector<std::thread> Thrds;
+    for (int i = 0; i < TotlThrds; ++i) {
+      Thrds.push_back(std::thread(LaunchKrnl2, Hmm1, NumElms, InitVal, HmmMem));
+    }
+
+    for (auto &thr : Thrds) {
+      if (thr.joinable()) {
+        thr.join();
+      }
+    }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+
+
+//  The following section tests the scenario wherein two threads each use
+//  different kernel but common HMM memory
+TEST_CASE("Unit_hipMallocManaged_TwoKrnlsComnHmmMem") {
+  IfTestPassed = true;
+  int managed = HmmAttrPrint();
+  if (managed == 1) {
+    int InitVal = 123, *Dptr = NULL, NumElms = 4096*4, TotThrds = 2;
+    int *HstPtr = reinterpret_cast<int*>(new int[NumElms]);
+    HIP_CHECK(hipMalloc(&Dptr, (NumElms * sizeof(int))));
+    for (int i = 0; i < NumElms; ++i) {
+      HstPtr[i] = InitVal;
+    }
+    HIP_CHECK(hipMemcpy(Dptr, HstPtr, (NumElms * sizeof(int)),
+                        hipMemcpyHostToDevice));
+    std::vector<std::thread> Thrds;
+    for (int i = 0; i < TotThrds; ++i) {
+      Thrds.push_back(std::thread(LaunchKrnl5, Dptr, NumElms, InitVal, i));
+    }
+
+    for (auto &thr : Thrds) {
+      if (thr.joinable()) {
+        thr.join();
+      }
+    }
+    delete[] HstPtr;
+    HIP_CHECK(hipFree(Dptr));
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+

tests/catch/unit/memory/hipMallocPitch.cc

@@ -0,0 +1,298 @@
+/*
+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.
+*/
+
+/*
+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 <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+static constexpr auto SMALLCHUNK_NUMW{4};
+static constexpr auto SMALLCHUNK_NUMH{4};
+static constexpr auto LARGECHUNK_NUMW{1025};
+static constexpr auto LARGECHUNK_NUMH{1000};
+static constexpr auto NUM_W{10};
+static constexpr auto NUM_H{10};
+static constexpr auto COLUMNS{8};
+static constexpr auto ROWS{8};
+static constexpr auto CHUNK_LOOP{100};
+
+
+template<typename T>
+__global__ void copy_var(T* A, T* B,
+                              size_t ROWS, size_t pitch_A) {
+  for (uint64_t i = 0; i< ROWS*pitch_A; i= i+pitch_A) {
+    A[i] = B[i];
+  }
+}
+template<typename T>
+static bool validateResult(T* A, T* B, size_t pitch_A) {
+  bool testResult = true;
+  for (uint64_t i=0; i < pitch_A*ROWS; i=i+pitch_A) {
+    if (A[i] != B[i]) {
+      testResult = false;
+      break;
+    }
+  }
+  return testResult;
+}
+/*
+ * This API verifies  memory allocations for small and
+ * bigger chunks of data.
+ * Two scenarios are verified in this API
+ * 1. SmallChunk: Allocates SMALLCHUNK_NUMW in a loop and
+ *    releases the memory and verifies the meminfo.
+ * 2. LargeChunk: Allocates LARGECHUNK_NUMW in a loop and
+ *    releases the memory and verifies the meminfo
+ *
+ * In both cases, the memory info before allocation and
+ * after releasing the memory should be the same
+ *
+ */
+template<typename T>
+static void MemoryAllocDiffSizes(int gpu) {
+  HIP_CHECK(hipSetDevice(gpu));
+  std::vector<size_t> array_size;
+  array_size.push_back(SMALLCHUNK_NUMH);
+  array_size.push_back(LARGECHUNK_NUMH);
+  for (auto &sizes : array_size) {
+    T* A_d[CHUNK_LOOP];
+    size_t pitch_A;
+    size_t width;
+    if (sizes == SMALLCHUNK_NUMH) {
+      width = SMALLCHUNK_NUMW * sizeof(T);
+    } else {
+      width = LARGECHUNK_NUMW * sizeof(T);
+    }
+    size_t tot, avail, ptot, pavail;
+    HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
+    for (int i = 0; i < CHUNK_LOOP; i++) {
+      HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d[i]),
+            &pitch_A, width, sizes));
+    }
+    for (int i = 0; i < CHUNK_LOOP; i++) {
+      HIP_CHECK(hipFree(A_d[i]));
+    }
+    HIP_CHECK(hipMemGetInfo(&avail, &tot));
+    if (pavail != avail) {
+      HIPASSERT(false);
+    }
+  }
+}
+
+/*Thread Function */
+static void threadFunc(int gpu) {
+  MemoryAllocDiffSizes<float>(gpu);
+}
+/*
+ * This testcase verifies the negative scenarios of hipMallocPitch API
+ */
+#if 0 //TODO: Review, fix and re-enable test
+TEST_CASE("Unit_hipMallocPitch_Negative") {
+  float* A_d;
+  size_t pitch_A;
+  size_t width{NUM_W * sizeof(float)};
+#if HT_NVIDIA
+  SECTION("NullPtr to Pitched Ptr") {
+    REQUIRE(hipMallocPitch(nullptr,
+            &pitch_A, width, NUM_H) != hipSuccess);
+  }
+
+  SECTION("nullptr to pitch") {
+    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
+                           nullptr, width, NUM_H) != hipSuccess);
+  }
+#endif
+  SECTION("Width 0 in hipMallocPitch") {
+    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
+                           &pitch_A, 0, NUM_H) == hipSuccess);
+  }
+
+  SECTION("Height 0 in hipMallocPitch") {
+    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
+                           &pitch_A, width, 0) == hipSuccess);
+  }
+
+  SECTION("Max int values") {
+    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
+                           &pitch_A, std::numeric_limits<int>::max(),
+                           std::numeric_limits<int>::max()) != hipSuccess);
+  }
+}
+#endif
+/*
+ * This testcase verifies the basic scenario of
+ * hipMallocPitch API for different datatypes
+ *
+ */
+TEMPLATE_TEST_CASE("Unit_hipMallocPitch_Basic",
+    "[hipMallocPitch]", int, unsigned int, float) {
+  TestType* A_d;
+  size_t pitch_A;
+  size_t width{NUM_W * sizeof(TestType)};
+  REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
+          &pitch_A, width, NUM_H) == hipSuccess);
+  HIP_CHECK(hipFree(A_d));
+}
+
+/*
+ * This testcase verifies hipMallocPitch API for small
+ * and big chunks of data.
+ */
+TEMPLATE_TEST_CASE("Unit_hipMallocPitch_SmallandBigChunks",
+    "[hipMallocPitch]", int, unsigned int, float) {
+  MemoryAllocDiffSizes<TestType>(0);
+}
+
+/*
+ * This testcase verifies the memory allocated by hipMallocPitch API
+ * by performing Memcpy2D on the allocated memory.
+ */
+TEMPLATE_TEST_CASE("Unit_hipMallocPitch_Memcpy2D", ""
+                   , int, float, double) {
+  HIP_CHECK(hipSetDevice(0));
+  TestType  *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, *A_d{nullptr},
+            *B_d{nullptr};
+  size_t pitch_A, pitch_B;
+  size_t width{NUM_W * sizeof(TestType)};
+
+  // Allocating memory
+  HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
+                                  &A_h, &B_h, &C_h, NUM_W*NUM_H, false);
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
+                          &pitch_A, width, NUM_H));
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&B_d),
+                          &pitch_B, width, NUM_H));
+
+  // Initialize the data
+  HipTest::setDefaultData<TestType>(NUM_W*NUM_H, A_h, B_h, C_h);
+
+  // Host to Device
+  HIP_CHECK(hipMemcpy2D(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
+                        COLUMNS*sizeof(TestType), ROWS, hipMemcpyHostToDevice));
+
+  // Performs D2D on same GPU device
+  HIP_CHECK(hipMemcpy2D(B_d, pitch_B, A_d,
+                        pitch_A, COLUMNS*sizeof(TestType),
+                        ROWS, hipMemcpyDeviceToDevice));
+
+  // hipMemcpy2D Device to Host
+  HIP_CHECK(hipMemcpy2D(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
+                        COLUMNS*sizeof(TestType), ROWS,
+                        hipMemcpyDeviceToHost));
+
+  // Validating the result
+  REQUIRE(HipTest::checkArray<TestType>(A_h, B_h, COLUMNS, ROWS) == true);
+
+
+  // DeAllocating the memory
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipFree(B_d));
+  HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr,
+                                A_h, B_h, C_h, false);
+}
+
+
+
+/*
+This testcase verifies the hipMallocPitch API in multithreaded
+scenario by launching threads in parallel on multiple GPUs
+and verifies the hipMallocPitch API with small and big chunks data
+*/
+
+TEST_CASE("Unit_hipMallocPitch_MultiThread", "") {
+  std::vector<std::thread> threadlist;
+  int devCnt = 0;
+
+  devCnt = HipTest::getDeviceCount();
+
+  size_t tot, avail, ptot, pavail;
+  HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
+  for (int i = 0; i < devCnt; i++) {
+    threadlist.push_back(std::thread(threadFunc, i));
+  }
+
+  for (auto &t : threadlist) {
+    t.join();
+  }
+  HIP_CHECK(hipMemGetInfo(&avail, &tot));
+
+  if (pavail != avail) {
+    WARN("Memory leak of hipMallocPitch API in multithreaded scenario");
+    REQUIRE(false);
+  }
+}
+
+/*
+ * This testcase verifies hipMallocPitch API by
+ *  1. Allocating Memory using hipMallocPitch API
+ *  2. Launching the kernel and copying the data from the allocated kernel
+ *     variable to another kernel variable.
+ *  3. Validating the result
+ */
+TEMPLATE_TEST_CASE("Unit_hipMallocPitch_KernelLaunch", ""
+                   , int, float, double) {
+  HIP_CHECK(hipSetDevice(0));
+  TestType  *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, *A_d{nullptr},
+            *B_d{nullptr};
+  size_t pitch_A, pitch_B;
+  size_t width{NUM_W * sizeof(TestType)};
+
+  // Allocating memory
+  HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
+                                  &A_h, &B_h, &C_h, NUM_W*NUM_H, false);
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
+                          &pitch_A, width, NUM_H));
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&B_d),
+                          &pitch_B, width, NUM_H));
+
+  // Host to Device
+  HIP_CHECK(hipMemcpy2D(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
+                        COLUMNS*sizeof(TestType), ROWS, hipMemcpyHostToDevice));
+
+
+  hipLaunchKernelGGL(copy_var<TestType>, dim3(1), dim3(1),
+          0, 0, static_cast<TestType*>(A_d),
+          static_cast<TestType*>(B_d), ROWS, pitch_A);
+
+
+  // hipMemcpy2D Device to Host
+  HIP_CHECK(hipMemcpy2D(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
+                        COLUMNS*sizeof(TestType), ROWS,
+                        hipMemcpyDeviceToHost));
+
+  // Validating the result
+  validateResult(A_h, B_h, pitch_A);
+
+  // DeAllocating the memory
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipFree(B_d));
+  HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr,
+                                A_h, B_h, C_h, false);
+}
+
+

tests/catch/unit/memory/hipMemCoherencyTst.cc

@@ -91,54 +91,39 @@ static void TstCoherency(int *Ptr, bool HmmMem) {
   }
 }
 
-static int HmmAttrPrint() {
-  int managed = 0;
-  INFO("The following are the attribute values related to HMM for"
-         " device 0:\n");
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-              hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
-  INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-                                 hipDeviceAttributeConcurrentManagedAccess, 0));
-  INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-                                 hipDeviceAttributePageableMemoryAccess, 0));
-  INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-              hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
-  INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
-         << managed);
-
-  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
-                                  0));
-  INFO("hipDeviceAttributeManagedMemory: " << managed);
-  return managed;
-}
-
-
 /* Test case description: The following test validates if fine grain
    behavior is observed or not with memory allocated using hipHostMalloc()*/
 // The following tests are disabled for Nvidia as they are not consistently
 // passing
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
-  int *Ptr = nullptr,  SIZE = sizeof(int);
+  int *Ptr = nullptr,  SIZE = sizeof(int), Pageable = 0;
   bool HmmMem = false;
   YES_COHERENT = false;
-  // Allocating hipHostMalloc() memory with hipHostMallocCoherent flag
-  SECTION("hipHostMalloc with hipHostMallocCoherent flag") {
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocCoherent));
-  }
-  SECTION("hipHostMalloc with Default flag") {
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
-  }
-  SECTION("hipHostMalloc with hipHostMallocMapped flag") {
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocMapped));
-  }
 
-  TstCoherency(Ptr, HmmMem);
-  HIP_CHECK(hipHostFree(Ptr));
-  REQUIRE(YES_COHERENT);
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  if (Pageable == 1) {
+    // Allocating hipHostMalloc() memory with hipHostMallocCoherent flag
+    SECTION("hipHostMalloc with hipHostMallocCoherent flag") {
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocCoherent));
+    }
+    SECTION("hipHostMalloc with Default flag") {
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
+    }
+    SECTION("hipHostMalloc with hipHostMallocMapped flag") {
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocMapped));
+    }
+
+    TstCoherency(Ptr, HmmMem);
+    HIP_CHECK(hipHostFree(Ptr));
+    REQUIRE(YES_COHERENT);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
+             "attribute. Hence skipping the test with Pass result.\n");
+  }
 }
 #endif
 
@@ -149,12 +134,19 @@ TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
 // passing
 #if HT_AMD
 TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
-  int *Ptr = nullptr, SIZE = sizeof(int);
+  int *Ptr = nullptr, SIZE = sizeof(int), Pageable = 0, managed = 0;
   bool HmmMem = true;
   YES_COHERENT = false;
 
-  int managed =  HmmAttrPrint();
-  if (managed == 1) {
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+
+  if (managed == 1 && Pageable == 1) {
     // Allocating hipMallocManaged() memory
     SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
       HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
@@ -166,8 +158,8 @@ TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
     HIP_CHECK(hipFree(Ptr));
     REQUIRE(YES_COHERENT);
   } else {
-    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-           "attribute. Hence skipping the testing with Pass result.\n");
+    SUCCEED("GPU 0 doesn't support ManagedMemory or PageableMemoryAccess"
+           "device attribute. Hence skipping the test with Pass result.\n");
   }
 }
 #endif
@@ -175,30 +167,40 @@ TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
 /* Test case description: The following test validates if memory access is fine
    with memory allocated using hipMallocManaged() and CoarseGrain Advise*/
 TEST_CASE("Unit_hipMallocManaged_CoherentTstWthAdvise") {
-  int *Ptr = nullptr, SIZE = sizeof(int);
+  int *Ptr = nullptr, SIZE = sizeof(int), managed = 0;
   YES_COHERENT = false;
-  // Allocating hipMallocManaged() memory
-  SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
-    HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
-  }
-  SECTION("hipMallocManaged with hipMemAttachHost flag") {
-    HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachHost));
-  }
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+
+  if (managed == 1) {
+    // Allocating hipMallocManaged() memory
+    SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
+      HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
+    }
+    SECTION("hipMallocManaged with hipMemAttachHost flag") {
+      HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachHost));
+    }
 #if HT_AMD
-  HIP_CHECK(hipMemAdvise(Ptr, SIZE, hipMemAdviseSetCoarseGrain, 0));
+    HIP_CHECK(hipMemAdvise(Ptr, SIZE, hipMemAdviseSetCoarseGrain, 0));
 #endif
-  // Initializing Ptr memory with 9
-  *Ptr = 9;
-  hipStream_t strm;
-  HIP_CHECK(hipStreamCreate(&strm));
-  SquareKrnl<<<1, 1, 0, strm>>>(Ptr);
-  HIP_CHECK(hipStreamSynchronize(strm));
-  if (*Ptr == 81) {
-    YES_COHERENT = true;
+    // Initializing Ptr memory with 9
+    *Ptr = 9;
+    hipStream_t strm;
+    HIP_CHECK(hipStreamCreate(&strm));
+    SquareKrnl<<<1, 1, 0, strm>>>(Ptr);
+    HIP_CHECK(hipStreamSynchronize(strm));
+    if (*Ptr == 81) {
+      YES_COHERENT = true;
+    }
+    HIP_CHECK(hipFree(Ptr));
+    HIP_CHECK(hipStreamDestroy(strm));
+    REQUIRE(YES_COHERENT);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+             "attribute. Hence skipping the test with Pass result.\n");
   }
-  HIP_CHECK(hipFree(Ptr));
-  HIP_CHECK(hipStreamDestroy(strm));
-  REQUIRE(YES_COHERENT);
 }
 
 
@@ -226,12 +228,18 @@ TEST_CASE("Unit_hipMalloc_CoherentTst") {
    hipExtMallocWithFlags()*/
 #if HT_AMD
 TEST_CASE("Unit_hipExtMallocWithFlags_CoherentTst") {
-  int *Ptr = nullptr, SIZE = sizeof(int), InitVal = 9;
+  int *Ptr = nullptr, SIZE = sizeof(int), InitVal = 9, Pageable = 0, managed = 0;
   bool FineGrain = true;
   YES_COHERENT = false;
 
-  int managed =  HmmAttrPrint();
-  if (managed == 1) {
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+  if (managed == 1 && Pageable == 1) {
     // Allocating hipExtMallocWithFlags() memory with flags
     SECTION("hipExtMallocWithFlags with hipDeviceMallocFinegrained flag") {
       HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&Ptr), SIZE*2,
@@ -264,8 +272,8 @@ TEST_CASE("Unit_hipExtMallocWithFlags_CoherentTst") {
     HIP_CHECK(hipFree(Ptr));
     REQUIRE(YES_COHERENT);
   } else {
-    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-           "attribute. Hence skipping the testing with Pass result.\n");
+    SUCCEED("GPU 0 doesn't support ManagedMemory or PageableMemoryAccess"
+           "device attribute. Hence skipping the test with Pass result.\n");
   }
 }
 #endif

tests/catch/unit/memory/hipMemPrefetchAsync.cc

@@ -0,0 +1,124 @@
+/*
+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.
+*/
+
+#include <hip_test_common.hh>
+// Kernel function
+__global__ void MemPrftchAsyncKernel(int* C_d, const int* A_d, size_t N) {
+  size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+  size_t stride = hipBlockDim_x * hipGridDim_x;
+  for (size_t i = offset; i < N; i += stride) {
+    C_d[i] = A_d[i] * A_d[i];
+  }
+}
+
+
+static int HmmAttrPrint() {
+  int managed = 0;
+  INFO("The following are the attribute values related to HMM for"
+         " device 0:\n");
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+              hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
+  INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                 hipDeviceAttributeConcurrentManagedAccess, 0));
+  INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+              hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
+  INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
+         << managed);
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+  return managed;
+}
+
+/*
+  Test Description: This test prefetches the memory to each of the available
+  devices and launch kernel followed by result verification
+  At the end the memory is prefetched to Host and kernel is launched followed
+  by result verification.
+*/
+
+TEST_CASE("Unit_hipMemPrefetchAsync") {
+  int MangdMem = HmmAttrPrint();
+  if (MangdMem == 1) {
+    bool IfTestPassed = true;
+    int A_CONST = 123, MEM_SIZE = (8192 * sizeof(int));
+    int *devPtr1 = NULL, *devPtr2 = NULL, NumDevs = 0, flag = 0;
+    hipStream_t strm;
+    HIP_CHECK(hipMallocManaged(&devPtr1, MEM_SIZE));
+    HIP_CHECK(hipMallocManaged(&devPtr2, MEM_SIZE));
+    HIP_CHECK(hipGetDeviceCount(&NumDevs));
+    // Initializing the memory
+    for (uint32_t k = 0; k < (MEM_SIZE/sizeof(int)); ++k) {
+      devPtr1[k] = A_CONST;
+      devPtr2[k] = 0;
+    }
+
+
+    for (int i = 0; i < NumDevs;  ++i) {
+      HIP_CHECK(hipSetDevice(i));
+      HIP_CHECK(hipStreamCreate(&strm));
+      HIP_CHECK(hipMemPrefetchAsync(devPtr1, MEM_SIZE, i, strm));
+      HIP_CHECK(hipStreamSynchronize(strm));
+      MemPrftchAsyncKernel<<<32, (MEM_SIZE/sizeof(int)/32)>>>(devPtr2, devPtr1,
+                                                       MEM_SIZE/sizeof(int));
+      for (uint32_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) {
+        if (devPtr1[m] != (A_CONST * A_CONST)) {
+          flag = 1;
+        }
+      }
+      HIP_CHECK(hipStreamDestroy(strm));
+      if (!flag) {
+        INFO("Test failed for device: " << i);
+        IfTestPassed = false;
+        flag = 0;
+      }
+    }
+    // The memory will be prefetched from last gpu in the system to the host
+    // memory and kernel is launched followed by result verification.
+    HIP_CHECK(hipStreamCreate(&strm));
+    HIP_CHECK(hipMemPrefetchAsync(devPtr1, MEM_SIZE, hipCpuDeviceId, strm));
+    HIP_CHECK(hipStreamSynchronize(strm));
+    MemPrftchAsyncKernel<<<32, (MEM_SIZE/sizeof(int)/32)>>>(devPtr2, devPtr1,
+                                                     MEM_SIZE/sizeof(int));
+    for (uint32_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) {
+      if (devPtr1[m] != (A_CONST * A_CONST)) {
+        flag = 1;
+      }
+    }
+    HIP_CHECK(hipStreamDestroy(strm));
+    if (!flag) {
+      INFO("Failed to prefetch the memory to System space.\n");
+      IfTestPassed = false;
+      flag = 0;
+    }
+
+    HIP_CHECK(hipFree(devPtr1));
+    HIP_CHECK(hipFree(devPtr2));
+    REQUIRE(IfTestPassed);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}

tests/catch/unit/memory/hipMemcpyParam2DAsync.cc

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
 in the Software without restriction, including without limitation the rights
@@ -393,6 +393,7 @@ TEST_CASE("Unit_hipMemcpyParam2DAsync_ExtentValidation") {
 
   // DeAllocating the Memory
   HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipStreamDestroy(stream));
   HipTest::freeArrays<char>(nullptr, nullptr, nullptr,
                                 A_h, B_h, C_h, false);
 }
@@ -483,6 +484,7 @@ TEST_CASE("Unit_hipMemcpyParam2DAsync_Negative") {
   // DeAllocating the memory
   HIP_CHECK(hipFree(A_d));
   HIP_CHECK(hipStreamSynchronize(stream));
+  HIP_CHECK(hipStreamDestroy(stream));
   HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
                                 A_h, B_h, C_h, false);
 }

tests/catch/unit/memory/hipPointerGetAttribute.cc

@@ -0,0 +1,355 @@
+/*
+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.
+*/
+/*
+Added Negative and Functional tests for hipPointerGetAttribute API
+
+Functional Scenarios:
+
+1. Allocate memory using different Allocation APIs and check whether
+   correct memory type and device oridinal are returned.
+
+2. Allocate device variable and get the pointer info by calling hipPointerGetAttribute API
+   with HIP_POINTER_ATTRIBUTE_DEVICE_POINTER/HIP_POINTER_ATTRIBUTE_START_ADDRESS
+   and Launch kernel with device variable and verify whether the pointer variable of
+   hipPointerGetAttribute is getting updated or not
+
+3. Allocate device memory in GPU-0 and get the pointer info in peer GPU
+
+4. Allocate device memory and get the buffer ID by calling
+   hipPointerGetAttribute API with HIP_POINTER_ATTRIBUTE_BUFFER_ID,
+   DeAllocate and Allocate the memory again and ensure that the buffer ID is unique
+
+5. Allocate host memory and get the device ordinal by calling
+   hipPointerGetAttribute API with HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL
+   and ensure that it matches with CUDA result(which returns 100)
+
+6. Allocate managed memory with different flags and trigger
+   hipPointerGetAttribute with the following flags HIP_POINTER_ATTRIBUTE_MAPPED and verify the behaviour
+*/
+
+#include <hip_test_common.hh>
+#include <string>
+static constexpr auto NUM_W{16};
+static constexpr auto NUM_H{16};
+static constexpr size_t N {10};
+#define INT_VAL 10
+#define VAL_DATA 99
+static __global__ void var_update(int* data) {
+  for (unsigned int i = 0; i < N; i++) {
+     data[i] = VAL_DATA;
+  }
+}
+
+/* Allocate memory using different Allocation APIs and check whether
+   correct memory type and device oridinal are returned */
+TEST_CASE("Unit_hipPointerGetAttribute_MemoryTypes") {
+  HIP_CHECK(hipSetDevice(0));
+  size_t pitch_A;
+  size_t width{NUM_W * sizeof(char)};
+  unsigned int datatype;
+  SECTION("Malloc Pitch Allocation") {
+    char *A_d;
+    HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
+          &pitch_A, width, NUM_H));
+    HIP_CHECK(hipPointerGetAttribute(&datatype,
+              HIP_POINTER_ATTRIBUTE_MEMORY_TYPE,
+              reinterpret_cast<hipDeviceptr_t>(A_d)));
+#if HT_NVIDIA
+    REQUIRE(datatype == CU_MEMORYTYPE_DEVICE);
+#else
+    REQUIRE(datatype == hipMemoryTypeDevice);
+#endif
+  }
+#if HT_AMD
+  SECTION("Malloc Array Allocation") {
+    hipArray *B_d;
+    hipChannelFormatDesc desc = hipCreateChannelDesc<char>();
+    HIP_CHECK(hipMallocArray(&B_d, &desc, NUM_W, NUM_H, hipArrayDefault));
+    HIP_CHECK(hipPointerGetAttribute(&datatype,
+                                     HIP_POINTER_ATTRIBUTE_MEMORY_TYPE,
+                                     reinterpret_cast<hipDeviceptr_t>(B_d)));
+#if HT_NVIDIA
+    REQUIRE(datatype == CU_MEMORYTYPE_ARRAY);
+#else
+    REQUIRE(datatype == hipMemoryTypeArray);
+#endif
+    HIP_CHECK(hipFreeArray(B_d));
+  }
+
+  SECTION("Malloc 3D Array Allocation") {
+    int width = 10, height = 10, depth = 10;
+    hipArray *arr;
+
+    hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8,
+        0, 0, 0, hipChannelFormatKindFloat);
+    HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height,
+            depth), hipArrayDefault));
+    HIP_CHECK(hipPointerGetAttribute(&datatype,
+                                     HIP_POINTER_ATTRIBUTE_MEMORY_TYPE,
+                                     reinterpret_cast<hipDeviceptr_t>(arr)));
+#if HT_NVIDIA
+    REQUIRE(datatype == CU_MEMORYTYPE_ARRAY);
+#else
+    REQUIRE(datatype == hipMemoryTypeArray);
+#endif
+    HIP_CHECK(hipFreeArray(arr));
+  }
+#endif
+}
+
+/*
+ * This testcase verifies the following scenario
+ * Initializes A_d with A_h and get pointer info using hipPointerGetAttribute
+ * The result of the API is passed to kernel for validation
+ * and modifies it in kernel.
+ * Validates the device variable to check whether the
+ * data is updated or not.
+ */
+TEST_CASE("Unit_hipPointerGetAttribute_KernelUpdation") {
+  HIP_CHECK(hipSetDevice(0));
+  size_t Nbytes = 0;
+  Nbytes = N * sizeof(int);
+  int* A_d, *A_h;
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  hipDeviceptr_t data = 0;
+  A_h = reinterpret_cast<int*>(malloc(Nbytes));
+  for (unsigned int i = 0; i < N; i++) {
+    A_h[i] = INT_VAL;
+  }
+  HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipPointerGetAttribute(&data,
+                                   HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,
+                                   reinterpret_cast<hipDeviceptr_t>(A_d)));
+  hipLaunchKernelGGL(var_update, dim3(1), dim3(1), 0, 0,
+                     reinterpret_cast<int *>(data));
+  HIP_CHECK(hipDeviceSynchronize());
+  HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
+  for (unsigned int i = 0; i < N; i++) {
+     REQUIRE(A_h[i] == VAL_DATA);
+  }
+  HIP_CHECK(hipFree(A_d));
+  free(A_h);
+}
+/*
+ * This testcase verifies the pointer info of device variable
+ * from peer GPU device.It validates the memory type and
+ * device ordinal in peer GPU
+ */
+TEST_CASE("Unit_hipPointerGetAttribute_PeerGPU") {
+  HIP_CHECK(hipSetDevice(0));
+  size_t Nbytes = 0;
+  Nbytes = N * sizeof(int);
+  int* A_d;
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  unsigned int data = 0;
+  int numDevices = 0;
+  int canAccessPeer = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
+    if (canAccessPeer) {
+      HIP_CHECK(hipSetDevice(1));
+      HIP_CHECK(hipPointerGetAttribute(&data,
+                HIP_POINTER_ATTRIBUTE_MEMORY_TYPE,
+                reinterpret_cast<hipDeviceptr_t>(A_d)));
+#if HT_NVIDIA
+      REQUIRE(data == CU_MEMORYTYPE_DEVICE);
+#else
+      REQUIRE(data == hipMemoryTypeDevice);
+#endif
+      HIP_CHECK(hipPointerGetAttribute(&data,
+                HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL,
+                reinterpret_cast<hipDeviceptr_t>(A_d)));
+      REQUIRE(data == 0);
+    } else {
+      SUCCEED("Machine does not seem to have P2P");
+    }
+  } else {
+    SUCCEED("skipped the testcase as no of devices is less than 2");
+  }
+  HIP_CHECK(hipFree(A_d));
+}
+
+/* Allocate device memory and get the buffer ID by calling
+   hipPointerGetAttribute API with HIP_POINTER_ATTRIBUTE_BUFFER_ID,
+   DeAllocate and Allocate the memory again and
+   ensure that the buffer ID is unique */
+TEST_CASE("Unit_hipPointerGetAttribute_BufferID") {
+  HIP_CHECK(hipSetDevice(0));
+  size_t Nbytes = 0;
+  Nbytes = N * sizeof(int);
+  int* A_d;
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  unsigned int bufid1, bufid2;
+  HIP_CHECK(hipPointerGetAttribute(&bufid1,
+            HIP_POINTER_ATTRIBUTE_BUFFER_ID,
+            reinterpret_cast<hipDeviceptr_t>(A_d)));
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  HIP_CHECK(hipPointerGetAttribute(&bufid2,
+            HIP_POINTER_ATTRIBUTE_BUFFER_ID,
+            reinterpret_cast<hipDeviceptr_t>(A_d)));
+  REQUIRE(bufid1 != bufid2);
+}
+
+/* Allocate host memory and get the device ordinal by calling
+   hipPointerGetAttribute API with HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL
+   and ensure that it matches with CUDA result
+*/
+#if HT_AMD
+TEST_CASE("Unit_hipPointerGetAttribute_HostDeviceOrdinal") {
+  size_t Nbytes = 0;
+  Nbytes = N * sizeof(int);
+  int* A_h;
+  unsigned int data = 0, data1 = 0;
+  A_h = reinterpret_cast<int*>(malloc(Nbytes));
+  for (unsigned int i = 0; i < N; i++) {
+    A_h[i] = INT_VAL;
+  }
+  REQUIRE(hipPointerGetAttribute(&data,
+          HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL,
+         reinterpret_cast<hipDeviceptr_t>(A_h)) == hipErrorInvalidValue);
+  REQUIRE(hipPointerGetAttribute(&data1,
+                                 HIP_POINTER_ATTRIBUTE_RANGE_SIZE,
+                                 reinterpret_cast<hipDeviceptr_t>(A_h))
+                                 == hipErrorInvalidValue);
+  free(A_h);
+}
+#endif
+
+/* Allocate managed memory with different flags and trigger
+   hipPointerGetAttribute with the following flags HIP_POINTER_ATTRIBUTE_MAPPED
+   and verify the behaviour */
+TEST_CASE("Unit_hipPointerGetAttribute_MappedMem") {
+  HIP_CHECK(hipSetDevice(0));
+  size_t Nbytes = 0;
+  Nbytes = N * sizeof(int);
+  int* A_d, *A_h;
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  A_h = reinterpret_cast<int*>(malloc(Nbytes));
+  for (unsigned int i = 0; i < N; i++) {
+    A_h[i] = INT_VAL;
+  }
+  HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+  int *ptr1 = 0, *ptr2 = 0;
+  unsigned int hostMalloc_mapped, mallocManaged;
+  HIP_CHECK(hipHostMalloc(&ptr1, Nbytes, hipHostMallocMapped));
+  HIP_CHECK(hipMallocManaged(&ptr2, Nbytes, hipMemAttachGlobal));
+  HIP_CHECK(hipPointerGetAttribute(&hostMalloc_mapped,
+                                   HIP_POINTER_ATTRIBUTE_MAPPED,
+                                   reinterpret_cast<hipDeviceptr_t>(A_d)));
+  HIP_CHECK(hipPointerGetAttribute(&mallocManaged,
+                                   HIP_POINTER_ATTRIBUTE_MAPPED,
+                                   reinterpret_cast<hipDeviceptr_t>(ptr2)));
+  REQUIRE(hostMalloc_mapped == 1);
+  REQUIRE(mallocManaged == 1);
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipHostFree(ptr1));
+  free(A_h);
+}
+
+/* This testcase verifies negative scenarios of hipPointerGetAttribute API */
+TEST_CASE("Unit_hipPointerGetAttribute_Negative") {
+  HIP_CHECK(hipSetDevice(0));
+  size_t Nbytes = 0;
+  constexpr size_t N {100};
+  Nbytes = N * sizeof(char);
+  char* A_d;
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  hipDeviceptr_t data = 0;
+  char *A_h;
+  A_h = reinterpret_cast<char*>(malloc(Nbytes));
+  SECTION("Pass nullptr to data") {
+    REQUIRE(hipPointerGetAttribute(nullptr,
+          HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,
+          reinterpret_cast<hipDeviceptr_t>(A_d))
+        == hipErrorInvalidValue);
+  }
+  SECTION("Pass nullptr to device attribute") {
+#if HT_AMD
+    REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,
+            nullptr) == hipErrorInvalidValue);
+#else
+    REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,
+          reinterpret_cast<hipDeviceptr_t>(nullptr)) == hipErrorInvalidValue);
+#endif
+  }
+  SECTION("DeAllocateMem and get the pointer info") {
+    char *B_d;
+    HIP_CHECK(hipMalloc(&B_d, Nbytes));
+    HIP_CHECK(hipFree(B_d));
+    REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,
+          reinterpret_cast<hipDeviceptr_t>(B_d)) == hipErrorInvalidValue);
+  }
+  SECTION("Get Start address of host pointer") {
+    char *A_h;
+    A_h = reinterpret_cast<char*>(malloc(Nbytes));
+    REQUIRE(hipPointerGetAttribute(&data,
+          HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR,
+          reinterpret_cast<hipDeviceptr_t>(A_h)) == hipErrorInvalidValue);
+  }
+  SECTION("Pass HIP_POINTER_ATTRIBUTE_HOST_POINTER to device pointer") {
+    REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_HOST_POINTER,
+          reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorInvalidValue);
+  }
+  SECTION("Pass BUFFER_ID attribute to host pointer") {
+    REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_BUFFER_ID,
+          reinterpret_cast<hipDeviceptr_t>(A_h))
+        == hipErrorInvalidValue);
+  }
+  SECTION("Pass invalid attribute") {
+    hipPointer_attribute attr{HIP_POINTER_ATTRIBUTE_DEVICE_POINTER};
+    REQUIRE(hipPointerGetAttribute(&data, attr,
+          reinterpret_cast<hipDeviceptr_t>(A_h)) == hipErrorInvalidValue);
+  }
+#if HT_AMD
+  SECTION("Pass HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE"
+      "not supported by HIP") {
+    REQUIRE(hipPointerGetAttribute(&data,
+          HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE,
+          reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorNotSupported);
+  }
+  SECTION("Pass HIP_POINTER_ATTRIBUTE_MEMPOOL_HANDLE not supported by HIP") {
+    REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_MEMPOOL_HANDLE,
+          reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorNotSupported);
+  }
+  SECTION("Pass HIP_POINTER_ATTRIBUTE_CONTEXT not supported by HIP") {
+    REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_CONTEXT,
+          reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorNotSupported);
+  }
+  SECTION("Pass HIP_POINTER_ATTRIBUTE_P2P_TOKENS  not supported by HIP") {
+    REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_P2P_TOKENS,
+          reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorNotSupported);
+  }
+  SECTION("Pass HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE"
+      "not supported by HIP") {
+    REQUIRE(hipPointerGetAttribute(&data,
+          HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE,
+          reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorNotSupported);
+  }
+  SECTION("Pass HIP_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES"
+      "not supported by HIP") {
+    REQUIRE(hipPointerGetAttribute(&data,
+          HIP_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES,
+          reinterpret_cast<hipDeviceptr_t>(A_d)) == hipErrorNotSupported);
+  }
+#endif
+  HIP_CHECK(hipFree(A_d));
+  free(A_h);
+}