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

Change-Id: I1da8b0d32f5ae49602b9ff062dc888cc6227b88f [ROCm/hip-tests commit: 6c35a8cb95]
2024-07-04 15:53:26 +05:30
@@ -89,7 +89,8 @@ if(HIP_PLATFORM MATCHES "amd")
      hipMallocMngdMultiThread.cc
      hipMemVmm.cc
      hipArray.cc
-      hipMemcpyDeviceToDeviceNoCU.cc)
+      hipMemcpyDeviceToDeviceNoCU.cc
+      hipGetProcAddress_Memory_APIs.cc)
  if(UNIX)
    # Should be compiled for NVIDIA as well after EXSWHTEC-346 is addressed
    # For windows build error occurs undefined symbol: hipPointerSetAttribute
@@ -70,7 +70,8 @@ if(HIP_PLATFORM MATCHES "amd")
 set(TEST_SRC
    ${TEST_SRC}
    hipExtModuleLaunchKernel.cc
-    hipHccModuleLaunchKernel.cc)
+    hipHccModuleLaunchKernel.cc
+    hipGetProcAddress_Module_APIs.cc)

 add_custom_target(empty_module.code
                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
@@ -92,6 +93,13 @@ add_custom_target(copyKernel.s
                  -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH}
                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH})

+add_custom_target(addKernel.code
+                  COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --genco ${OFFLOAD_ARCH_STR}
+                  ${CMAKE_CURRENT_SOURCE_DIR}/addKernel.cc
+                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/addKernel.code
+                  -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH}
+                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH})
+
 if(UNIX)
 set(TEST_SRC
    ${TEST_SRC}
@@ -174,6 +182,7 @@ add_dependencies(ModuleTest get_tex_ref_module)
 if(HIP_PLATFORM MATCHES "amd")
 add_dependencies(build_tests empty_module.code)
 add_dependencies(build_tests copyKernel.code copyKernel.s)
+add_dependencies(build_tests addKernel.code)
 if(UNIX)
 add_dependencies(build_tests copiousArgKernel.code copiousArgKernel0.code copiousArgKernel1.code copiousArgKernel2.code
 copiousArgKernel3.code copiousArgKernel16.code copiousArgKernel17.code)
@@ -0,0 +1,35 @@
+/*
+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/hip_runtime.h>
+
+__device__ int globalDevData = 10;
+
+extern "C" __global__ void addKernel(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;
+  }
+}
+
+texture<float, 2> tex;
+
+extern "C" __global__ void sampleModuleKernel() {
+}
@@ -0,0 +1,885 @@
+/*
+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/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() {
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - This test will get the function pointer of different module management
+ *  - (load/unload/GetAttribute/launch) related APIs from the hipGetProcAddress API
+ *  - and then validates the basic functionality of that particular API
+ *  - using the funtion pointer.
+ * Test source
+ * ------------------------
+ *  - unit/module/hipGetProcAddress_Module_APIs.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipGetProcAddress_ModuleApis") {
+  void* hipModuleLoad_ptr = nullptr;
+  void* hipModuleUnload_ptr = nullptr;
+  void* hipModuleGetFunction_ptr = nullptr;
+  void* hipModuleLaunchKernel_ptr = nullptr;
+  void* hipGetFuncBySymbol_ptr = nullptr;
+  void* hipFuncGetAttributes_ptr = nullptr;
+  void* hipFuncGetAttribute_ptr = nullptr;
+  void* hipModuleGetGlobal_ptr = nullptr;
+  void* hipExtModuleLaunchKernel_ptr = nullptr;
+  void* hipHccModuleLaunchKernel_ptr = nullptr;
+
+  int currentHipVersion = 0;
+  HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+  HIP_CHECK(hipGetProcAddress("hipModuleLoad",
+                              &hipModuleLoad_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipModuleUnload",
+                              &hipModuleUnload_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipModuleGetFunction",
+                              &hipModuleGetFunction_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipModuleLaunchKernel",
+                              &hipModuleLaunchKernel_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipGetFuncBySymbol",
+                              &hipGetFuncBySymbol_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipFuncGetAttributes",
+                              &hipFuncGetAttributes_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipFuncGetAttribute",
+                              &hipFuncGetAttribute_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipModuleGetGlobal",
+                              &hipModuleGetGlobal_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipExtModuleLaunchKernel",
+                              &hipExtModuleLaunchKernel_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipHccModuleLaunchKernel",
+                              &hipHccModuleLaunchKernel_ptr,
+                              currentHipVersion, 0, nullptr));
+
+  hipError_t (*dyn_hipModuleLoad_ptr)(hipModule_t *, const char *) =
+    reinterpret_cast<hipError_t (*)(hipModule_t *, const char *)>
+    (hipModuleLoad_ptr);
+  hipError_t (*dyn_hipModuleUnload_ptr)(hipModule_t) =
+    reinterpret_cast<hipError_t (*)(hipModule_t)>
+    (hipModuleUnload_ptr);
+  hipError_t (*dyn_hipModuleGetFunction_ptr)(
+    hipFunction_t *, hipModule_t, const char *) =
+    reinterpret_cast<hipError_t (*)(hipFunction_t *,
+    hipModule_t, const char *)>
+    (hipModuleGetFunction_ptr);
+  hipError_t (*dyn_hipModuleLaunchKernel_ptr)(
+    hipFunction_t,
+    unsigned int, unsigned int, unsigned int,
+    unsigned int, unsigned int, unsigned int,
+    unsigned int, hipStream_t,
+    void **, void **) =
+    reinterpret_cast<hipError_t (*)(hipFunction_t,
+    unsigned int, unsigned int, unsigned int,
+    unsigned int, unsigned int, unsigned int,
+    unsigned int, hipStream_t,
+    void **, void **) > (hipModuleLaunchKernel_ptr);
+  hipError_t (*dyn_hipGetFuncBySymbol_ptr)(hipFunction_t *, const void *) =
+    reinterpret_cast<hipError_t (*)(hipFunction_t *, const void *)>
+    (hipGetFuncBySymbol_ptr);
+  hipError_t (*dyn_hipFuncGetAttributes_ptr)(
+    struct hipFuncAttributes *, const void *) =
+    reinterpret_cast<hipError_t (*)(struct hipFuncAttributes *, const void *)>
+    (hipFuncGetAttributes_ptr);
+  hipError_t (*dyn_hipFuncGetAttribute_ptr)(
+    int *, hipFunction_attribute, hipFunction_t) =
+    reinterpret_cast<hipError_t (*)(int *, hipFunction_attribute,
+    hipFunction_t)>(hipFuncGetAttribute_ptr);
+  hipError_t (*dyn_hipModuleGetGlobal_ptr)(
+    hipDeviceptr_t *, size_t *, hipModule_t, const char *) =
+    reinterpret_cast<hipError_t (*)(hipDeviceptr_t *, size_t *,
+    hipModule_t, const char *)>
+    (hipModuleGetGlobal_ptr);
+
+  hipError_t (*dyn_hipExtModuleLaunchKernel_ptr)(hipFunction_t,
+                                      uint32_t, uint32_t, uint32_t,
+                                      uint32_t, uint32_t, uint32_t,
+                                      size_t, hipStream_t,
+                                      void **, void **,
+                                      hipEvent_t, hipEvent_t, uint32_t) =
+    reinterpret_cast<hipError_t (*)(hipFunction_t,
+                                      uint32_t, uint32_t, uint32_t,
+                                      uint32_t, uint32_t, uint32_t,
+                                      size_t, hipStream_t,
+                                      void **, void **,
+                                      hipEvent_t, hipEvent_t, uint32_t)>
+                                      (hipExtModuleLaunchKernel_ptr);
+
+  hipError_t (*dyn_hipHccModuleLaunchKernel_ptr)(hipFunction_t,
+                                      uint32_t, uint32_t, uint32_t,
+                                      uint32_t, uint32_t, uint32_t,
+                                      size_t, hipStream_t,
+                                      void **, void **,
+                                      hipEvent_t, hipEvent_t) =
+    reinterpret_cast<hipError_t (*)(hipFunction_t,
+                                      uint32_t, uint32_t, uint32_t,
+                                      uint32_t, uint32_t, uint32_t,
+                                      size_t, hipStream_t,
+                                      void **, void **,
+                                      hipEvent_t, hipEvent_t)>
+                                      (hipHccModuleLaunchKernel_ptr);
+
+  // Validating hipModuleLoad API
+  hipModule_t module;
+  HIP_CHECK(dyn_hipModuleLoad_ptr(&module, "addKernel.code"));
+  REQUIRE(module != nullptr);
+
+  // Validating hipModuleGetFunction API
+  hipFunction_t function;
+  HIP_CHECK(dyn_hipModuleGetFunction_ptr(&function, module, "addKernel"));
+  REQUIRE(function != nullptr);
+
+  // Validating  hipModuleLaunchKernel 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 kernelParameters {
+    void *arr;
+    int size;
+  };
+  kernelParameters kernelParam{};
+  kernelParam.arr = devArr;
+  kernelParam.size = N;
+
+  auto size = sizeof(kernelParam);
+  void* kernel_parameter[] = { HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam,
+                               HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
+                               HIP_LAUNCH_PARAM_END };
+
+  HIP_CHECK(dyn_hipModuleLaunchKernel_ptr(function,
+            blocksPerGrid.x, blocksPerGrid.y, blocksPerGrid.z,
+            threadsPerBlock.x, threadsPerBlock.y, threadsPerBlock.z,
+            0, 0, nullptr, kernel_parameter));
+
+  HIP_CHECK(hipMemcpy(hostArr, devArr, Nbytes, hipMemcpyDeviceToHost));
+  REQUIRE(validateHostArray(hostArr, N, 12) == true);
+
+  // Validating hipExtModuleLaunchKernel API
+  HIP_CHECK(dyn_hipExtModuleLaunchKernel_ptr(function,
+            blocksPerGrid.x, blocksPerGrid.y, blocksPerGrid.z,
+            threadsPerBlock.x, threadsPerBlock.y, threadsPerBlock.z,
+            0, 0,
+            nullptr, kernel_parameter,
+            nullptr, nullptr, 0));
+
+  HIP_CHECK(hipMemcpy(hostArr, devArr, Nbytes, hipMemcpyDeviceToHost));
+  REQUIRE(validateHostArray(hostArr, N, 14) == true);
+
+  // Validating hipHccModuleLaunchKernel API
+  HIP_CHECK(dyn_hipHccModuleLaunchKernel_ptr(function,
+            blocksPerGrid.x, blocksPerGrid.y, blocksPerGrid.z,
+            threadsPerBlock.x, threadsPerBlock.y, threadsPerBlock.z,
+            0, 0,
+            nullptr, kernel_parameter,
+            nullptr, nullptr));
+
+  HIP_CHECK(hipMemcpy(hostArr, devArr, Nbytes, hipMemcpyDeviceToHost));
+  REQUIRE(validateHostArray(hostArr, N, 16) == true);
+
+  // Validating hipGetFuncBySymbol API
+  hipFunction_t functionWithOrgApi, functionWithFuncPtr;
+  HIP_CHECK(hipGetFuncBySymbol(&functionWithOrgApi,
+            reinterpret_cast<const void*>(addOneKernel)));
+  REQUIRE(functionWithOrgApi != nullptr);
+
+  HIP_CHECK(dyn_hipGetFuncBySymbol_ptr(&functionWithFuncPtr,
+            reinterpret_cast<const void*>(addOneKernel)));
+  REQUIRE(functionWithFuncPtr != nullptr);
+
+  REQUIRE(functionWithFuncPtr == functionWithOrgApi);
+
+  // Validating hipFuncGetAttributes API
+  struct hipFuncAttributes attrWithOrgApi, attrWithFuncPtr;
+
+  HIP_CHECK(hipFuncGetAttributes(&attrWithOrgApi,
+            reinterpret_cast<const void*>(addOneKernel)));
+  HIP_CHECK(dyn_hipFuncGetAttributes_ptr(&attrWithFuncPtr,
+            reinterpret_cast<const void*>(addOneKernel)));
+
+  REQUIRE(attrWithFuncPtr.binaryVersion == attrWithOrgApi.binaryVersion);
+  REQUIRE(attrWithFuncPtr.cacheModeCA == attrWithOrgApi.cacheModeCA);
+  REQUIRE(attrWithFuncPtr.constSizeBytes == attrWithOrgApi.constSizeBytes);
+  REQUIRE(attrWithFuncPtr.localSizeBytes == attrWithOrgApi.localSizeBytes);
+  REQUIRE(attrWithFuncPtr.maxDynamicSharedSizeBytes ==
+          attrWithOrgApi.maxDynamicSharedSizeBytes);
+  REQUIRE(attrWithFuncPtr.maxThreadsPerBlock ==
+          attrWithOrgApi.maxThreadsPerBlock);
+  REQUIRE(attrWithFuncPtr.numRegs == attrWithOrgApi.numRegs);
+  REQUIRE(attrWithFuncPtr.preferredShmemCarveout ==
+          attrWithOrgApi.preferredShmemCarveout);
+  REQUIRE(attrWithFuncPtr.ptxVersion == attrWithOrgApi.ptxVersion);
+  REQUIRE(attrWithFuncPtr.sharedSizeBytes == attrWithOrgApi.sharedSizeBytes);
+
+  // Validating hipFuncGetAttribute API
+  hipFunction_attribute attributes[] = {
+                        HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK,
+                        HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,
+                        HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES,
+                        HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES,
+                        HIP_FUNC_ATTRIBUTE_NUM_REGS,
+                        HIP_FUNC_ATTRIBUTE_PTX_VERSION,
+                        HIP_FUNC_ATTRIBUTE_BINARY_VERSION,
+                        HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA,
+                        HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES,
+                        HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT};
+
+  for ( auto attribute : attributes ) {
+    int valuewithOrgAPI = 0, valueWithFuncPointer = 0;
+
+    HIP_CHECK(hipFuncGetAttribute(&valuewithOrgAPI, attribute, function));
+    HIP_CHECK(dyn_hipFuncGetAttribute_ptr(&valueWithFuncPointer, attribute,
+                                          function));
+
+    REQUIRE(valueWithFuncPointer == valuewithOrgAPI);
+  }
+
+  // Validating hipModuleGetGlobal API
+  hipDeviceptr_t dptrWithOrgApi = nullptr;
+  size_t bytesWithOrgApi = 0;
+  HIP_CHECK(hipModuleGetGlobal(&dptrWithOrgApi, &bytesWithOrgApi,
+                                module, "globalDevData"));
+  REQUIRE(dptrWithOrgApi != nullptr);
+
+  hipDeviceptr_t dptrWithFuncPtr = nullptr;
+  size_t bytesWithFuncPtr = 0;
+  HIP_CHECK(dyn_hipModuleGetGlobal_ptr(&dptrWithFuncPtr, &bytesWithFuncPtr,
+                                       module, "globalDevData") );
+  REQUIRE(dptrWithFuncPtr != nullptr);
+  REQUIRE(bytesWithFuncPtr == 4);
+
+  REQUIRE(dptrWithFuncPtr == dptrWithOrgApi);
+  REQUIRE(bytesWithFuncPtr == bytesWithOrgApi);
+
+  // Validating hipModuleUnload API
+  HIP_CHECK(dyn_hipModuleUnload_ptr(module));
+  REQUIRE(dyn_hipModuleUnload_ptr(module) == hipErrorNotFound);
+
+  free(hostArr);
+  HIP_CHECK(hipFree(devArr));
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - This test will get the function pointer of different module management
+ *  - (load data) related APIs from the hipGetProcAddress API
+ *  - and then validates the basic functionality of that particular API
+ *  - using the funtion pointer.
+ * Test source
+ * ------------------------
+ *  - unit/module/hipGetProcAddress_Module_APIs.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipGetProcAddress_ModuleApisLoadData") {
+  void* hipModuleLoadData_ptr = nullptr;
+  void* hipModuleLoadDataEx_ptr = nullptr;
+
+  int currentHipVersion = 0;
+  HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+  HIP_CHECK(hipGetProcAddress("hipModuleLoadData",
+                              &hipModuleLoadData_ptr,
+                              currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress("hipModuleLoadDataEx",
+                              &hipModuleLoadDataEx_ptr,
+                              currentHipVersion, 0, nullptr));
+
+  hipError_t (*dyn_hipModuleLoadData_ptr)(hipModule_t *, const void *) =
+    reinterpret_cast<hipError_t (*)(hipModule_t *, const void *)>
+                                      (hipModuleLoadData_ptr);
+  hipError_t (*dyn_hipModuleLoadDataEx_ptr)(hipModule_t *, const void *,
+             unsigned int, hipJitOption *, void **) =
+    reinterpret_cast<hipError_t (*)(hipModule_t *, const void *,
+                                    unsigned int, hipJitOption *, void **)>
+                                    (hipModuleLoadDataEx_ptr);
+
+  const auto rtc = CreateRTCCharArray(
+                   R"(extern "C" __global__ void simpleKernel() {})");
+
+  // Validating hipModuleLoadData API
+  {
+    hipModule_t module = nullptr;
+
+    HIP_CHECK(dyn_hipModuleLoadData_ptr(&module, rtc.data()));
+    REQUIRE(module != nullptr);
+
+    hipFunction_t function;
+    HIP_CHECK(hipModuleGetFunction(&function, module, "simpleKernel"));
+    REQUIRE(function != nullptr);
+    HIP_CHECK(hipModuleLaunchKernel(function,
+              1, 1, 1, 1, 1, 1,
+              0, 0, nullptr, nullptr));
+
+    HIP_CHECK(hipModuleUnload(module));
+  }
+
+  // Validating hipModuleLoadDataEx API
+  {
+    hipModule_t module = nullptr;
+
+    HIP_CHECK(dyn_hipModuleLoadDataEx_ptr(&module, rtc.data(),
+                                          0, nullptr, nullptr));
+    REQUIRE(module != nullptr);
+
+    hipFunction_t function;
+    HIP_CHECK(hipModuleGetFunction(&function, module, "simpleKernel"));
+    REQUIRE(function != nullptr);
+    HIP_CHECK(hipModuleLaunchKernel(function,
+              1, 1, 1, 1, 1, 1,
+              0, 0, nullptr, nullptr));
+
+    HIP_CHECK(hipModuleUnload(module));
+  }
+}
+
+ /**
+ * Test Description
+ * ------------------------
+ *  - This test will get the function pointer of different module management
+ *  - (Cooperative Kernels) related APIs from the hipGetProcAddress API
+ *  - and then validates the basic functionality of that particular API
+ *  - using the funtion pointer.
+ * Test source
+ * ------------------------
+ *  - unit/module/hipGetProcAddress_Module_APIs.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipGetProcAddress_ModuleApisCooperativeKernels") {
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
+    HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
+    return;
+  }
+
+  void *hipModuleLaunchCooperativeKernel_ptr = nullptr;
+  void *hipModuleLaunchCooperativeKernelMultiDevice_ptr = nullptr;
+  void *hipLaunchCooperativeKernel_ptr = nullptr;
+  void *hipLaunchCooperativeKernelMultiDevice_ptr = nullptr;
+  void *hipExtLaunchMultiKernelMultiDevice_ptr = nullptr;
+
+  int currentHipVersion = 0;
+  HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+  HIP_CHECK(hipGetProcAddress(
+            "hipModuleLaunchCooperativeKernel",
+            &hipModuleLaunchCooperativeKernel_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipModuleLaunchCooperativeKernelMultiDevice",
+            &hipModuleLaunchCooperativeKernelMultiDevice_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipLaunchCooperativeKernel",
+            &hipLaunchCooperativeKernel_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipLaunchCooperativeKernelMultiDevice",
+            &hipLaunchCooperativeKernelMultiDevice_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipExtLaunchMultiKernelMultiDevice",
+            &hipExtLaunchMultiKernelMultiDevice_ptr,
+            currentHipVersion, 0, nullptr));
+
+  hipError_t (*dyn_hipModuleLaunchCooperativeKernel_ptr)(
+             hipFunction_t,
+             unsigned int, unsigned int, unsigned int,
+             unsigned int, unsigned int, unsigned int,
+             unsigned int, hipStream_t, void **) =
+    reinterpret_cast<hipError_t (*)(hipFunction_t,
+             unsigned int, unsigned int, unsigned int,
+             unsigned int, unsigned int, unsigned int,
+             unsigned int, hipStream_t, void **)>
+             (hipModuleLaunchCooperativeKernel_ptr);
+
+  hipError_t (*dyn_hipModuleLaunchCooperativeKernelMultiDevice_ptr)(
+             hipFunctionLaunchParams *, unsigned int, unsigned int) =
+    reinterpret_cast<hipError_t (*)(hipFunctionLaunchParams *,
+             unsigned int, unsigned int)>
+            (hipModuleLaunchCooperativeKernelMultiDevice_ptr);
+
+  hipError_t (*dyn_hipLaunchCooperativeKernel_ptr)(
+             const void *, dim3, dim3, void **, unsigned int, hipStream_t) =
+    reinterpret_cast<hipError_t (*)(const void *, dim3, dim3, void **,
+             unsigned int, hipStream_t)>
+             (hipLaunchCooperativeKernel_ptr);
+
+  hipError_t (*dyn_hipLaunchCooperativeKernelMultiDevice_ptr)(
+             hipLaunchParams *, int, unsigned int) =
+    reinterpret_cast<hipError_t (*)(hipLaunchParams *, int, unsigned int)>
+             (hipLaunchCooperativeKernelMultiDevice_ptr);
+
+  hipError_t (*dyn_hipExtLaunchMultiKernelMultiDevice_ptr)(
+             hipLaunchParams *, int, unsigned int) =
+    reinterpret_cast<hipError_t (*)(hipLaunchParams *, int, unsigned int)>
+             (hipExtLaunchMultiKernelMultiDevice_ptr);
+
+  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 kernelParameters {
+    void *arr;
+    int size;
+  };
+  kernelParameters kernelParam;
+  kernelParam.arr = devArr;
+  kernelParam.size = N;
+  void *kernel_parameter[] = {&kernelParam.arr, &kernelParam.size};
+
+  // Validating hipModuleLaunchCooperativeKernel API
+  {
+    hipModule_t module;
+    HIP_CHECK(hipModuleLoad(&module, "addKernel.code"));
+    REQUIRE(module != nullptr);
+
+    hipFunction_t function;
+    HIP_CHECK(hipModuleGetFunction(&function, module, "addKernel"));
+    REQUIRE(function != nullptr);
+
+    HIP_CHECK(dyn_hipModuleLaunchCooperativeKernel_ptr(function,
+              blocksPerGrid.x, blocksPerGrid.y, blocksPerGrid.z,
+              threadsPerBlock.x, threadsPerBlock.y, threadsPerBlock.z,
+              0, 0, kernel_parameter));
+
+    HIP_CHECK(hipMemcpy(hostArr, devArr, Nbytes, hipMemcpyDeviceToHost));
+    REQUIRE(validateHostArray(hostArr, N, 12) == true);
+    HIP_CHECK(hipModuleUnload(module));
+  }
+
+  // Validating hipModuleLaunchCooperativeKernelMultiDevice API
+  {
+    int deviceCount = 0;
+    HIP_CHECK(hipGetDeviceCount(&deviceCount));
+
+    hipModule_t *module = new hipModule_t[deviceCount];
+    hipFunction_t *function = new hipFunction_t[deviceCount];
+    hipStream_t *streamArr = new hipStream_t[deviceCount];
+
+    for (int i = 0; i < deviceCount; ++i) {
+      HIP_CHECK(hipSetDevice(i));
+      HIP_CHECK(hipStreamCreate(&streamArr[i]));
+
+      HIP_CHECK(hipModuleLoad(&module[i], "addKernel.code"));
+      REQUIRE(module[i] != nullptr);
+
+      HIP_CHECK(hipModuleGetFunction(&function[i], module[i],
+                                     "sampleModuleKernel"));
+      REQUIRE(function[i] != nullptr);
+    }
+
+    HIP_CHECK(hipSetDevice(0));
+
+    ::std::vector<hipFunctionLaunchParams> params(deviceCount);
+
+    for (int i = 0; i < deviceCount; ++i) {
+      params[i].function = function[i];
+      params[i].gridDimX = 1;
+      params[i].gridDimY = 1;
+      params[i].gridDimZ = 1;
+      params[i].blockDimX = 1;
+      params[i].blockDimY = 1;
+      params[i].blockDimZ = 1;
+      params[i].kernelParams = nullptr;
+      params[i].sharedMemBytes = 0;
+      params[i].hStream = streamArr[i];
+    }
+
+    HIP_CHECK(dyn_hipModuleLaunchCooperativeKernelMultiDevice_ptr(
+              params.data(), deviceCount, 0));
+
+    for (int i = 0; i < deviceCount; ++i) {
+      HIP_CHECK(hipStreamSynchronize(params[i].hStream));
+    }
+
+    for (int i = 0; i < deviceCount; ++i) {
+      HIP_CHECK(hipStreamDestroy(streamArr[i]));
+      HIP_CHECK(hipModuleUnload(module[i]));
+    }
+  }
+
+  // Validating hipLaunchCooperativeKernel API
+  {
+    HIP_CHECK(dyn_hipLaunchCooperativeKernel_ptr(
+              reinterpret_cast<void *>(addOneKernel),
+              dim3(1, 1, 1), dim3(1, 1, 1),
+              kernel_parameter, 0, 0));
+    HIP_CHECK(hipMemcpy(hostArr, devArr, Nbytes, hipMemcpyDeviceToHost));
+    REQUIRE(validateHostArray(hostArr, N, 13) == true);
+  }
+
+  // Validating hipLaunchCooperativeKernelMultiDevice API
+  {
+    int deviceCount = 0;
+    HIP_CHECK(hipGetDeviceCount(&deviceCount));
+
+    hipStream_t *streamArr = new hipStream_t[deviceCount];
+
+    for (int i = 0; i < deviceCount; ++i) {
+      HIP_CHECK(hipSetDevice(i));
+      HIP_CHECK(hipStreamCreate(&streamArr[i]));
+    }
+
+    std::vector<hipLaunchParams> params(deviceCount);
+
+    for (int i = 0; i < deviceCount; ++i) {
+      params[i].func = reinterpret_cast<void *>(sampleKernel);
+      params[i].gridDim = {1, 1, 1};
+      params[i].blockDim = {1, 1, 1};
+      params[i].args = nullptr;
+      params[i].sharedMem = 0;
+      params[i].stream = streamArr[i];
+    }
+
+    HIP_CHECK(dyn_hipLaunchCooperativeKernelMultiDevice_ptr(
+              params.data(), deviceCount, 0));
+
+    for (int i = 0; i < deviceCount; ++i) {
+      HIP_CHECK(hipStreamSynchronize(params[i].stream));
+    }
+
+    for (int i = 0; i < deviceCount; ++i) {
+      HIP_CHECK(hipStreamDestroy(streamArr[i]));
+    }
+  }
+
+  // Validating hipExtLaunchMultiKernelMultiDevice API
+  {
+    int deviceCount = 0;
+    HIP_CHECK(hipGetDeviceCount(&deviceCount));
+
+    hipStream_t *streamArr = new hipStream_t[deviceCount];
+
+    for (int i = 0; i < deviceCount; ++i) {
+      HIP_CHECK(hipSetDevice(i));
+      HIP_CHECK(hipStreamCreate(&streamArr[i]));
+    }
+
+    std::vector<hipLaunchParams> params(deviceCount);
+
+    for (int i = 0; i < deviceCount; ++i) {
+      params[i].func = reinterpret_cast<void *>(sampleKernel);
+      params[i].gridDim = {1, 1, 1};
+      params[i].blockDim = {1, 1, 1};
+      params[i].args = nullptr;
+      params[i].sharedMem = 0;
+      params[i].stream = streamArr[i];
+    }
+
+    HIP_CHECK(dyn_hipExtLaunchMultiKernelMultiDevice_ptr(
+              params.data(), deviceCount, 0));
+
+    for (int i = 0; i < deviceCount; ++i) {
+      HIP_CHECK(hipStreamSynchronize(params[i].stream));
+    }
+
+    for (int i = 0; i < deviceCount; ++i) {
+      HIP_CHECK(hipStreamDestroy(streamArr[i]));
+    }
+  }
+
+  free(hostArr);
+  HIP_CHECK(hipFree(devArr));
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - This test will get the function pointer of different Occupancy
+ *  - related APIs from the hipGetProcAddress API
+ *  - and then validates the basic functionality of that particular API
+ *  - using the funtion pointer.
+ * Test source
+ * ------------------------
+ *  - unit/module/hipGetProcAddress_Module_APIs.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipGetProcAddress_ModuleApisOccupancy") {
+  void* hipModuleOccupancyMaxPotentialBlockSize_ptr = nullptr;
+  void* hipModuleOccupancyMaxPotentialBlockSizeWithFlags_ptr = nullptr;
+  void* hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_ptr = nullptr;
+  void* hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_ptr =
+        nullptr;
+  void* hipOccupancyMaxActiveBlocksPerMultiprocessor_ptr = nullptr;
+  void* hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_ptr = nullptr;
+  void* hipOccupancyMaxPotentialBlockSize_ptr = nullptr;
+
+  int currentHipVersion = 0;
+  HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
+
+  HIP_CHECK(hipGetProcAddress(
+            "hipModuleOccupancyMaxPotentialBlockSize",
+            &hipModuleOccupancyMaxPotentialBlockSize_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipModuleOccupancyMaxPotentialBlockSizeWithFlags",
+            &hipModuleOccupancyMaxPotentialBlockSizeWithFlags_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipModuleOccupancyMaxActiveBlocksPerMultiprocessor",
+            &hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags",
+            &hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipOccupancyMaxActiveBlocksPerMultiprocessor",
+            &hipOccupancyMaxActiveBlocksPerMultiprocessor_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags",
+            &hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_ptr,
+            currentHipVersion, 0, nullptr));
+  HIP_CHECK(hipGetProcAddress(
+            "hipOccupancyMaxPotentialBlockSize",
+            &hipOccupancyMaxPotentialBlockSize_ptr,
+            currentHipVersion, 0, nullptr));
+
+  hipError_t(*dyn_hipModuleOccupancyMaxPotentialBlockSize_ptr)(
+  int *, int *, hipFunction_t, size_t, int) =
+  reinterpret_cast<hipError_t (*)(int *, int *, hipFunction_t, size_t, int)>
+  (hipModuleOccupancyMaxPotentialBlockSize_ptr);
+
+  hipError_t(*dyn_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_ptr)(
+  int *, int *, hipFunction_t, size_t, int, unsigned int) =
+  reinterpret_cast<hipError_t (*)(int *, int *, hipFunction_t,
+  size_t, int, unsigned int)>
+  (hipModuleOccupancyMaxPotentialBlockSizeWithFlags_ptr);
+
+  hipError_t(*dyn_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_ptr)(
+  int *, hipFunction_t, int, size_t) =
+  reinterpret_cast<hipError_t (*)(int *, hipFunction_t, int, size_t)>
+  (hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_ptr);
+
+  hipError_t(
+  *dyn_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_ptr)(
+  int *, hipFunction_t, int, size_t, unsigned int) =
+  reinterpret_cast<hipError_t (*)(int *, hipFunction_t, int,
+  size_t, unsigned int)>
+  (hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_ptr);
+
+  hipError_t(*dyn_hipOccupancyMaxActiveBlocksPerMultiprocessor_ptr)(
+  int *, const void *, int, size_t) =
+  reinterpret_cast<hipError_t (*)(int *, const void *, int, size_t)>
+  (hipOccupancyMaxActiveBlocksPerMultiprocessor_ptr);
+
+  hipError_t(*dyn_hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_ptr)(
+  int *, const void *, int, size_t, unsigned int) =
+  reinterpret_cast<hipError_t (*)(int *, const void *,
+  int, size_t, unsigned int)>
+  (hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_ptr);
+
+  hipError_t(*dyn_hipOccupancyMaxPotentialBlockSize_ptr)(
+  int *, int *, const void *, size_t, int) =
+  reinterpret_cast<hipError_t (*)(int *, int *, const void *, size_t, int)>
+  (hipOccupancyMaxPotentialBlockSize_ptr);
+
+  hipModule_t module;
+  HIP_CHECK(hipModuleLoad(&module, "addKernel.code"));
+  REQUIRE(module != nullptr);
+  hipFunction_t function;
+  HIP_CHECK(hipModuleGetFunction(&function, module, "addKernel"));
+  REQUIRE(function != nullptr);
+
+  int gridSize = 0, blockSize = 0;
+  int gridSizeWithFuncPtr = 0, blockSizeWithFuncPtr = 0;
+
+  // Validating hipModuleOccupancyMaxPotentialBlockSize API
+  {
+    HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize,
+                                                      function, 0, 0));
+    HIP_CHECK(dyn_hipModuleOccupancyMaxPotentialBlockSize_ptr(
+              &gridSizeWithFuncPtr, &blockSizeWithFuncPtr, function, 0, 0));
+
+    REQUIRE(gridSizeWithFuncPtr == gridSize);
+    REQUIRE(blockSizeWithFuncPtr == blockSize);
+  }
+
+  // Validating hipModuleOccupancyMaxPotentialBlockSizeWithFlags API
+  {
+    gridSize = 0; blockSize = 0;
+    gridSizeWithFuncPtr = 0; blockSizeWithFuncPtr = 0;
+    HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSizeWithFlags(
+              &gridSize, &blockSize, function, 0, 0, 0));
+    HIP_CHECK(dyn_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_ptr(
+              &gridSizeWithFuncPtr, &blockSizeWithFuncPtr, function, 0, 0, 0));
+
+    REQUIRE(gridSizeWithFuncPtr == gridSize);
+    REQUIRE(blockSizeWithFuncPtr == blockSize);
+  }
+
+  int numBlocks = 0, numBlocksWithFuncPtr = 0;
+  // Validating hipModuleOccupancyMaxActiveBlocksPerMultiprocessor API
+  {
+    HIP_CHECK(hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(
+              &numBlocks, function, blockSize, 0));
+    HIP_CHECK(dyn_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_ptr(
+              &numBlocksWithFuncPtr, function, blockSize, 0));
+
+    REQUIRE(numBlocksWithFuncPtr == numBlocks);
+  }
+
+  // Validating hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags API
+  {
+    numBlocks = 0; numBlocksWithFuncPtr = 0;
+    HIP_CHECK(hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+              &numBlocks, function, blockSize, 0, 0));
+    HIP_CHECK(
+    dyn_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_ptr(
+     &numBlocksWithFuncPtr, function, blockSize, 0, 0));
+
+    REQUIRE(numBlocksWithFuncPtr == numBlocks);
+  }
+
+  // Validating hipOccupancyMaxActiveBlocksPerMultiprocessor API
+  {
+    numBlocks = 0; numBlocksWithFuncPtr = 0;
+    HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(
+              &numBlocks, reinterpret_cast<const void *>(addOneKernel),
+              blockSize, 0));
+    HIP_CHECK(dyn_hipOccupancyMaxActiveBlocksPerMultiprocessor_ptr(
+              &numBlocksWithFuncPtr,
+              reinterpret_cast<const void *>(addOneKernel), blockSize, 0));
+
+    REQUIRE(numBlocksWithFuncPtr == numBlocks);
+  }
+
+  // Validating hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags API
+  {
+    numBlocks = 0; numBlocksWithFuncPtr = 0;
+    HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+              &numBlocks, reinterpret_cast<const void *>(addOneKernel),
+              blockSize, 0, 0));
+    HIP_CHECK(dyn_hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_ptr(
+              &numBlocksWithFuncPtr,
+              reinterpret_cast<const void *>(addOneKernel),
+              blockSize, 0, 0));
+
+    REQUIRE(numBlocksWithFuncPtr == numBlocks);
+  }
+
+  // Validating hipOccupancyMaxPotentialBlockSize API
+  {
+    gridSize = 0; blockSize = 0;
+    gridSizeWithFuncPtr = 0; blockSizeWithFuncPtr = 0;
+    HIP_CHECK(hipOccupancyMaxPotentialBlockSize(
+              &gridSize, &blockSize,
+              reinterpret_cast<const void *>(addOneKernel), 0, 0));
+    HIP_CHECK(dyn_hipOccupancyMaxPotentialBlockSize_ptr(
+              &gridSizeWithFuncPtr, &blockSizeWithFuncPtr,
+              reinterpret_cast<const void *>(addOneKernel), 0, 0));
+
+    REQUIRE(gridSizeWithFuncPtr == gridSize);
+    REQUIRE(blockSizeWithFuncPtr == blockSize);
+  }
+
+  HIP_CHECK(hipModuleUnload(module));
+}