SWDEV-467854 - [catch2][dtest] Functional testcases for hipGetProcAddress API - Part 2

Change-Id: I4619f8e69b75d1193e98adce2707bf7876b1ba59

catch/unit/device/CMakeLists.txt

@@ -33,6 +33,9 @@ set(TEST_SRC
 if(HIP_PLATFORM MATCHES "amd" AND BUILD_SHARED_LIBS)
   set(AMD_SRC
       hipGetProcAddressDevMgmt.cc)
+  if(UNIX)
+    set(AMD_SRC ${AMD_SRC} hipGetProcAddressIpcApis.cc)
+  endif()
   set(TEST_SRC ${TEST_SRC} ${AMD_SRC})
 endif()
 

catch/unit/device/hipGetProcAddressDevMgmt.cc

@@ -19,6 +19,7 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip_test_helper.hh>
 #include <hip_test_defgroups.hh>
+#include <hip/hip_deprecated.h>
 
 /**
  * @addtogroup hipGetProcAddress hipGetProcAddress
@@ -50,29 +51,33 @@ void CreateMemPool(int device, hipMemPool_t &mem_pool) {
  *  - HIP_VERSION >= 6.2
  */
 TEST_CASE("Unit_hipGetProcAddress_ValidateDeviceApis") {
-  void* hipGetDeviceCount_ptr;
-  void* hipRuntimeGetVersion_ptr;
-  void* hipDeviceGetLimit_ptr;
-  void* hipDeviceSetLimit_ptr;
-  void* hipDeviceComputeCapability_ptr;
-  void* hipDeviceGet_ptr;
-  void* hipDeviceGetPCIBusId_ptr;
-  void* hipDeviceGetByPCIBusId_ptr;
-  void* hipDeviceGetDefaultMemPool_ptr;
-  void* hipDeviceGetName_ptr;
-  void* hipDeviceGetUuid_ptr;
-  void* hipGetDeviceFlags_ptr;
-  void* hipSetDeviceFlags_ptr;
-  void* hipDeviceReset_ptr;
-  void* hipDriverGetVersion_ptr;
-  void* hipDeviceGetCacheConfig_ptr;
-  void* hipDeviceSetCacheConfig_ptr;
-  void* hipDeviceTotalMem_ptr;
-  void* hipGetDeviceProperties_ptr;
-  void* hipChooseDevice_ptr;
-  void* hipDeviceSetSharedMemConfig_ptr;
-  void* hipDeviceGetSharedMemConfig_ptr;
-  void* hipDeviceGetAttribute_ptr;
+  void* hipGetDeviceCount_ptr = nullptr;
+  void* hipRuntimeGetVersion_ptr = nullptr;
+  void* hipDeviceGetLimit_ptr = nullptr;
+  void* hipDeviceSetLimit_ptr = nullptr;
+  void* hipDeviceComputeCapability_ptr = nullptr;
+  void* hipDeviceGet_ptr = nullptr;
+  void* hipDeviceGetPCIBusId_ptr = nullptr;
+  void* hipDeviceGetByPCIBusId_ptr = nullptr;
+  void* hipDeviceGetDefaultMemPool_ptr = nullptr;
+  void* hipDeviceGetName_ptr = nullptr;
+  void* hipDeviceGetUuid_ptr = nullptr;
+  void* hipGetDeviceFlags_ptr = nullptr;
+  void* hipSetDeviceFlags_ptr = nullptr;
+  void* hipDeviceReset_ptr = nullptr;
+  void* hipDriverGetVersion_ptr = nullptr;
+  void* hipDeviceGetCacheConfig_ptr = nullptr;
+  void* hipDeviceSetCacheConfig_ptr = nullptr;
+  void* hipDeviceTotalMem_ptr = nullptr;
+  void* hipGetDeviceProperties_ptr = nullptr;
+  void* hipGetDevicePropertiesR0000_ptr = nullptr;
+  void* hipGetDevicePropertiesR0600_ptr = nullptr;
+  void* hipChooseDevice_ptr = nullptr;
+  void* hipChooseDeviceR0000_ptr = nullptr;
+  void* hipChooseDeviceR0600_ptr = nullptr;
+  void* hipDeviceSetSharedMemConfig_ptr = nullptr;
+  void* hipDeviceGetSharedMemConfig_ptr = nullptr;
+  void* hipDeviceGetAttribute_ptr = nullptr;
 
   int currentHipVersion = 0;
   HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
@@ -134,8 +139,20 @@ TEST_CASE("Unit_hipGetProcAddress_ValidateDeviceApis") {
   HIP_CHECK(hipGetProcAddress("hipGetDeviceProperties",
                               &hipGetDeviceProperties_ptr,
                               currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipGetDevicePropertiesR0000",
+                              &hipGetDevicePropertiesR0000_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipGetDevicePropertiesR0600",
+                              &hipGetDevicePropertiesR0600_ptr,
+                              currentHipVersion, 0, nullptr));
   HIP_CHECK(hipGetProcAddress("hipChooseDevice", &hipChooseDevice_ptr,
                               currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipChooseDeviceR0000",
+                              &hipChooseDeviceR0000_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipChooseDeviceR0600",
+                              &hipChooseDeviceR0600_ptr,
+                              currentHipVersion, 0, nullptr));
   HIP_CHECK(hipGetProcAddress("hipDeviceSetSharedMemConfig",
                               &hipDeviceSetSharedMemConfig_ptr,
                               currentHipVersion, 0, nullptr));
@@ -201,9 +218,23 @@ TEST_CASE("Unit_hipGetProcAddress_ValidateDeviceApis") {
   hipError_t (*dyn_hipGetDeviceProperties_ptr)(hipDeviceProp_t*, int) =
   reinterpret_cast<hipError_t (*)(hipDeviceProp_t*, int)>
   (hipGetDeviceProperties_ptr);
+  hipError_t(*dyn_hipGetDevicePropertiesR0000_ptr)
+            (hipDeviceProp_tR0000*, int) =
+  reinterpret_cast<hipError_t (*)(hipDeviceProp_tR0000*, int)>
+  (hipGetDevicePropertiesR0000_ptr);
+  hipError_t(*dyn_hipGetDevicePropertiesR0600_ptr)
+            (hipDeviceProp_tR0600*, int) =
+  reinterpret_cast<hipError_t (*)(hipDeviceProp_tR0600*, int)>
+  (hipGetDevicePropertiesR0600_ptr);
   hipError_t (*dyn_hipChooseDevice_ptr)(int*, hipDeviceProp_t*) =
   reinterpret_cast<hipError_t (*)(int*, hipDeviceProp_t*)>
   (hipChooseDevice_ptr);
+  hipError_t (*dyn_hipChooseDeviceR0000_ptr)(int*, hipDeviceProp_tR0000*) =
+  reinterpret_cast<hipError_t (*)(int*, hipDeviceProp_tR0000*)>
+  (hipChooseDevice_ptr);
+  hipError_t (*dyn_hipChooseDeviceR0600_ptr)(int*, hipDeviceProp_tR0600*) =
+  reinterpret_cast<hipError_t (*)(int*, hipDeviceProp_tR0600*)>
+  (hipChooseDevice_ptr);
   hipError_t (*dyn_hipDeviceSetSharedMemConfig_ptr)(hipSharedMemConfig) =
   reinterpret_cast<hipError_t (*)(hipSharedMemConfig)>
   (hipDeviceSetSharedMemConfig_ptr);
@@ -336,12 +367,37 @@ TEST_CASE("Unit_hipGetProcAddress_ValidateDeviceApis") {
   HIP_CHECK(hipGetDeviceProperties(&prop, 0));
   HIP_CHECK(dyn_hipGetDeviceProperties_ptr(&prop_ptr, 0));
   REQUIRE(prop.major == prop_ptr.major);
+
+  // hipGetDevicePropertiesR0000 API
+  hipDeviceProp_tR0000 propR0000, propR0000_ptr;
+  HIP_CHECK(hipGetDevicePropertiesR0000(&propR0000, 0));
+  HIP_CHECK(dyn_hipGetDevicePropertiesR0000_ptr(&propR0000_ptr, 0));
+  REQUIRE(propR0000.major == propR0000_ptr.major);
+
+  // hipGetDevicePropertiesR0600 API
+  hipDeviceProp_tR0600 propR0600, propR0600_ptr;
+  HIP_CHECK(hipGetDevicePropertiesR0600(&propR0600, 0));
+  HIP_CHECK(dyn_hipGetDevicePropertiesR0600_ptr(&propR0600_ptr, 0));
+  REQUIRE(propR0600.major == propR0600_ptr.major);
+
   // hipChooseDevice API
   int dev, dev_ptr;
   HIP_CHECK(hipChooseDevice(&dev, &prop));
   HIP_CHECK(dyn_hipChooseDevice_ptr(&dev_ptr, &prop));
   REQUIRE(dev == dev_ptr);
 
+  // hipChooseDeviceR0000 API
+  int devR0000, devR0000_ptr;
+  HIP_CHECK(hipChooseDeviceR0000(&devR0000, &propR0000));
+  HIP_CHECK(dyn_hipChooseDeviceR0000_ptr(&devR0000_ptr, &propR0000));
+  REQUIRE(devR0000 == devR0000_ptr);
+
+  // hipChooseDeviceR0600 API
+  int devR0600, devR0600_ptr;
+  HIP_CHECK(hipChooseDeviceR0600(&devR0600, &propR0600));
+  HIP_CHECK(dyn_hipChooseDeviceR0600_ptr(&devR0600_ptr, &propR0600));
+  REQUIRE(devR0600 == devR0600_ptr);
+
   // hipDeviceSetSharedMemConfig API
   HIP_CHECK(hipDeviceSetSharedMemConfig(hipSharedMemBankSizeFourByte));
   HIP_CHECK(dyn_hipDeviceSetSharedMemConfig_ptr(hipSharedMemBankSizeFourByte));
@@ -376,11 +432,11 @@ TEST_CASE("Unit_hipGetProcAddress_ValidateDeviceApis") {
  */
 
 TEST_CASE("Unit_hipGetProcAddress_PeerDeviceAccessAPIs") {
-  void* hipDeviceCanAccessPeer_ptr;
-  void* hipSetDevice_ptr;
-  void* hipGetDevice_ptr;
-  void* hipDeviceEnablePeerAccess_ptr;
-  void* hipDeviceDisablePeerAccess_ptr;
+  void* hipDeviceCanAccessPeer_ptr = nullptr;
+  void* hipSetDevice_ptr = nullptr;
+  void* hipGetDevice_ptr = nullptr;
+  void* hipDeviceEnablePeerAccess_ptr = nullptr;
+  void* hipDeviceDisablePeerAccess_ptr = nullptr;
 
   int currentHipVersion = 0;
   HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
@@ -467,8 +523,8 @@ bool CheckMemPoolSupport(const int device) {
 }
 
 TEST_CASE("Unit_hipGetProcAddress_SetGetMemPoolAPIs") {
-  void* hipDeviceSetMemPool_ptr;
-  void* hipDeviceGetMemPool_ptr;
+  void* hipDeviceSetMemPool_ptr = nullptr;
+  void* hipDeviceGetMemPool_ptr = nullptr;
   int currentHipVersion = 0;
   HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
   HIP_CHECK(hipGetProcAddress("hipDeviceSetMemPool",

catch/unit/device/hipGetProcAddressHelpers.hh

@@ -0,0 +1,229 @@
+/*
+Copyright (c) 2024 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.
+*/
+
+#ifndef CATCH_UNIT_DEVICE_HIPGETPROCADDRESSHELPERS_HH_
+#define CATCH_UNIT_DEVICE_HIPGETPROCADDRESSHELPERS_HH_
+
+#include <hip_test_common.hh>
+
+/**
+ * Local Function to fill the array with given value
+ */
+void fillHostArray(int *arr, int size, int value) {
+  for ( int i = 0; i < size; i++ ) {
+    arr[i] = value;
+  }
+}
+
+/**
+ * Local Function to validate the array with given reference value
+ */
+bool validateHostArray(int *arr, int size, int refValue) {
+  for ( int i = 0; i < size; i++ ) {
+    if ( arr[i] != refValue ) {
+      return false;
+    }
+  }
+  return true;
+}
+
+/**
+ * Local Function to fill the character array with given value
+ */
+void fillCharHostArray(char *arr, int size, int value) {
+  for ( int i = 0; i < size; i++ ) {
+    arr[i] = value;
+  }
+}
+
+/**
+ * Local Function to validate the array with given reference value
+ */
+bool validateCharHostArray(char *arr, int size, int refValue) {
+  for ( int i = 0; i < size; i++ ) {
+    if ( arr[i] != refValue ) {
+      return false;
+    }
+  }
+  return true;
+}
+
+/**
+ * Kernel to validate the array with given reference value
+ */
+__global__ void verifyArray(int *arr, int size, int refValue, int* status) {
+  for ( int i = 0; i < size; i++ ) {
+    if ( arr[i] != refValue ) {
+      *status = 0;
+      return;
+    }
+  }
+  *status = 1;
+  return;
+}
+
+/**
+ * Local Function to validate the device array with given reference value
+ */
+bool validateDeviceArray(int *arr, int size, int refValue) {
+  int *devStatus = nullptr;
+  HIP_CHECK(hipMalloc(&devStatus, sizeof(int)));
+  REQUIRE(devStatus != nullptr);
+
+  verifyArray<<<1, 1>>>(arr, size, refValue, devStatus);
+  int status;
+  HIP_CHECK(hipMemcpy(&status, devStatus, sizeof(int), hipMemcpyDeviceToHost));
+
+  HIP_CHECK(hipFree(devStatus));
+
+  if ( status == 1 ) {
+    return true;
+  } else {
+    return false;
+  }
+}
+
+/**
+ * Kernel to validate the character array with given reference value
+ */
+__global__ void verifyCharArray(char *arr, int size,
+                                int refValue, int* status) {
+  for ( int i = 0; i < size; i++ ) {
+    if ( arr[i] != refValue ) {
+      *status = 0;
+      return;
+    }
+  }
+  *status = 1;
+  return;
+}
+
+/**
+ * Local Function to validate the character device array with
+ * given reference value
+ */
+bool validateCharDeviceArray(char *arr, int size, int refValue) {
+  int *devStatus = nullptr;
+  HIP_CHECK(hipMalloc(&devStatus, sizeof(int)));
+  REQUIRE(devStatus != nullptr);
+
+  verifyCharArray<<< 1, 1 >>>(arr, size, refValue, devStatus);
+  int status;
+  HIP_CHECK(hipMemcpy(&status, devStatus, sizeof(int), hipMemcpyDeviceToHost));
+
+  HIP_CHECK(hipFree(devStatus));
+
+  if ( status == 1 ) {
+    return true;
+  } else {
+    return false;
+  }
+}
+
+/**
+ * Kernel to fill the array with given value
+ */
+__global__ void fillArray(int *arr, int size, int value) {
+  for ( int i = 0; i < size; i++ ) {
+    arr[i] = value;
+  }
+}
+
+/**
+ * Local Function to fill the device array with given value
+ */
+void fillDeviceArray(int *arr, int size, int value) {
+  fillArray<<<1, 1>>>(arr, size, value);
+}
+
+/**
+ * Local Function to validate the array of different types
+ */
+template<class T>
+bool validateArrayT(T *arr, int size, T value) {
+  for ( int i = 0; i < size; i++ ) {
+    if ( arr[i] != value ) {
+      return false;
+    }
+  }
+  return true;
+}
+
+/**
+ * Kernel to increment 1 value to all the elements in array
+ */
+__global__ void addOneKernel(int *a, int size) {
+  int offset = blockDim.x * blockIdx.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  for (int i = offset; i < size; i+= stride) {
+    a[i] += 1;
+  }
+}
+
+/**
+ * Kernel to increment 2 value to all the elements in array
+ */
+__global__ void addTwoKernel(int *a, int size) {
+  int offset = blockDim.x * blockIdx.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  for (int i = offset; i < size; i+= stride) {
+    a[i] += 2;
+  }
+}
+
+/**
+ * A Simple Kernel useful for validating some kernel related scenarios
+ */
+__global__ void simpleKernel() {
+}
+
+/**
+ * In addTenKernel, integer a value increased by 10
+ */
+__global__ void addTenKernel(int *a) {
+  *a = *a + 10;
+}
+
+/**
+ * Host function adds 10 to the integer a
+ */
+void addTen(void* data) {
+  int *a = reinterpret_cast<int *>(data);
+  *a = *a + 10;
+}
+
+/**
+ * Host function adds 20 to the integer a
+ */
+void addTwenty(void* data) {
+  int *a = reinterpret_cast<int *>(data);
+  *a = *a + 20;
+}
+
+/**
+ * Local function used in call back scenarios
+ */
+void callBackFunction(hipStream_t stream, hipError_t status, void* userData) {
+  REQUIRE(stream != nullptr);
+  REQUIRE(status == hipSuccess);
+  int *a = reinterpret_cast<int *>(userData);
+  *a += 100;
+}
+
+#endif  // CATCH_UNIT_DEVICE_HIPGETPROCADDRESSHELPERS_HH_

catch/unit/device/hipGetProcAddressIpcApis.cc

@@ -0,0 +1,252 @@
+/*
+Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_helper.hh>
+#include <hip_test_defgroups.hh>
+#include <utils.hh>
+#include <unistd.h>
+#include <sys/types.h>
+#include <sys/wait.h>
+#include "hipGetProcAddressHelpers.hh"
+
+/**
+ * Test Description
+ * ------------------------
+ *  - This test will get the function pointer of different
+ *  - memory IPC related APIs from the hipGetProcAddress API
+ *  - and then validates the basic functionality of that particular API
+ *  - using the funtion pointer.
+ * Test source
+ * ------------------------
+ *  - unit/device/hipGetProcAddress_IPC_APIs.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipGetProcAddress_IPC_Memory") {
+  int N = 40;
+  int Nbytes = N * sizeof(int);
+
+  int fd[2];
+  REQUIRE(pipe(fd) == 0);
+
+  auto pid = fork();
+
+  // Validating hipIpcGetMemHandle API
+  if ( pid != 0 ) {  // parent process
+    void* hipIpcGetMemHandle_ptr = nullptr;
+
+    int currentHipVersion = 0;
+    HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+    HIP_CHECK(hipGetProcAddress(
+              "hipIpcGetMemHandle",
+              &hipIpcGetMemHandle_ptr,
+              currentHipVersion, 0, nullptr));
+
+    hipError_t (*dyn_hipIpcGetMemHandle_ptr)(hipIpcMemHandle_t *, void *) =
+      reinterpret_cast<hipError_t (*)(hipIpcMemHandle_t *, void *)>
+      (hipIpcGetMemHandle_ptr);
+
+    int* srcHostMem = reinterpret_cast<int *>(malloc(Nbytes));
+    REQUIRE(srcHostMem != nullptr);
+    fillHostArray(srcHostMem, N, 10);
+
+    int* devMemSrc = nullptr;
+    HIP_CHECK(hipMalloc(&devMemSrc, Nbytes));
+    REQUIRE(devMemSrc != nullptr);
+    HIP_CHECK(hipMemcpy(devMemSrc, srcHostMem, Nbytes, hipMemcpyHostToDevice));
+
+    hipIpcMemHandle_t handle;
+    HIP_CHECK(dyn_hipIpcGetMemHandle_ptr(&handle, devMemSrc));
+
+    REQUIRE(write(fd[1], &handle, sizeof(handle)) >= 0);
+    REQUIRE(close(fd[1]) == 0);
+
+    REQUIRE(wait(NULL) >= 0);
+
+    HIP_CHECK(hipFree(devMemSrc));
+    free(srcHostMem);
+  } else {  // child process
+    // Validating hipIpcOpenMemHandle, hipIpcCloseMemHandle API's
+    void* hipIpcOpenMemHandle_ptr = nullptr;
+    void* hipIpcCloseMemHandle_ptr = nullptr;
+
+    int currentHipVersion = 0;
+    HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+    HIP_CHECK(hipGetProcAddress(
+              "hipIpcOpenMemHandle",
+              &hipIpcOpenMemHandle_ptr,
+              currentHipVersion, 0, nullptr));
+    HIP_CHECK(hipGetProcAddress(
+              "hipIpcCloseMemHandle",
+              &hipIpcCloseMemHandle_ptr,
+              currentHipVersion, 0, nullptr));
+
+    hipError_t (*dyn_hipIpcOpenMemHandle_ptr)(void **, hipIpcMemHandle_t,
+                                              unsigned int) =
+      reinterpret_cast<hipError_t (*)(void **, hipIpcMemHandle_t,
+                                      unsigned int)>
+                                     (hipIpcOpenMemHandle_ptr);
+    hipError_t (*dyn_hipIpcCloseMemHandle_ptr)(void *) =
+      reinterpret_cast<hipError_t (*)(void *)>
+      (hipIpcCloseMemHandle_ptr);
+
+    hipIpcMemHandle_t handle;
+    REQUIRE(read(fd[0], &handle, sizeof(handle)) >= 0);
+    REQUIRE(close(fd[0]) == 0);
+
+    int *devPtr = nullptr;
+    HIP_CHECK(dyn_hipIpcOpenMemHandle_ptr(reinterpret_cast<void **>(&devPtr),
+              handle, hipIpcMemLazyEnablePeerAccess));
+    REQUIRE(devPtr != nullptr);
+
+    addOneKernel<<< 1, 1 >>>(devPtr , N);
+
+    int* dstHostMem = reinterpret_cast<int *>(malloc(Nbytes));
+    REQUIRE(dstHostMem != nullptr);
+
+    HIP_CHECK(hipMemcpy(dstHostMem, devPtr, Nbytes, hipMemcpyDeviceToHost));
+    REQUIRE(validateHostArray(dstHostMem, N, 11) == true);
+
+    HIP_CHECK(dyn_hipIpcCloseMemHandle_ptr(devPtr));
+
+    free(dstHostMem);
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - This test will get the function pointer of different
+ *  - Event IPC related APIs from the hipGetProcAddress API
+ *  - and then validates the basic functionality of that particular API
+ *  - using the funtion pointer.
+ * Test source
+ * ------------------------
+ *  - unit/device/hipGetProcAddress_IPC_APIs.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipGetProcAddress_IPC_Event") {
+  int fd[2];
+  REQUIRE(pipe(fd) == 0);
+
+  auto pid = fork();
+
+  // Validating hipIpcGetEventHandle API
+  if ( pid != 0 ) {  // parent process
+    void* hipIpcGetEventHandle_ptr = nullptr;
+
+    int currentHipVersion = 0;
+    HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+    HIP_CHECK(hipGetProcAddress(
+              "hipIpcGetEventHandle",
+              &hipIpcGetEventHandle_ptr,
+              currentHipVersion, 0, nullptr));
+
+    hipError_t (*dyn_hipIpcGetEventHandle_ptr)(hipIpcEventHandle_t *,
+                                               hipEvent_t) =
+      reinterpret_cast<hipError_t (*)(hipIpcEventHandle_t *, hipEvent_t)>
+                                     (hipIpcGetEventHandle_ptr);
+
+    hipEvent_t start = nullptr;
+    HIP_CHECK(hipEventCreateWithFlags(&start, hipEventInterprocess |
+                                              hipEventDisableTiming));
+    REQUIRE(start != nullptr);
+
+    hipIpcEventHandle_t handle;
+    HIP_CHECK(dyn_hipIpcGetEventHandle_ptr(&handle, start));
+
+    REQUIRE(write(fd[1], &handle, sizeof(hipIpcEventHandle_t)) >= 0);
+    REQUIRE(close(fd[1]) == 0);
+
+    REQUIRE(wait(NULL) >= 0);
+
+    HIP_CHECK(hipEventDestroy(start));
+  } else {  // child process
+    // Validating hipIpcOpenMemHandle API
+    void* hipIpcOpenEventHandle_ptr = nullptr;
+
+    int currentHipVersion = 0;
+    HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+    HIP_CHECK(hipGetProcAddress(
+              "hipIpcOpenEventHandle",
+              &hipIpcOpenEventHandle_ptr,
+              currentHipVersion, 0, nullptr));
+
+    hipError_t (*dyn_hipIpcOpenEventHandle_ptr)(hipEvent_t *,
+                                                hipIpcEventHandle_t) =
+      reinterpret_cast<hipError_t (*)(hipEvent_t *, hipIpcEventHandle_t)>
+                                     (hipIpcOpenEventHandle_ptr);
+
+    hipIpcEventHandle_t handle;
+    REQUIRE(read(fd[0], &handle, sizeof(handle)) >= 0);
+    REQUIRE(close(fd[0]) == 0);
+
+    hipEvent_t start = nullptr;
+    HIP_CHECK(dyn_hipIpcOpenEventHandle_ptr(&start, handle));
+    REQUIRE(start != nullptr);
+
+    hipEvent_t stop = nullptr;
+    HIP_CHECK(hipEventCreate(&stop));
+    REQUIRE(stop != nullptr);
+
+    int N = 40;
+    int Nbytes = N * sizeof(int);
+
+    int* hostMem = reinterpret_cast<int *>(malloc(Nbytes));
+    REQUIRE(hostMem != nullptr);
+    fillHostArray(hostMem, N, 10);
+
+    int* devMem = nullptr;
+    HIP_CHECK(hipMalloc(&devMem, Nbytes));
+    REQUIRE(devMem != nullptr);
+
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreate(&stream));
+
+    HIP_CHECK(hipEventRecord(start, stream));
+
+    HIP_CHECK(hipMemcpyAsync(devMem, hostMem, Nbytes,
+                             hipMemcpyHostToDevice, stream));
+    addOneKernel<<< 1, 1 >>>(devMem , N);
+    HIP_CHECK(hipMemcpyAsync(hostMem, devMem, Nbytes,
+                             hipMemcpyDeviceToHost, stream));
+
+    HIP_CHECK(hipEventRecord(stop, stream));
+    HIP_CHECK(hipEventSynchronize(stop));
+
+    REQUIRE(validateHostArray(hostMem, N, 11) == true);
+
+    float time = 0.0f;
+    HIP_CHECK(hipEventElapsedTime(&time, start, stop));
+    REQUIRE(time > 0.0f);
+
+    HIP_CHECK(hipEventDestroy(stop));
+    HIP_CHECK(hipStreamDestroy(stream));
+    free(hostMem);
+    HIP_CHECK(hipFree(devMem));
+  }
+}