SWDEV-308763 - Make hiprtc a separate lib and rewrite it using comgr

Change-Id: I28ac50ec897accb2a2b4590d3c25965d907426fb

hipamd/src/hiprtc/cmake/HIPRTC.cmake

@@ -0,0 +1,93 @@
+# Copyright (c) 2021 - 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.
+
+
+###############################################################################
+# HIPRTC.cmake
+###############################################################################
+
+# This file includes macros required to generate the hiprtc builtins library.
+
+function(get_hiprtc_macros HIPRTC_DEFINES)
+  set(${HIPRTC_DEFINES}
+"#pragma clang diagnostic push\n\
+#pragma clang diagnostic ignored \"-Wreserved-id-macro\"\n\
+#define __device__ __attribute__((device))\n\
+#define __host__ __attribute__((host))\n\
+#define __global__ __attribute__((global))\n\
+#define __constant__ __attribute__((constant))\n\
+#define __shared__ __attribute__((shared))\n\
+#define __align__(x) __attribute__((aligned(x)))\n\
+#define __noinline__ __attribute__((noinline))\n\
+#define __forceinline__ inline __attribute__((always_inline))\n\
+
+#define launch_bounds_impl0(requiredMaxThreadsPerBlock)                                       \\\n\
+    __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock)))\n\
+#define launch_bounds_impl1(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)           \\\n\
+    __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock),                \\\n\
+                   amdgpu_waves_per_eu(minBlocksPerMultiprocessor)))\n\
+#define select_impl_(_1, _2, impl_, ...) impl_\n\
+#define __launch_bounds__(...)                                                                \\\n\
+    select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)           \n\
+#pragma clang diagnostic pop\n\
+#define HIP_INCLUDE_HIP_HIP_RUNTIME_H"
+  PARENT_SCOPE)
+endfunction(get_hiprtc_macros)
+
+# To allow concatenating above macros during build time, call this file in script mode.
+if(HIPRTC_ADD_MACROS)
+  message(STATUS "Appending hiprtc macros to ${HIPRTC_PREPROCESSED_FILE}.")
+  get_hiprtc_macros(HIPRTC_DEFINES)
+  FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_DEFINES}")
+endif()
+
+macro(generate_hiprtc_header HiprtcHeader)
+  FILE(WRITE ${HiprtcHeader}
+"#pragma push_macro(\"CHAR_BIT\")\n\
+#pragma push_macro(\"INT_MAX\")\n\
+#define CHAR_BIT __CHAR_BIT__\n\
+#define INT_MAX __INTMAX_MAX__\n\
+#include \"hip/hip_runtime.h\"\n\
+#include \"hip/hip_fp16.h\"\n\
+#pragma pop_macro(\"CHAR_BIT\")\n\
+#pragma pop_macro(\"INT_MAX\")")
+endmacro(generate_hiprtc_header)
+
+macro(generate_hiprtc_mcin HiprtcMcin HiprtcPreprocessedInput)
+  if(WIN32)
+    set(HIPRTC_TYPE_LINUX_ONLY "")
+  else()
+    set(HIPRTC_TYPE_LINUX_ONLY
+      "  .type __hipRTC_header,@object\n"
+      "  .type __hipRTC_header_size,@object")
+  endif()
+  FILE(WRITE ${HiprtcMcin}
+"// Automatically generated script for HIPRTC.\n\
+${HIPRTC_TYPE_LINUX_ONLY}\n\
+  .section .hipRTC_header,\"a\"\n\
+  .globl __hipRTC_header\n\
+  .globl __hipRTC_header_size\n\
+  .p2align 3\n\
+__hipRTC_header:\n\
+  .incbin \"${HiprtcPreprocessedInput}\"\n\
+__hipRTC_header_size:\n\
+  .long __hipRTC_header_size - __hipRTC_header\n")
+endmacro(generate_hiprtc_mcin)
+