From a224eef6247a33740484f0106de702b2fcc50d40 Mon Sep 17 00:00:00 2001
From: kjayapra-amd <karthik.jayaprakash@amd.com>
Date: Thu, 8 Dec 2022 11:44:33 -0800
Subject: [PATCH] SWDEV-359061 - Tests for hiprtcLink File and Data.

Change-Id: I2e75f3d4e50e410db31c8906cad2453da158816d
---
 catch/unit/rtc/CMakeLists.txt    |   3 +
 catch/unit/rtc/hiprtcHelper.cc   | 318 +++++++++++++++++++++++++++++++
 catch/unit/rtc/hiprtcHelper.hpp  |  49 +++++
 catch/unit/rtc/hiprtcLinkData.cc | 106 +++++++++++
 catch/unit/rtc/hiprtcLinkFile.cc | 111 +++++++++++
 5 files changed, 587 insertions(+)
 create mode 100644 catch/unit/rtc/hiprtcHelper.cc
 create mode 100644 catch/unit/rtc/hiprtcHelper.hpp
 create mode 100644 catch/unit/rtc/hiprtcLinkData.cc
 create mode 100644 catch/unit/rtc/hiprtcLinkFile.cc

diff --git a/catch/unit/rtc/CMakeLists.txt b/catch/unit/rtc/CMakeLists.txt
index 534289455..179471360 100644
--- a/catch/unit/rtc/CMakeLists.txt
+++ b/catch/unit/rtc/CMakeLists.txt
@@ -3,6 +3,9 @@ set(TEST_SRC
     saxpy.cc
     warpsize.cc
     hipRtcFunctional.cc
+    hiprtcHelper.cc
+    hiprtcLinkFile.cc
+    hiprtcLinkData.cc
 )
 
 # AMD only tests
