From 2be5ef4dfba7392fe92dfffe27114ccc1836b45a Mon Sep 17 00:00:00 2001
From: Marko Arandjelovic <Marko.Arandjelovic@amd.com>
Date: Wed, 18 Dec 2024 16:37:55 +0000
Subject: [PATCH] SWDEV-504084 - Added multidevice testcases for
 hipModuleGetFunction/Global

Change-Id: I53d6d63f93c800efa2ec0ff2be5ae593756549e0


[ROCm/hip-tests commit: cddb658314408e7e2fa100be034935c0f6166fa0]
---
 .../catch/unit/module/hipModuleGetFunction.cc | 19 +++++++++++++++++-
 .../catch/unit/module/hipModuleGetGlobal.cc   | 20 ++++++++++++++++++-
 2 files changed, 37 insertions(+), 2 deletions(-)

diff --git a/projects/hip-tests/catch/unit/module/hipModuleGetFunction.cc b/projects/hip-tests/catch/unit/module/hipModuleGetFunction.cc
index 676a61c33d..c4decba9d1 100644
--- a/projects/hip-tests/catch/unit/module/hipModuleGetFunction.cc
+++ b/projects/hip-tests/catch/unit/module/hipModuleGetFunction.cc
@@ -71,4 +71,21 @@ TEST_CASE("Unit_hipModuleGetFunction_Negative_Parameters") {
   SECTION("kname == __device__ kernel") {
     HIP_CHECK_ERROR(hipModuleGetFunction(&kernel, GetModule(), "DeviceKernel"), hipErrorNotFound);
   }
-}
\ No newline at end of file
+}
+
+// Test description: Loading kernel function from different device than the one on which the module
+// is loaded
+TEST_CASE("Unit_hipModuleGetFunction_DiffDevice") {
+  int numDevices = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices < 2) {
+    SUCCEED("skipped the testcase as no of devices is less than 2");
+    return;
+  }
+
+  hipFunction_t kernel = nullptr;
+  auto module = GetModule();
+  HIP_CHECK(hipSetDevice(1));
+  HIP_CHECK(hipModuleGetFunction(&kernel, module, "GlobalKernel"));
+  REQUIRE(kernel != nullptr);
+}
diff --git a/projects/hip-tests/catch/unit/module/hipModuleGetGlobal.cc b/projects/hip-tests/catch/unit/module/hipModuleGetGlobal.cc
index 32f46f0caf..a4a5423d14 100644
--- a/projects/hip-tests/catch/unit/module/hipModuleGetGlobal.cc
+++ b/projects/hip-tests/catch/unit/module/hipModuleGetGlobal.cc
@@ -142,4 +142,22 @@ TEST_CASE("Unit_hipModuleGetGlobal_Negative_Name_Is_Empty_String") {
 TEST_CASE("Unit_hipModuleGetGlobal_Negative_Dptr_And_Bytes_Are_Nullptr") {
   hipModule_t module = GetModule();
   HIP_CHECK_ERROR(hipModuleGetGlobal(nullptr, nullptr, module, "int_var"), hipErrorInvalidValue);
-}
\ No newline at end of file
+}
+
+// Test description: Loading device ptr from different device than the one on which the module
+// is loaded
+TEST_CASE("Unit_hipModuleGetGlobal_DiffDevice") {
+  int numDevices = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices < 2) {
+    SUCCEED("skipped the testcase as no of devices is less than 2");
+    return;
+  }
+  auto module = GetModule();
+  HIP_CHECK(hipSetDevice(1));
+  hipDeviceptr_t global;
+  size_t global_size = 0;
+  HIP_CHECK(hipModuleGetGlobal(&global, &global_size, module, "int_var"));
+  REQUIRE(global != 0);
+  REQUIRE(sizeof(int) == global_size);
+}