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

Change-Id: I4619f8e69b75d1193e98adce2707bf7876b1ba59


[ROCm/hip-tests commit: 7fb9301d6b]

projects/hip-tests/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()
 

projects/hip-tests/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",

projects/hip-tests/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_

projects/hip-tests/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));
+  }
+}

projects/hip-tests/catch/unit/executionControl/CMakeLists.txt

@@ -11,6 +11,7 @@ if(HIP_PLATFORM MATCHES "amd")
         hipExtLaunchKernel.cc
         hipExtLaunchMultiKernelMultiDevice.cc
         launch_api.cc
+        hipGetProcAddressLaunchCbExecCtrlApis.cc
     )
 else()
     # These functions are currently unimplemented on AMD

projects/hip-tests/catch/unit/executionControl/hipGetProcAddressLaunchCbExecCtrlApis.cc

@@ -0,0 +1,546 @@
+/*
+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 <utils.hh>
+#include "../device/hipGetProcAddressHelpers.hh"
+
+/**
+ * Test Description
+ * ------------------------
+ *  - This test will get the function pointer of different
+ *  - kernel launch APIs from the hipGetProcAddress API
+ *  - and then validates the basic functionality of that particular API
+ *  - using the funtion pointer.
+ * Test source
+ * ------------------------
+ *  - unit/executionControl/hipGetProcAddress_Launch_Callback_ExecCtrl_APIs.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipGetProcAddress_KernelLaunchApis") {
+  void* hipConfigureCall_ptr = nullptr;
+  void* hipSetupArgument_ptr = nullptr;
+  void* hipLaunchByPtr_ptr = nullptr;
+  void* __hipPushCallConfiguration_ptr = nullptr;
+  void* __hipPopCallConfiguration_ptr = nullptr;
+  void* hipExtLaunchKernel_ptr = nullptr;
+  void* hipDrvMemcpy2DUnaligned_ptr = nullptr;
+  void* hipLaunchKernel_ptr = nullptr;
+  void* hipLaunchHostFunc_ptr = nullptr;
+
+  int currentHipVersion = 0;
+  HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+  HIP_CHECK(hipGetProcAddress(
+            "hipConfigureCall",
+            &hipConfigureCall_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipSetupArgument",
+            &hipSetupArgument_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipLaunchByPtr",
+            &hipLaunchByPtr_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "__hipPushCallConfiguration",
+            &__hipPushCallConfiguration_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "__hipPopCallConfiguration",
+            &__hipPopCallConfiguration_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipExtLaunchKernel",
+            &hipExtLaunchKernel_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipDrvMemcpy2DUnaligned",
+            &hipDrvMemcpy2DUnaligned_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipLaunchKernel",
+            &hipLaunchKernel_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipLaunchHostFunc",
+            &hipLaunchHostFunc_ptr,
+            currentHipVersion, 0, nullptr));
+
+  hipError_t (*dyn_hipConfigureCall_ptr)(dim3, dim3, size_t, hipStream_t) =
+    reinterpret_cast<hipError_t (*)(dim3, dim3, size_t, hipStream_t)>
+                                   (hipConfigureCall_ptr);
+  hipError_t (*dyn_hipSetupArgument_ptr)(const void *, size_t, size_t) =
+    reinterpret_cast<hipError_t (*)(const void *, size_t, size_t)>
+                                   (hipSetupArgument_ptr);
+  hipError_t (*dyn_hipLaunchByPtr_ptr)(const void *) =
+    reinterpret_cast<hipError_t (*)(const void *)>
+                                   (hipLaunchByPtr_ptr);
+  hipError_t (*dyn__hipPushCallConfiguration_ptr)(dim3, dim3,
+                                                  size_t, hipStream_t) =
+    reinterpret_cast<hipError_t (*)(dim3, dim3, size_t, hipStream_t)>
+                                   (__hipPushCallConfiguration_ptr);
+  hipError_t (*dyn__hipPopCallConfiguration_ptr)(dim3 *, dim3 *,
+                                                 size_t *, hipStream_t *) =
+    reinterpret_cast<hipError_t (*)(dim3 *, dim3 *, size_t *, hipStream_t *)>
+                                   (__hipPopCallConfiguration_ptr);
+  hipError_t (*dyn_hipExtLaunchKernel_ptr)(const void *, dim3, dim3, void **,
+                                           size_t, hipStream_t, hipEvent_t,
+                                           hipEvent_t, int) =
+    reinterpret_cast<hipError_t (*)(const void *, dim3, dim3, void **, size_t,
+                                    hipStream_t, hipEvent_t, hipEvent_t, int)>
+                                   (hipExtLaunchKernel_ptr);
+  hipError_t (*dyn_hipDrvMemcpy2DUnaligned_ptr)(const hip_Memcpy2D *) =
+    reinterpret_cast<hipError_t (*)(const hip_Memcpy2D *)>
+                                   (hipDrvMemcpy2DUnaligned_ptr);
+
+  hipError_t (*dyn_hipLaunchKernel_ptr)(const void *, dim3, dim3, void **,
+                                        size_t, hipStream_t) =
+    reinterpret_cast<hipError_t (*)(const void *, dim3, dim3, void **,
+                                    size_t, hipStream_t)>
+                                   (hipLaunchKernel_ptr);
+
+  hipError_t (*dyn_hipLaunchHostFunc_ptr)(hipStream_t, hipHostFn_t, void *) =
+    reinterpret_cast<hipError_t (*)(hipStream_t, hipHostFn_t, void *)>
+    (hipLaunchHostFunc_ptr);
+
+  // Validating hipConfigureCall, hipSetupArgument, hipLaunchByPtr API's
+  {
+    // hipConfigureCall
+    dim3 blocksPerGrid(1, 1, 1);
+    dim3 threadsPerBlock(1, 1, 1);
+
+    HIP_CHECK(dyn_hipConfigureCall_ptr(blocksPerGrid, threadsPerBlock,
+                                       0, nullptr));
+
+    // hipSetupArgument
+    int hostInt = 10;
+    int *devInt = nullptr;
+    HIP_CHECK(hipMalloc(&devInt, sizeof(int)));
+    REQUIRE(devInt != nullptr);
+    HIP_CHECK(hipMemcpy(devInt, &hostInt, sizeof(int), hipMemcpyHostToDevice));
+
+    HIP_CHECK(dyn_hipSetupArgument_ptr(&devInt, sizeof(int*), 0));
+
+    // hipLaunchByPtr
+    HIP_CHECK(dyn_hipLaunchByPtr_ptr(reinterpret_cast<void*>(addTenKernel)));
+
+    HIP_CHECK(hipMemcpy(&hostInt, devInt, sizeof(int), hipMemcpyDeviceToHost));
+    REQUIRE(hostInt == 20);
+
+    HIP_CHECK(hipFree(devInt));
+  }
+
+  // Validating __hipPushCallConfiguration, __hipPopCallConfiguration APIs
+  {
+    // __hipPushCallConfiguration
+    dim3 blocksPerGrid(4, 6, 8);
+    dim3 threadsPerBlock(20, 30, 40);
+    HIP_CHECK(dyn__hipPushCallConfiguration_ptr(blocksPerGrid, threadsPerBlock,
+                                                512, nullptr));
+
+    // __hipPopCallConfiguration
+    dim3 gridDim, blockDim;
+    size_t sharedMem = -1;
+    hipStream_t stream;
+
+    HIP_CHECK(dyn__hipPopCallConfiguration_ptr(&gridDim, &blockDim,
+                                               &sharedMem, &stream));
+
+    REQUIRE(gridDim.x == blocksPerGrid.x);
+    REQUIRE(gridDim.y == blocksPerGrid.y);
+    REQUIRE(gridDim.z == blocksPerGrid.z);
+    REQUIRE(blockDim.x == threadsPerBlock.x);
+    REQUIRE(blockDim.y == threadsPerBlock.y);
+    REQUIRE(blockDim.z == threadsPerBlock.z);
+    REQUIRE(sharedMem == 512);
+    REQUIRE(stream == nullptr);
+  }
+
+  // Validating hipExtLaunchKernel API
+  {
+    dim3 blocksPerGrid(1, 1, 1);
+    dim3 threadsPerBlock(1, 1, 1);
+
+    int hostInt = 10;
+    int *devInt = nullptr;
+    HIP_CHECK(hipMalloc(&devInt, sizeof(int)));
+    REQUIRE(devInt != nullptr);
+    HIP_CHECK(hipMemcpy(devInt, &hostInt, sizeof(int), hipMemcpyHostToDevice));
+
+    void *kernel_args[] = {&devInt};
+    HIP_CHECK(dyn_hipExtLaunchKernel_ptr(reinterpret_cast<void*>(addTenKernel),
+              blocksPerGrid, threadsPerBlock, kernel_args, 0,
+              nullptr, nullptr, nullptr, 0));
+
+    HIP_CHECK(hipMemcpy(&hostInt, devInt, sizeof(int), hipMemcpyDeviceToHost));
+    REQUIRE(hostInt == 20);
+
+    HIP_CHECK(hipFree(devInt));
+  }
+
+  // Validating hipDrvMemcpy2DUnaligned API
+  {
+    size_t width = 256;
+    size_t height = 256;
+    const int N = width * height;
+    int value = 10;
+    size_t pitch;
+
+    // Host to Host
+    {
+      char* sHostMem = reinterpret_cast<char *>(malloc( N * sizeof(char)));
+      REQUIRE(sHostMem != nullptr);
+      fillCharHostArray(sHostMem, N, value);
+
+      char* dHostMem = reinterpret_cast<char *>(malloc( N * sizeof(char)));
+      REQUIRE(dHostMem != nullptr);
+
+      hip_Memcpy2D desc = {};
+      desc.srcMemoryType = hipMemoryTypeHost;
+      desc.srcHost = sHostMem;
+      desc.srcPitch = width;
+
+      desc.dstMemoryType = hipMemoryTypeHost;
+      desc.dstHost = dHostMem;
+      desc.dstPitch = width;
+      desc.WidthInBytes = width * sizeof(char);
+      desc.Height = height;
+
+      HIP_CHECK(dyn_hipDrvMemcpy2DUnaligned_ptr(&desc));
+
+      REQUIRE(validateCharHostArray(dHostMem, N, value) == true);
+
+      free(sHostMem);
+      free(dHostMem);
+    }
+
+    // Host to Device
+    {
+      char* hostMem = reinterpret_cast<char *>(malloc( N * sizeof(char)));
+      REQUIRE(hostMem != nullptr);
+      fillCharHostArray(hostMem, N, value);
+
+      char *devMem = nullptr;
+      HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&devMem),
+                               &pitch, width, height));
+      REQUIRE(devMem != nullptr);
+
+      hip_Memcpy2D desc = {};
+      desc.srcMemoryType = hipMemoryTypeHost;
+      desc.srcHost = hostMem;
+      desc.srcPitch = width;
+      desc.dstMemoryType = hipMemoryTypeDevice;
+      desc.dstDevice = hipDeviceptr_t(devMem);
+      desc.dstPitch = pitch;
+      desc.WidthInBytes = width*sizeof(char);
+      desc.Height = height;
+
+      HIP_CHECK(dyn_hipDrvMemcpy2DUnaligned_ptr(&desc));
+
+      REQUIRE(validateCharDeviceArray(devMem, N, value) == true);
+
+      free(hostMem);
+      HIP_CHECK(hipFree(devMem));
+    }
+
+    // Device To Host
+    {
+      char *devMem = nullptr;
+      HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&devMem),
+                               &pitch, width, height));
+      REQUIRE(devMem != nullptr);
+      HIP_CHECK(hipMemset2D(devMem, pitch, value, width, height));
+
+      char* hostMem = reinterpret_cast<char *>(malloc( N * sizeof(char)));
+      REQUIRE(hostMem != nullptr);
+
+      hip_Memcpy2D desc = {};
+      desc.srcMemoryType = hipMemoryTypeDevice;
+      desc.srcDevice = (hipDeviceptr_t)devMem;
+      desc.srcPitch = width;
+
+      desc.dstMemoryType = hipMemoryTypeHost;
+      desc.dstHost = hostMem;
+      desc.dstPitch = width;
+      desc.WidthInBytes = width * sizeof(char);
+      desc.Height = height;
+
+      HIP_CHECK(dyn_hipDrvMemcpy2DUnaligned_ptr(&desc));
+
+      REQUIRE(validateCharHostArray(hostMem, N, value) == true);
+
+      HIP_CHECK(hipFree(devMem));
+      free(hostMem);
+    }
+
+    // Device To Device
+    {
+      char *sDevMem = nullptr;
+      size_t sPitch;
+      HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&sDevMem),
+                               &sPitch, width, height));
+      REQUIRE(sDevMem != nullptr);
+      HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height));
+
+      char *dDevMem = nullptr;
+      size_t dPitch;
+      HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&dDevMem),
+                               &dPitch, width, height));
+      REQUIRE(dDevMem != nullptr);
+
+      hip_Memcpy2D desc = {};
+      desc.srcMemoryType = hipMemoryTypeDevice;
+      desc.srcDevice = hipDeviceptr_t(sDevMem);
+      desc.srcPitch = sPitch;
+      desc.dstMemoryType = hipMemoryTypeDevice;
+      desc.dstDevice = hipDeviceptr_t(dDevMem);
+      desc.dstPitch = dPitch;
+      desc.WidthInBytes = width*sizeof(char);
+      desc.Height = height;
+
+      HIP_CHECK(dyn_hipDrvMemcpy2DUnaligned_ptr(&desc));
+
+      REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true);
+
+      HIP_CHECK(hipFree(sDevMem));
+      HIP_CHECK(hipFree(dDevMem));
+    }
+  }
+
+  // Validating hipLaunchKernel API
+  {
+    const int N = 10;
+    const int Nbytes = 10 * sizeof(int);
+
+    int *hostArr = reinterpret_cast<int *>(malloc(Nbytes));
+    REQUIRE(hostArr != nullptr);
+    fillHostArray(hostArr, N, 10);
+
+    int *devArr = nullptr;
+    HIP_CHECK(hipMalloc(&devArr, Nbytes));
+    REQUIRE(devArr != nullptr);
+    HIP_CHECK(hipMemcpy(devArr, hostArr, Nbytes, hipMemcpyHostToDevice));
+
+    dim3 blocksPerGrid(1, 1, 1);
+    dim3 threadsPerBlock(1, 1, N);
+
+    struct kernelArgs {
+      void *arr;
+      int size;
+    };
+    kernelArgs kernelArg;
+    kernelArg.arr = devArr;
+    kernelArg.size = N;
+    void *kernel_args[] = {&kernelArg.arr, &kernelArg.size};
+
+    HIP_CHECK(dyn_hipLaunchKernel_ptr(reinterpret_cast<void*>(addOneKernel),
+              blocksPerGrid, threadsPerBlock, kernel_args, 0, nullptr));
+
+    HIP_CHECK(hipMemcpy(hostArr, devArr, Nbytes, hipMemcpyDeviceToHost));
+    REQUIRE(validateHostArray(hostArr, N, 11) == true);
+
+    free(hostArr);
+    HIP_CHECK(hipFree(devArr));
+  }
+
+  // Validating hipLaunchHostFunc API
+  {
+     int data = 30;
+     hipHostFn_t fn = addTen;
+
+     hipStream_t stream;
+     HIP_CHECK(hipStreamCreate(&stream));
+
+     HIP_CHECK(dyn_hipLaunchHostFunc_ptr(stream, fn,
+               reinterpret_cast<void *>(&data)));
+     HIP_CHECK(hipStreamSynchronize(stream));
+
+     REQUIRE(data == 40);
+     HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - This test will get the function pointer of different
+ *  - call back activity APIs from the hipGetProcAddress API
+ *  - and then validates the basic functionality of that particular API
+ *  - using the funtion pointer.
+ * Test source
+ * ------------------------
+ *  - unit/executionControl/hipGetProcAddress_Launch_Callback_ExecCtrl_APIs.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipGetProcAddress_CallbackActivityAPIs") {
+  void* hipGetStreamDeviceId_ptr = nullptr;
+  void* hipApiName_ptr = nullptr;
+  void* hipKernelNameRef_ptr = nullptr;
+
+  int currentHipVersion = 0;
+  HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+  HIP_CHECK(hipGetProcAddress(
+            "hipGetStreamDeviceId",
+            &hipGetStreamDeviceId_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipApiName",
+            &hipApiName_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipKernelNameRef",
+            &hipKernelNameRef_ptr,
+            currentHipVersion, 0, nullptr));
+
+  int (*dyn_hipGetStreamDeviceId_ptr)(hipStream_t) =
+    reinterpret_cast<int (*)(hipStream_t)>
+    (hipGetStreamDeviceId_ptr);
+  const char * (*dyn_hipApiName_ptr)(uint32_t) =
+    reinterpret_cast<const char * (*)(uint32_t)>
+    (hipApiName_ptr);
+  const char * (*dyn_hipKernelNameRef_ptr)(const hipFunction_t) =
+    reinterpret_cast<const char * (*)(const hipFunction_t)>
+    (hipKernelNameRef_ptr);
+
+  // Validating hipGetStreamDeviceId API
+  {
+    int deviceId = 0;
+    HIP_CHECK(hipSetDevice(deviceId));
+
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreate(&stream));
+
+    REQUIRE(hipGetStreamDeviceId(stream) ==
+            dyn_hipGetStreamDeviceId_ptr(stream));
+
+    REQUIRE(dyn_hipGetStreamDeviceId_ptr(stream) == deviceId);
+
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+
+  // Validating hipApiName API
+  {
+    uint32_t id = 1;
+    REQUIRE(strcmp(hipApiName(id), dyn_hipApiName_ptr(id)) == 0);
+  }
+
+  // Validating hipKernelNameRef API
+  {
+    hipFunction_t funcPointer;
+    REQUIRE(hipGetFuncBySymbol(&funcPointer,
+            reinterpret_cast<const void*>(addTenKernel)) == hipSuccess);
+
+    REQUIRE(strcmp(hipKernelNameRef(funcPointer),
+                   dyn_hipKernelNameRef_ptr(funcPointer)) == 0);
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - This test will get the function pointer of different
+ *  - Execution Control APIs from the hipGetProcAddress API
+ *  - and then validates the basic functionality of that particular API
+ *  - using the funtion pointer.
+ * Test source
+ * ------------------------
+ *  - unit/executionControl/hipGetProcAddress_Launch_Callback_ExecCtrl_APIs.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipGetProcAddress_ExecutionControlAPIs") {
+  void* hipFuncSetAttribute_ptr = nullptr;
+  void* hipFuncSetCacheConfig_ptr = nullptr;
+  void* hipFuncSetSharedMemConfig_ptr = nullptr;
+
+  int currentHipVersion = 0;
+  HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+  HIP_CHECK(hipGetProcAddress(
+            "hipFuncSetAttribute",
+            &hipFuncSetAttribute_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipFuncSetCacheConfig",
+            &hipFuncSetCacheConfig_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipFuncSetSharedMemConfig",
+            &hipFuncSetSharedMemConfig_ptr,
+            currentHipVersion, 0, nullptr));
+
+  hipError_t (*dyn_hipFuncSetAttribute_ptr)(const void *,
+                                            hipFuncAttribute, int) =
+    reinterpret_cast<hipError_t (*)(const void *,
+                                    hipFuncAttribute, int)>
+                                   (hipFuncSetAttribute_ptr);
+
+  hipError_t (*dyn_hipFuncSetCacheConfig_ptr)(const void *,
+                                              hipFuncCache_t) =
+    reinterpret_cast<hipError_t (*)(const void *,
+                                    hipFuncCache_t)>
+                                   (hipFuncSetCacheConfig_ptr);
+
+  hipError_t (*dyn_hipFuncSetSharedMemConfig_ptr)(const void *,
+                                                  hipSharedMemConfig) =
+    reinterpret_cast<hipError_t (*)(const void *, hipSharedMemConfig)>
+                                   (hipFuncSetSharedMemConfig_ptr);
+
+  // Validating hipFuncSetAttribute API
+  hipFuncAttribute funcAttributes[] =
+                   { hipFuncAttributeMaxDynamicSharedMemorySize,
+                     hipFuncAttributePreferredSharedMemoryCarveout,
+                     hipFuncAttributeMax};
+  int value = 90;
+  for ( auto funcAttribute : funcAttributes ) {
+    HIP_CHECK(dyn_hipFuncSetAttribute_ptr(
+              reinterpret_cast<void*>(addOneKernel),
+              funcAttribute, value));
+  }
+
+  // Validating hipFuncSetCacheConfig API
+  hipFuncCache_t funcCacheConfig[] = { hipFuncCachePreferNone,
+                                       hipFuncCachePreferShared,
+                                       hipFuncCachePreferL1,
+                                       hipFuncCachePreferEqual
+                                     };
+  for ( auto config : funcCacheConfig ) {
+    HIP_CHECK(dyn_hipFuncSetCacheConfig_ptr(
+              reinterpret_cast<void*>(addOneKernel), config));
+  }
+
+  // Validating hipFuncSetSharedMemConfig API
+  hipSharedMemConfig sharedMemConfig[] = { hipSharedMemBankSizeDefault,
+                                           hipSharedMemBankSizeFourByte,
+                                           hipSharedMemBankSizeEightByte
+                                         };
+  for ( auto config : sharedMemConfig ) {
+    HIP_CHECK(dyn_hipFuncSetSharedMemConfig_ptr(
+              reinterpret_cast<void*>(addOneKernel), config));
+  }
+}

projects/hip-tests/catch/unit/graph/CMakeLists.txt

@@ -88,6 +88,7 @@ if(HIP_PLATFORM MATCHES "amd")
   set(AMD_SRC
     hipStreamCaptureExtModuleLaunchKernel.cc
     hipStreamBeginCaptureToGraph.cc
+    hipGetProcAddressGraphApis.cc
   )
   set(TEST_SRC ${TEST_SRC} ${AMD_SRC})
 endif()

projects/hip-tests/catch/unit/memory/CMakeLists.txt

@@ -106,7 +106,7 @@ if(HIP_PLATFORM MATCHES "amd")
       hipMemVmm.cc
       hipArray.cc
       hipMemcpyDeviceToDeviceNoCU.cc
-      hipGetProcAddress_Memory_APIs.cc)
+      hipGetProcAddressMemoryApis.cc)
   if(UNIX)
     # Should be compiled for NVIDIA as well after EXSWHTEC-346 is addressed
     # For windows build error occurs undefined symbol: hipPointerSetAttribute

projects/hip-tests/catch/unit/memory/hipGetProcAddressMemoryApis.cc

@@ -21,159 +21,7 @@ THE SOFTWARE.
 #include <hip_test_helper.hh>
 #include <utils.hh>
 #include "hipMallocManagedCommon.hh"
-
-/**
- * @addtogroup hipGetProcAddress hipGetProcAddress
- * @{
- * @ingroup DeviceTest
- * `hipGetProcAddress(const char* symbol, void** pfn,
- *                    int  hipVersion, uint64_t flags,
- *                    hipDriverProcAddressQueryResult* symbolStatus);` -
- * Gets the symbol's function address.
- */
-
-/**
- * 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 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);
-}
-
-/**
- * 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;
-  }
-}
-
-/**
- * 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;
-}
+#include "../device/hipGetProcAddressHelpers.hh"
 
 /**
  * Test Description

projects/hip-tests/catch/unit/module/CMakeLists.txt

@@ -77,7 +77,7 @@ set(TEST_SRC
     ${TEST_SRC}
     hipExtModuleLaunchKernel.cc
     hipHccModuleLaunchKernel.cc
-    hipGetProcAddress_Module_APIs.cc)
+    hipGetProcAddressModuleApis.cc)
 
 add_custom_target(empty_module.code
                   COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}

projects/hip-tests/catch/unit/module/hipGetProcAddressModuleApis.cc

@@ -22,54 +22,7 @@ THE SOFTWARE.
 #include <utils.hh>
 #include "hip/hip_ext.h"
 #include "hip_module_common.hh"
-
-/**
- * @addtogroup hipGetProcAddress hipGetProcAddress
- * @{
- * @ingroup DeviceTest
- * `hipGetProcAddress(const char* symbol, void** pfn,
-                      int  hipVersion, uint64_t flags,
-                      hipDriverProcAddressQueryResult* symbolStatus);` -
- * Gets the symbol's function address.
- */
-
-/**
- * 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;
-}
-
-/**
- * Kernel to add one for each element 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;
-  }
-}
-
-/**
- * A simple kernel to check some APIs functionality
- */
-__global__ void sampleKernel() {
-}
+#include "../device/hipGetProcAddressHelpers.hh"
 
 /**
  * Test Description
@@ -628,7 +581,7 @@ TEST_CASE("Unit_hipGetProcAddress_ModuleApisCooperativeKernels") {
     std::vector<hipLaunchParams> params(deviceCount);
 
     for (int i = 0; i < deviceCount; ++i) {
-      params[i].func = reinterpret_cast<void *>(sampleKernel);
+      params[i].func = reinterpret_cast<void *>(simpleKernel);
       params[i].gridDim = {1, 1, 1};
       params[i].blockDim = {1, 1, 1};
       params[i].args = nullptr;
@@ -663,7 +616,7 @@ TEST_CASE("Unit_hipGetProcAddress_ModuleApisCooperativeKernels") {
     std::vector<hipLaunchParams> params(deviceCount);
 
     for (int i = 0; i < deviceCount; ++i) {
-      params[i].func = reinterpret_cast<void *>(sampleKernel);
+      params[i].func = reinterpret_cast<void *>(simpleKernel);
       params[i].gridDim = {1, 1, 1};
       params[i].blockDim = {1, 1, 1};
       params[i].args = nullptr;

projects/hip-tests/catch/unit/streamperthread/CMakeLists.txt

@@ -8,6 +8,11 @@ set(TEST_SRC
   hipStreamPerThrdCompilerOptn.cc
 )
 
+if(HIP_PLATFORM MATCHES "amd")
+   set(TEST_SRC ${TEST_SRC}
+       hipGetProcAddressSptApis.cc)
+endif()
+
 if(HIP_PLATFORM MATCHES "amd")
         set_source_files_properties(hipStreamPerThrdCompilerOptn.cc PROPERTIES COMPILE_OPTIONS "-fgpu-default-stream=per-thread")
 endif()

projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt

@@ -48,9 +48,10 @@ if(HIP_PLATFORM MATCHES "amd")
       hipMemMap.cc
       hipMemRelease.cc
       hipMemUnmap.cc
-      hipMemVmm_old.cc)
+      hipMemVmm_old.cc
+      hipGetProcAddressVmmApis.cc)
 endif()
 
 hip_add_exe_to_target(NAME VirtualMemoryManagementTest
   TEST_SRC ${TEST_SRC}
-  TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
+  TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC})

projects/hip-tests/catch/unit/virtualMemoryManagement/hipGetProcAddressVmmApis.cc

@@ -0,0 +1,248 @@
+/*
+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 <utils.hh>
+#include "hip_vmm_common.hh"
+#include "../device/hipGetProcAddressHelpers.hh"
+
+/**
+ * Test Description
+ * ------------------------
+ *  - This test will get the function pointer of different
+ *  - Virtual Memory Management APIs from the hipGetProcAddress API
+ *  - and then validates the basic functionality of that particular API
+ *  - using the funtion pointer.
+ * Test source
+ * ------------------------
+ *  - unit/virtualMemoryManagement/hipGetProcAddress_VMMAPIs.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipGetProcAddress_VMM") {
+  int value = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&value,
+            hipDeviceAttributeVirtualMemoryManagementSupported, 0));
+  if (value == 0) {
+    HipTest::HIP_SKIP_TEST("Machine does not support VMM. Skipping Test..");
+    return;
+  }
+
+  void* hipMemGetAllocationGranularity_ptr = nullptr;
+  void* hipMemCreate_ptr = nullptr;
+  void* hipMemAddressReserve_ptr = nullptr;
+  void* hipMemMap_ptr = nullptr;
+  void* hipMemSetAccess_ptr = nullptr;
+  void* hipMemUnmap_ptr = nullptr;
+  void* hipMemAddressFree_ptr = nullptr;
+  void* hipMemRelease_ptr = nullptr;
+  void* hipMemGetAccess_ptr = nullptr;
+  void* hipMemGetAllocationPropertiesFromHandle_ptr = nullptr;
+  void* hipMemRetainAllocationHandle_ptr = nullptr;
+
+  int currentHipVersion = 0;
+  HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+  HIP_CHECK(hipGetProcAddress("hipMemGetAllocationGranularity",
+                              &hipMemGetAllocationGranularity_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipMemCreate",
+                              &hipMemCreate_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipMemAddressReserve",
+                              &hipMemAddressReserve_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipMemMap",
+                              &hipMemMap_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipMemSetAccess",
+                              &hipMemSetAccess_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipMemUnmap",
+                              &hipMemUnmap_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipMemAddressFree",
+                              &hipMemAddressFree_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipMemRelease",
+                              &hipMemRelease_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipMemGetAccess",
+                              &hipMemGetAccess_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipMemGetAllocationPropertiesFromHandle",
+                              &hipMemGetAllocationPropertiesFromHandle_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipMemRetainAllocationHandle",
+                              &hipMemRetainAllocationHandle_ptr,
+                              currentHipVersion, 0, nullptr));
+
+  hipError_t(*dyn_hipMemGetAllocationGranularity_ptr)(
+              size_t *, const hipMemAllocationProp *,
+              hipMemAllocationGranularity_flags) =
+    reinterpret_cast<hipError_t (*)(size_t *,
+              const hipMemAllocationProp *,
+              hipMemAllocationGranularity_flags)>
+             (hipMemGetAllocationGranularity_ptr);
+
+  hipError_t(*dyn_hipMemCreate_ptr)(
+             hipMemGenericAllocationHandle_t *,
+             size_t, const hipMemAllocationProp *,
+             uint64_t) =
+    reinterpret_cast<hipError_t (*)(hipMemGenericAllocationHandle_t *,
+                                    size_t, const hipMemAllocationProp *,
+                                    uint64_t)>
+                                    (hipMemCreate_ptr);
+
+  hipError_t(*dyn_hipMemAddressReserve_ptr)(
+             void **, size_t, size_t, void *, uint64_t) =
+    reinterpret_cast<hipError_t (*)(void **, size_t, size_t, void *,
+                                    uint64_t)>
+                                    (hipMemAddressReserve_ptr);
+
+  hipError_t(*dyn_hipMemMap_ptr)(
+             void *, size_t, size_t,
+             hipMemGenericAllocationHandle_t, uint64_t) =
+    reinterpret_cast<hipError_t (*)(void *, size_t, size_t,
+                                    hipMemGenericAllocationHandle_t,
+                                    uint64_t)>
+                                    (hipMemMap_ptr);
+
+  hipError_t(*dyn_hipMemSetAccess_ptr)(
+             void *, size_t, const hipMemAccessDesc *, size_t) =
+    reinterpret_cast<hipError_t (*)(void *, size_t,
+                                    const hipMemAccessDesc *, size_t)>
+                                    (hipMemSetAccess_ptr);
+
+  hipError_t(*dyn_hipMemUnmap_ptr)(
+             void *, size_t) =
+    reinterpret_cast<hipError_t (*)(void *, size_t)>
+    (hipMemUnmap_ptr);
+
+  hipError_t(*dyn_hipMemAddressFree_ptr)(
+             void *, size_t) =
+    reinterpret_cast<hipError_t (*)(void *, size_t)>
+                                   (hipMemAddressFree_ptr);
+
+  hipError_t(*dyn_hipMemRelease_ptr)(hipMemGenericAllocationHandle_t) =
+    reinterpret_cast<hipError_t (*)(hipMemGenericAllocationHandle_t)>
+                                   (hipMemRelease_ptr);
+
+  hipError_t(*dyn_hipMemGetAccess_ptr)(
+             uint64_t *, const hipMemLocation *, void *) =
+    reinterpret_cast<hipError_t (*)(uint64_t *,
+                                    const hipMemLocation *, void *)>
+                                    (hipMemGetAccess_ptr);
+
+  hipError_t(*dyn_hipMemGetAllocationPropertiesFromHandle_ptr)(
+             hipMemAllocationProp *, hipMemGenericAllocationHandle_t) =
+    reinterpret_cast<hipError_t (*)(hipMemAllocationProp *,
+      hipMemGenericAllocationHandle_t)>
+     (hipMemGetAllocationPropertiesFromHandle_ptr);
+
+  hipError_t(*dyn_hipMemRetainAllocationHandle_ptr)(
+             hipMemGenericAllocationHandle_t *, void *) =
+    reinterpret_cast<hipError_t (*)(hipMemGenericAllocationHandle_t *, void *)>
+     (hipMemRetainAllocationHandle_ptr);
+
+  const int N = 10;
+  const int Nbytes = 10 * sizeof(int);
+  int *hostMem = reinterpret_cast<int *>(malloc(Nbytes));
+  REQUIRE(hostMem != nullptr);
+  fillHostArray(hostMem, N, 10);
+
+  hipMemAllocationProp prop{};
+  prop.type = hipMemAllocationTypePinned;
+  prop.requestedHandleType = hipMemHandleTypeNone;
+  prop.location.type = hipMemLocationTypeDevice;
+  prop.location.id = 0;
+
+  // Validating hipMemGetAllocationGranularity API
+  size_t granularity = 0;
+  size_t granularityWithFuncPtr = 0;
+
+  HIP_CHECK(hipMemGetAllocationGranularity(&granularity,
+            &prop, hipMemAllocationGranularityMinimum));
+  HIP_CHECK(dyn_hipMemGetAllocationGranularity_ptr(&granularityWithFuncPtr,
+            &prop, hipMemAllocationGranularityMinimum));
+
+  REQUIRE(granularity > 0);
+  REQUIRE(granularityWithFuncPtr > 0);
+  REQUIRE(granularityWithFuncPtr == granularity);
+
+  // Validating hipMemCreate API
+  hipMemGenericAllocationHandle_t handle = nullptr;
+  HIP_CHECK(dyn_hipMemCreate_ptr(&handle, granularity, &prop, 0));
+  REQUIRE(handle != nullptr);
+
+  // Validating hipMemAddressReserve API
+  void* ptr = nullptr;
+  HIP_CHECK(dyn_hipMemAddressReserve_ptr(&ptr, granularity, 0, 0, 0));
+  REQUIRE(ptr != nullptr);
+
+  // Validating hipMemMap API
+  HIP_CHECK(dyn_hipMemMap_ptr(ptr, granularity, 0, handle, 0));
+  REQUIRE(ptr != nullptr);
+
+  // Validating hipMemSetAccess API
+  hipMemAccessDesc desc;
+  desc.location.type = hipMemLocationTypeDevice;
+  desc.location.id = 0;
+  desc.flags = hipMemAccessFlagsProtReadWrite;
+  HIP_CHECK(dyn_hipMemSetAccess_ptr(ptr, granularity, &desc, 1));
+  REQUIRE(ptr != nullptr);
+
+  // Performing some operations on ptr, to validate it
+  HIP_CHECK(hipMemcpy(ptr, hostMem, Nbytes, hipMemcpyHostToDevice));
+  addOneKernel<<< 1, 1 >>>(reinterpret_cast<int *>(ptr), N);
+  HIP_CHECK(hipMemcpy(hostMem, ptr, Nbytes, hipMemcpyDeviceToHost));
+  validateHostArray(hostMem, N, 11);
+
+  // Validating hipMemGetAccess API
+  uint64_t flags;
+  hipMemLocation location = {hipMemLocationTypeDevice, 0};
+  HIP_CHECK(dyn_hipMemGetAccess_ptr(&flags, &location, ptr));
+  REQUIRE(flags == hipMemAccessFlagsProtReadWrite);
+
+  // Validating hipMemGetAllocationPropertiesFromHandle API
+  hipMemAllocationProp requiredProp;
+  HIP_CHECK(dyn_hipMemGetAllocationPropertiesFromHandle_ptr(
+            &requiredProp, handle));
+  REQUIRE(requiredProp.type == hipMemAllocationTypePinned);
+  REQUIRE(requiredProp.requestedHandleType == hipMemHandleTypeNone);
+  REQUIRE(requiredProp.location.type == hipMemLocationTypeDevice);
+  REQUIRE(requiredProp.location.id == 0);
+
+  // Validating hipMemRetainAllocationHandle API
+  hipMemGenericAllocationHandle_t requiredHandle = nullptr;
+  HIP_CHECK(dyn_hipMemRetainAllocationHandle_ptr(&requiredHandle, ptr));
+  REQUIRE(requiredHandle == handle);
+
+  // Validating hipMemUnmap, hipMemAddressFree, hipMemRelease API's
+  HIP_CHECK(dyn_hipMemUnmap_ptr(ptr, granularity));
+  HIP_CHECK(dyn_hipMemAddressFree_ptr(ptr, granularity));
+  HIP_CHECK(dyn_hipMemRelease_ptr(handle));
+
+  // Performing operation on ptr, to check it is invalidated or not
+  REQUIRE(hipMemcpy(ptr, hostMem, Nbytes, hipMemcpyHostToDevice)
+          == hipErrorInvalidValue);
+  free(hostMem);
+}