From 8417e0b4426a6bf6bd7f670a473e54fd138f7236 Mon Sep 17 00:00:00 2001
From: Satyanvesh Dittakavi <Satyanvesh.Dittakavi@amd.com>
Date: Mon, 6 May 2024 11:24:46 +0000
Subject: [PATCH] SWDEV-458637 - Adds test for bitcode compilation with
 undefined function

Change-Id: Ic94d2a4a99ce6993a9a11e74192e13c90a99598e
---
 catch/unit/rtc/CMakeLists.txt                |  1 +
 catch/unit/rtc/hiprtc_Bitcode_UndefinedFn.cc | 74 ++++++++++++++++++++
 2 files changed, 75 insertions(+)
 create mode 100644 catch/unit/rtc/hiprtc_Bitcode_UndefinedFn.cc

diff --git a/catch/unit/rtc/CMakeLists.txt b/catch/unit/rtc/CMakeLists.txt
index 50945f9486..df60e2fce3 100644
--- a/catch/unit/rtc/CMakeLists.txt
+++ b/catch/unit/rtc/CMakeLists.txt
@@ -24,6 +24,7 @@ set(AMD_TEST_SRC
     hiprtc_TextureTypes_HeaderTst.cc
     hipRtcComplexHeader.cc
     hipRtcPtrdiff_t.cc
+    hiprtc_Bitcode_UndefinedFn.cc
 )
 
 add_custom_target(copyRtcHeaders ALL
diff --git a/catch/unit/rtc/hiprtc_Bitcode_UndefinedFn.cc b/catch/unit/rtc/hiprtc_Bitcode_UndefinedFn.cc
new file mode 100644
index 0000000000..cfbe7dccd6
--- /dev/null
+++ b/catch/unit/rtc/hiprtc_Bitcode_UndefinedFn.cc
@@ -0,0 +1,74 @@
+/*
+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/hiprtc.h>
+#include <hip/hip_runtime.h>
+
+#include <cassert>
+#include <cstddef>
+#include <memory>
+#include <iostream>
+#include <iterator>
+#include <vector>
+
+// This test verifies successful compilation of source to bitcode (with -fgpu-rdc option) when the
+// hiprtc kernel has an undefined function which will be linked in the later stages using hiprtc
+// linker APIs
+
+TEST_CASE("Unit_hiprtc_bitcode_undefined_function") {
+  using namespace std;
+
+  static constexpr auto kernel {
+  R"(
+      __device__ void foo();
+
+      extern "C"
+      __global__ void gpu_kernel() {
+          foo();
+      }
+  )"};
+
+  hiprtcProgram prog;
+  hiprtcCreateProgram(&prog,            // prog
+                      kernel,           // buffer
+                      "gpu_kernel.cu",  // name
+                      0, nullptr, nullptr);
+  hipDeviceProp_t props;
+  int device = 0;
+  HIP_CHECK(hipGetDeviceProperties(&props, device));
+
+  std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
+  const char* options[] = {sarg.c_str(), "-fgpu-rdc"};
+
+  hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)};
+
+  size_t logSize;
+  HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
+  if (logSize) {
+    string log(logSize, '\0');
+    HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
+    std::cout << log << '\n';
+  }
+
+  REQUIRE(compileResult == HIPRTC_SUCCESS);
+
+  HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
+
+}