diff --git a/catch/unit/rtc/hiprtcHelper.cc b/catch/unit/rtc/hiprtcHelper.cc
new file mode 100644
index 000000000..be2d0ab51
--- /dev/null
+++ b/catch/unit/rtc/hiprtcHelper.cc
@@ -0,0 +1,318 @@
+#include "hiprtcHelper.hpp"
+
+#include <hip_test_common.hh>
+
+#include <iostream>
+#include <fstream>
+
+#if defined(_WIN32)
+#include <windows.h>
+#else
+#include <unistd.h>
+#include <fcntl.h>
+#include <sys/mman.h>
+#include <sys/stat.h>
+#endif
+
+#include <hip/hiprtc.h>
+#include <hip/hip_runtime.h>
+
+#define kernel_name "vcpy_kernel"
+#define kernel_name0 "testKernel"
+
+bool CommitBCToFile(char* executable, size_t exe_size, const std::string& bit_code_file) {
+  std::fstream bc_file;
+  bc_file.open(bit_code_file, std::ios::out | std::ios::binary);
+  if (!bc_file) {
+    std::cout << "File not created" << std::endl;
+  }
+
+  // std::cout<<"EXE SIZE: "<<exe_size<<std::endl;
+  bc_file.write(executable, exe_size);
+
+  bc_file.close();
+  return true;
+}
+
+bool TestCompileProgram(const char* source_code, hiprtcProgram* prog_ptr, std::string prog_name,
+                        std::string func_name) {
+  hiprtcCreateProgram(prog_ptr, source_code, prog_name.c_str(), 0, nullptr, nullptr);
+  hiprtcAddNameExpression(*prog_ptr, func_name.c_str());
+
+  std::cout<<"Func Name: "<<func_name<<std::endl;
+
+  int device = 0;
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, device));
+  std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
+  std::string srdc = std::string("-fgpu-rdc");
+
+  const char* options[] = {sarg.c_str(), srdc.c_str()};
+  hiprtcResult compileResult(hiprtcCompileProgram(*prog_ptr, 2, options));
+
+  size_t logSize;
+  hiprtcGetProgramLogSize(*prog_ptr, &logSize);
+
+  if (logSize) {
+    std::string log(logSize, '\0');
+    hiprtcGetProgramLog(*prog_ptr, &log[0]);
+    std::cout << log << std::endl;
+  }
+
+  if (compileResult != HIPRTC_SUCCESS) {
+    std::cout << "Compilation failed." << std::endl;
+  }
+ 
+  return true;
+}
+
+bool TestCompileRDC(char** bit_code_pptr, size_t* bit_code_size_ptr, const char* routine_ptr,
+                    std::string routine_name) {
+  hiprtcProgram prog;
+  hiprtcCreateProgram(&prog, routine_ptr, routine_name.c_str(), 0, nullptr, nullptr);
+
+  int device = 0;
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, device));
+  std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
+  std::string srdc = std::string("-fgpu-rdc");
+
+  const char* options[] = {sarg.c_str(), srdc.c_str()};
+
+  hiprtcResult compileResult(hiprtcCompileProgram(prog, 2, options));
+
+  size_t logSize;
+  hiprtcGetProgramLogSize(prog, &logSize);
+
+  if (logSize) {
+    std::string log(logSize, '\0');
+    hiprtcGetProgramLog(prog, &log[0]);
+    std::cout << log << std::endl;
+  }
+
+  if (compileResult != HIPRTC_SUCCESS) {
+    std::cout << "Compilation failed." << std::endl;
+  }
+
+  hiprtcGetBitcodeSize(prog, bit_code_size_ptr);
+  *bit_code_pptr = new char[*bit_code_size_ptr];
+  hiprtcGetBitcode(prog, *bit_code_pptr);
+
+  return true;
+}
+
+bool TestModuleLoadData(void* cuOut) {
+  bool test_passed = true;
+
+  size_t LEN = 64;
+  size_t SIZE = sizeof(float) * LEN;
+  float* A = nullptr;
+  float* B = nullptr;
+
+  hipDeviceptr_t Ad, Bd;
+  A = new float[LEN];
+  B = new float[LEN];
+
+  for (size_t idx = 0; idx < LEN; ++idx) {
+    A[idx] = idx;
+    B[idx] = 0;
+  }
+
+  for (size_t idx = 0; idx < LEN; ++idx) {
+    // std::cout<<"AT_INIT --> idx: "<<idx<<" A: "<<A[idx]<<" B: "<<B[idx]<<std::endl;
+  }
+
+  HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
+  HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
+  HIPCHECK(hipMemcpyHtoD(Ad, A, SIZE));
+  HIPCHECK(hipMemcpyHtoD(Bd, B, SIZE));
+
+  hipStream_t hip_stream;
+  HIPCHECK(hipStreamCreate(&hip_stream));
+
+  struct {
+    void* _Ad;
+    void* _Bd;
+  } args;
+  args._Ad = reinterpret_cast<void*>(Ad);
+  args._Bd = reinterpret_cast<void*>(Bd);
+  size_t args_size = sizeof(args);
+
+  void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE,
+                    &args_size, HIP_LAUNCH_PARAM_END};
+
+  hipModule_t hip_module = nullptr;
+  hipFunction_t hip_function = nullptr;
+  HIPCHECK(hipModuleLoadData(&hip_module, cuOut));
+  HIPCHECK(hipModuleGetFunction(&hip_function, hip_module, kernel_name));
+  HIPCHECK(hipModuleLaunchKernel(hip_function, 1, 1, 1, LEN, 1, 1, 0, hip_stream, NULL,
+                                 reinterpret_cast<void**>(&config)));
+  HIPCHECK(hipStreamSynchronize(hip_stream));
+
+  HIPCHECK(hipMemcpyDtoH(B, Bd, SIZE));
+
+  for (size_t idx = 0; idx < LEN; ++idx) {
+    if (A[idx] != B[idx]) {
+      test_passed = false;
+      std::cout << "FAIL --> idx: " << idx << " A: " << A[idx] << " B: " << B[idx] << std::endl;
+      break;
+    } else {
+      // std::cout<<"PASS --> idx: "<<idx<<" A: "<<A[idx]<<" B: "<<B[idx]<<std::endl;
+    }
+  }
+
+  HIPCHECK(hipStreamDestroy(hip_stream));
+  HIPCHECK(hipModuleUnload(hip_module));
+
+  return test_passed;
+}
+
+bool TestModuleLoad2Data(void* cuOut) {
+  bool test_passed = true;
+
+  int* x = nullptr;
+  HIPCHECK(hipMalloc(&x, sizeof(int)));
+  HIPCHECK(hipMemset(x, 0x00, sizeof(int)));
+
+  struct {
+    void* _x;
+  } args;
+  args._x = reinterpret_cast<void*>(x);
+  size_t args_size = sizeof(args);
+
+  void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE,
+                    &args_size, HIP_LAUNCH_PARAM_END};
+
+  hipModule_t hip_module = nullptr;
+  hipFunction_t hip_function = nullptr;
+  HIPCHECK(hipModuleLoadData(&hip_module, cuOut));
+  HIPCHECK(hipModuleGetFunction(&hip_function, hip_module, kernel_name0));
+  HIPCHECK(hipModuleLaunchKernel(hip_function, 1, 1, 1, 1, 1, 1, 0, NULL, NULL,
+                                 reinterpret_cast<void**>(&config)));
+  HIPCHECK(hipDeviceSynchronize());
+
+  HIPCHECK(hipModuleUnload(hip_module));
+  HIPCHECK(hipFree(x));
+  return test_passed;
+}
+
+bool OpenFileHandle(const char* fname, FileDesc* fd_ptr, size_t* sz_ptr) {
+  if ((fd_ptr == nullptr) || (sz_ptr == nullptr)) {
+    std::cout << "Invalid arguments, fname: " << fname << " fd_ptr: " << fd_ptr
+              << "sz_ptr: " << sz_ptr << std::endl;
+    return false;
+  }
+
+#if defined(_WIN32)
+  *fd_ptr = INVALID_HANDLE_VALUE;
+  *fd_ptr =
+      CreateFileA(fname, GENERIC_READ, 0x1, NULL, OPEN_EXISTING, FILE_ATTRIBUTE_READONLY, NULL);
+  if (*fd_ptr == INVALID_HANDLE_VALUE) {
+    return false;
+  }
+
+  *sz_ptr = GetFileSize(*fd_ptr, NULL);
+
+#else
+  // open system function call, return false on fail
+  struct stat stat_buf;
+  *fd_ptr = open(fname, O_RDONLY);
+  if (*fd_ptr < 0) {
+    return false;
+  }
+
+  // Retrieve stat info and size
+  if (fstat(*fd_ptr, &stat_buf) != 0) {
+    close(*fd_ptr);
+    return false;
+  }
+
+  *sz_ptr = stat_buf.st_size;
+#endif
+
+  return true;
+}
+
+bool CloseFileHandle(FileDesc fdesc) {
+#if defined(_WIN32)
+  // return false on failure
+  if (CloseHandle(fdesc) < 0) {
+    return false;
+  }
+#else
+  // Return false if close system call fails
+  if (close(fdesc) < 0) {
+    return false;
+  }
+#endif
+  return true;
+}
+
+bool MemoryMapFileDesc(FileDesc fdesc, size_t fsize, size_t foffset, const void** mmap_pptr) {
+#if defined(_WIN32)
+  if (fdesc == INVALID_HANDLE_VALUE) {
+    return false;
+  }
+
+  HANDLE map_handle = CreateFileMappingA(fdesc, NULL, PAGE_READONLY, 0, 0, NULL);
+  if (map_handle == INVALID_HANDLE_VALUE) {
+    CloseHandle(map_handle);
+    return false;
+  }
+
+  *mmap_pptr = MapViewOfFile(map_handle, FILE_MAP_READ, 0, 0, 0);
+#else
+  if (fdesc <= 0) {
+    return false;
+  }
+
+  // If the offset is not aligned then align it
+  // and recalculate the new size
+  if (foffset > 0) {
+    size_t old_foffset = foffset;
+    foffset = alignUp(foffset, 4096);
+    fsize += (foffset - old_foffset);
+  }
+
+  *mmap_pptr = mmap(NULL, fsize, PROT_READ, MAP_SHARED, fdesc, foffset);
+#endif
+  return true;
+}
+
+bool MemoryUnmapFile(const void* mmap_ptr, size_t mmap_size) {
+#if defined(_WIN32)
+  if (!UnmapViewOfFile(mmap_ptr)) {
+    std::cout << "Unmap file failed: " << mmap_ptr << std::endl;
+    return false;
+  }
+#else
+  if (munmap(const_cast<void*>(mmap_ptr), mmap_size) != 0) {
+    std::cout << "Unmap file failed: " << mmap_ptr << std::endl;
+    return false;
+  }
+#endif
+  return true;
+}
+
+bool GetMapPtr(const char* fname, FileDesc* fd_ptr, size_t* sz_ptr, const void** mmap_pptr) {
+  if (!OpenFileHandle(fname, fd_ptr, sz_ptr)) {
+    std::cout << "Opening File handle Failed: " << fname << std::endl;
+    return false;
+  }
+
+  if (!MemoryMapFileDesc(*fd_ptr, *sz_ptr, 0, mmap_pptr)) {
+    std::cout << "Memmap failed: fd_ptr: " << *fd_ptr << "fd_size" << *sz_ptr << std::endl;
+    return false;
+  }
+
+  if (!CloseFileHandle(*fd_ptr)) {
+    std::cout << "Closing the file handle Failed: " << *fd_ptr << std::endl;
+    return false;
+  }
+
+  return true;
+}
+
+bool DelMapPtr(const void* mmap_ptr, size_t mmap_size) {
+  return MemoryUnmapFile(mmap_ptr, mmap_size);
+}
diff --git a/catch/unit/rtc/hiprtcHelper.hpp b/catch/unit/rtc/hiprtcHelper.hpp
new file mode 100644
index 000000000..c9c2ad24b
--- /dev/null
+++ b/catch/unit/rtc/hiprtcHelper.hpp
@@ -0,0 +1,49 @@
+#ifndef HIPRTC_HELPER_HPP
+#define HIPRTC_HELPER_HPP
+
+#include <hip/hiprtc.h>
+
+// Header for File Desc
+#if defined(_WIN32)
+  typedef void* FileDesc;
+#else
+  typedef int FileDesc;
+#endif
+
+// Templated Helper function
+template <typename T> inline T alignDown(T value, size_t alignment) {
+  return (T)(value & ~(alignment - 1));
+}
+
+template <typename T> inline T* alignDown(T* value, size_t alignment) {
+  return (T*)alignDown((intptr_t)value, alignment);
+}
+
+template <typename T> inline T alignUp(T value, size_t alignment) {
+  return alignDown((T)(value + alignment - 1), alignment);
+}
+
+template <typename T> inline T* alignUp(T* value, size_t alignment) {
+  return (T*)alignDown((intptr_t)(value + alignment - 1), alignment);
+}
+
+// Exported Functions
+bool CommitBCToFile(char* executable, size_t exe_size, const std::string& bit_code_file);
+
+bool TestModuleLoadData(void* cuOut);
+bool TestModuleLoad2Data(void* cuOut);
+bool TestCompileRDC(char** bit_code_pptr, size_t* bit_code_size_ptr,
+                    const char* routine_ptr, std::string routine_name);
+bool TestCompileProgram(const char* source_code, hiprtcProgram* prog_ptr, std::string prog_name,
+                        std::string func_name);
+
+bool OpenFileHandle(const char* fname, FileDesc* fd_ptr, size_t* sz_ptr);
+bool CloseFileHandle(FileDesc fdesc);
+
+bool MemoryMapFileDesc(FileDesc fdesc, size_t fsize, size_t foffset, const void** mmap_pptr);
+bool MemoryUnmapFile(const void* mmap_ptr, size_t mmap_size);
+
+bool GetMapPtr(const char* fname, FileDesc* fd_ptr, size_t* sz_ptr, const void** mmap_pptr);
+bool DelMapPtr(const void* mmap_ptr, size_t mmap_size);
+
+#endif /* HIPRTC_HELPER_HPP */
diff --git a/catch/unit/rtc/hiprtcLinkData.cc b/catch/unit/rtc/hiprtcLinkData.cc
new file mode 100644
index 000000000..ad10ac23a
--- /dev/null
+++ b/catch/unit/rtc/hiprtcLinkData.cc
@@ -0,0 +1,106 @@
+/*
+ * Copyright (c) 2022 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 WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ * */
+
+#include <hip_test_common.hh>
+
+#include "hiprtcHelper.hpp"
+#include <hip/hiprtc.h>
+#include <hip/hip_runtime.h>
+
+static constexpr auto vcpy{
+  R"(
+    extern "C" __global__ void vcpy_kernel(float* a, float* b) {
+      int tx = __ockl_get_local_id(0);
+      b[tx] = a[tx];
+    }
+)"};
+
+bool TestAddLinkData(char* bit_code_ptr, size_t bit_code_size,
+                     hiprtcJITInputType input_type, void** cuOut, size_t* out_size) {
+
+    hiprtcLinkState hiprtc_link_state;
+    hiprtcJIT_option options[6];
+    void* option_vals[6];
+    float wall_time;
+    size_t log_size = 8192;
+    char error_log[8192];
+    char info_log[8192];
+
+    options[0] = HIPRTC_JIT_WALL_TIME;
+    option_vals[0] = (void*)(&wall_time);
+
+    options[1] = HIPRTC_JIT_INFO_LOG_BUFFER;
+    option_vals[1] = (void*)info_log;
+
+    options[2] = HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
+    option_vals[2] = (void*)(log_size);
+
+    options[3] = HIPRTC_JIT_ERROR_LOG_BUFFER;
+    option_vals[3] = (void*)error_log;
+
+    options[4] = HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;
+    option_vals[4] = (void*)(log_size);
+
+    options[5] = HIPRTC_JIT_LOG_VERBOSE;
+    option_vals[5] = (void*)1;
+
+    if (HIPRTC_SUCCESS != hiprtcLinkCreate(6, options, option_vals, &hiprtc_link_state)) {
+      std::cout<<"hipLinkCreate failed "<<std::endl;
+      return false;
+    }
+
+    std::string link_data_name("link_data");
+    if (HIPRTC_SUCCESS != hiprtcLinkAddData(hiprtc_link_state, input_type, bit_code_ptr,
+                                            bit_code_size, link_data_name.c_str(), 0, 0, 0)) {
+      std::cout<<"hipLinkAddFile failed "<<std::endl;
+      return false;
+    }
+
+    if (HIPRTC_SUCCESS != hiprtcLinkComplete(hiprtc_link_state, cuOut, out_size)) {
+      std::cout<<"hipLinkComplete failed "<<std::endl;
+      return false;
+    }
+
+    if (HIPRTC_SUCCESS != hiprtcLinkDestroy(hiprtc_link_state)) {
+      std::cout<<"hipLinkDestroy failed "<<std::endl;
+      return false;
+    }
+
+    return true;
+}
+
+TEST_CASE("Unit_hiprtcLinkData_bitcode") {
+  // Create a bit code via RTC Compile APIs with fgpu-rdc options.
+  char* bit_code = nullptr;
+  size_t bit_code_size = 0;
+  REQUIRE(TestCompileRDC(&bit_code, &bit_code_size, vcpy, "VectorCopy"));
+
+  // Test the link file API
+  void* executable = nullptr;
+  size_t exe_size = 0;
+  REQUIRE(TestAddLinkData(bit_code, bit_code_size, HIPRTC_JIT_INPUT_LLVM_BITCODE,
+                          &executable, &exe_size));
+
+  // Test the executable generated by the API.
+  REQUIRE(TestModuleLoadData(executable));
+}
+
diff --git a/catch/unit/rtc/hiprtcLinkFile.cc b/catch/unit/rtc/hiprtcLinkFile.cc
new file mode 100644
index 000000000..490a4b93c
--- /dev/null
+++ b/catch/unit/rtc/hiprtcLinkFile.cc
@@ -0,0 +1,111 @@
+/*
+ * Copyright (c) 2022 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 WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ * */
+
+#include <hip_test_common.hh>
+
+#include <iostream>
+#include <fstream>
+#include <string>
+
+#include "hiprtcHelper.hpp"
+#include <hip/hiprtc.h>
+#include <hip/hip_runtime.h>
+
+static constexpr auto vcpy{
+  R"(
+    extern "C" __global__ void vcpy_kernel(float* a, float* b) {
+      int tx = __ockl_get_local_id(0);
+      b[tx] = a[tx];
+    }
+)"};
+
+bool TestAddLinkFile(std::string bc_file_path, hiprtcJITInputType input_type, void** cuOut,
+                                                                           size_t* cuSize) {
+    hiprtcLinkState hiprtc_link_state;
+    hiprtcJIT_option options[6];
+    void* option_vals[6];
+    float wall_time;
+    size_t log_size = 8192;
+    char error_log[8192];
+    char info_log[8192];
+
+    options[0] = HIPRTC_JIT_WALL_TIME;
+    option_vals[0] = (void*)(&wall_time);
+
+    options[1] = HIPRTC_JIT_INFO_LOG_BUFFER;
+    option_vals[1] = (void*)info_log;
+
+    options[2] = HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
+    option_vals[2] = (void*)(log_size);
+
+    options[3] = HIPRTC_JIT_ERROR_LOG_BUFFER;
+    option_vals[3] = (void*)error_log;
+
+    options[4] = HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;
+    option_vals[4] = (void*)(log_size);
+
+    options[5] = HIPRTC_JIT_LOG_VERBOSE;
+    option_vals[5] = (void*)1;
+
+    if (HIPRTC_SUCCESS != hiprtcLinkCreate(6, options, option_vals, &hiprtc_link_state)) {
+      std::cout<<"hipLinkCreate failed "<<std::endl;
+      return false;
+    }
+
+    std::string link_data_name("link_data");
+    if (HIPRTC_SUCCESS != hiprtcLinkAddFile(hiprtc_link_state, input_type, bc_file_path.c_str(),
+                                             0, 0, 0)) {
+      std::cout<<"hipLinkAddFile failed "<<std::endl;
+      return false;
+    }
+
+    if (HIPRTC_SUCCESS != hiprtcLinkComplete(hiprtc_link_state, cuOut, cuSize)) {
+      std::cout<<"hipLinkComplete failed "<<std::endl;
+      return false;
+    }
+
+    if (HIPRTC_SUCCESS != hiprtcLinkDestroy(hiprtc_link_state)) {
+      std::cout<<"hipLinkDestroy failed "<<std::endl;
+      return false;
+    }
+
+    return true;
+}
+
+TEST_CASE("Unit_hiprtcLinkFile_bitcode") {
+  // Create a bit code via RTC Compile APIs with fgpu-rdc options.
+  char* bit_code = nullptr;
+  size_t bit_code_size = 0;
+  REQUIRE(TestCompileRDC(&bit_code, &bit_code_size, vcpy, "VectorCopy.hip"));
+
+  // Write that Bit code to File.
+  std::string bc_file_path("vcpy_kernel.bc");
+  REQUIRE(CommitBCToFile(bit_code, bit_code_size, bc_file_path));
+
+  // Test the link file API
+  void* executable = nullptr;
+  size_t exe_size = 0;
+  REQUIRE(TestAddLinkFile(bc_file_path, HIPRTC_JIT_INPUT_LLVM_BITCODE, &executable, &exe_size));
+
+  // Test the executable generated by the API.
+  REQUIRE(TestModuleLoadData(executable));
+}