diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt
index 4e5639406f..0c4b928756 100755
--- a/hipamd/CMakeLists.txt
+++ b/hipamd/CMakeLists.txt
@@ -171,28 +171,80 @@ endif()
 
 # Enable preprocessed hiprtc-builtins library
 if(__HIP_ENABLE_RTC)
-  find_package(LLVM REQUIRED CONFIG
-    PATHS
-      /opt/rocm/llvm)
-  # find_package(LLVM) returns the lib/cmake/llvm location. We require the root.
-  if(NOT DEFINED HIP_LLVM_ROOT)
-    set(HIP_LLVM_ROOT "${LLVM_DIR}/../../..")
+  message(STATUS "HIP RTC enabled.")
+  include(HIPRTC RESULT_VARIABLE HIPRTC_CMAKE)
+  # Requires clang and llvm-mc to create this library.
+  find_package(LLVM REQUIRED CONFIG PATHS /opt/rocm/llvm)
+  find_package(Clang REQUIRED CONFIG PATHS /opt/rocm/llvm)
+  set(HIPRTC_GEN_DIR "${CMAKE_CURRENT_BINARY_DIR}/hip_rtc_gen")
+  set(HIPRTC_GEN_HEADER "${HIPRTC_GEN_DIR}/hipRTC_header.h")
+  set(HIPRTC_GEN_MCIN "${HIPRTC_GEN_DIR}/hipRTC_header.mcin")
+  set(HIPRTC_GEN_PREPROCESSED "${HIPRTC_GEN_DIR}/hipRTC")
+  set(HIPRTC_GEN_OBJ "${HIPRTC_GEN_DIR}/hipRTC_header${CMAKE_CXX_OUTPUT_EXTENSION}")
+
+  # Generate required HIPRTC files.
+  FILE(MAKE_DIRECTORY ${HIPRTC_GEN_DIR})
+  generate_hiprtc_header("${HIPRTC_GEN_HEADER}")
+  generate_hiprtc_mcin("${HIPRTC_GEN_MCIN}" "${HIPRTC_GEN_PREPROCESSED}")
+
+  # Generate HIPRTC Builtins Preprocessed Object.
+  # Note: second command appends define macros at build time.
+  # FIXME: --hip-version forced to 3.6 to use clang headers, until Windows versioning is fixed.
+  add_custom_command(
+    OUTPUT ${HIPRTC_GEN_PREPROCESSED}
+    COMMAND $<TARGET_FILE:clang> -O3 --rocm-path=${PROJECT_SOURCE_DIR}/include/.. -std=c++14 -nogpulib --hip-version=3.6 -isystem ${PROJECT_SOURCE_DIR}/include -isystem ${PROJECT_BINARY_DIR}/include -isystem ${CMAKE_CURRENT_SOURCE_DIR}/include --cuda-device-only -D__HIPCC_RTC__ -x hip ${HIPRTC_GEN_HEADER} -E -o ${HIPRTC_GEN_PREPROCESSED}
+    COMMAND ${CMAKE_COMMAND} -DHIPRTC_ADD_MACROS=1 -DHIPRTC_PREPROCESSED_FILE=${HIPRTC_GEN_PREPROCESSED} -P ${HIPRTC_CMAKE}
+    DEPENDS clang ${HIPRTC_GEN_HEADER})
+  add_custom_command(
+    OUTPUT ${HIPRTC_GEN_OBJ}
+    COMMAND $<TARGET_FILE:llvm-mc> -o ${HIPRTC_GEN_OBJ} ${HIPRTC_GEN_MCIN} --filetype=obj
+    DEPENDS llvm-mc ${HIPRTC_GEN_PREPROCESSED} ${HIPRTC_GEN_MCIN})
+
+  # Create hiprtc-builtins library.
+  add_library(hiprtc-builtins ${HIPRTC_GEN_OBJ})
+  set_target_properties(hiprtc-builtins PROPERTIES
+    CXX_STANDARD 14
+    CXX_STANDARD_REQUIRED ON
+    CXX_EXTENSIONS OFF
+    POSITION_INDEPENDENT_CODE ON
+    LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib
+    LINKER_LANGUAGE CXX
+    VERSION ${HIP_LIB_VERSION_STRING})
+
+  # Windows and Linux have different naming conventions.
+  if(WIN32)
+    # Windows uses DEF file to determine which symbols to expose.
+    target_sources(hiprtc-builtins PRIVATE src/hiprtc-builtins.def)
+    set_target_properties(hiprtc-builtins PROPERTIES
+      OUTPUT_NAME "hiprtc-builtins64_${HIP_LIB_VERSION_MAJOR}${HIP_LIB_VERSION_MINOR}")
+    # Since ${HIPRTC_GEN_OBJ} was manually generated with llvm-mc, /MT did not embed
+    # libcmt.lib inside of the obj. So we need to manually set it as defaultlib.
+    target_link_options(hiprtc-builtins PRIVATE "LINKER:/DEFAULTLIB:libcmt")
+  else()
+    # SOVERSION is only supported on Linux.
+    set_target_properties(hiprtc-builtins PROPERTIES
+      OUTPUT_NAME "hiprtc-builtins"
+      SOVERSION ${HIP_LIB_VERSION_MAJOR})
   endif()
 
+  # Test the header file works with simple compilation.
+  add_custom_command(
+    OUTPUT ${HIPRTC_GEN_DIR}/tmp.bc
+    COMMAND $<TARGET_FILE:clang> -O3 --rocm-path=${PROJECT_SOURCE_DIR}/include/.. -std=c++14 -nogpulib -nogpuinc -emit-llvm -c -isystem ${PROJECT_SOURCE_DIR}/include -isystem ${PROJECT_BINARY_DIR}/include -isystem ${CMAKE_CURRENT_SOURCE_DIR}/include --cuda-device-only -D__HIPCC_RTC__ --offload-arch=gfx906 -x hip-cpp-output ${HIPRTC_GEN_PREPROCESSED} -o ${HIPRTC_GEN_DIR}/tmp.bc
+    DEPENDS clang ${HIPRTC_GEN_PREPROCESSED})
+
+  # FIXME: As a workaround, add hiprtc object into amdhip64, until we can
+  # figure out how to link hiprtc-builtins into amdhip64. CMake approach is not working:
+  # target_link_libraries(amdhip64 PUBLIC hiprtc-builtins)
+  target_link_libraries(amdhip64 PRIVATE ${HIPRTC_GEN_OBJ})
+  target_compile_definitions(amdhip64 PRIVATE __HIP_ENABLE_RTC)
+  add_dependencies(amdhip64 hiprtc-builtins)
+  # Windows will install the library to the bin directory instead.
   if(WIN32)
-    set(HIPRTC_LIB_NAME "hiprtc-builtins64_${HIP_LIB_VERSION_MAJOR}${HIP_LIB_VERSION_MINOR}.dll")
+    install(TARGETS hiprtc-builtins DESTINATION ${CMAKE_INSTALL_BINDIR})
   else()
-    set(HIPRTC_LIB_NAME "libhiprtc-builtins.so.${HIP_LIB_VERSION_MAJOR}.${HIP_LIB_VERSION_MINOR}")
+    install(TARGETS hiprtc-builtins DESTINATION lib)
   endif()
-  execute_process(
-	  COMMAND sh -c "mkdir -p ${PROJECT_BINARY_DIR}/lib; ${HIP_COMMON_DIR}/bin/hip_embed_pch.sh ${PROJECT_BINARY_DIR}/include ${PROJECT_SOURCE_DIR}/include ${CMAKE_CURRENT_SOURCE_DIR}/include ${HIP_LLVM_ROOT} -r ${PROJECT_BINARY_DIR}/lib/${HIPRTC_LIB_NAME}"
-    COMMAND_ECHO STDERR
-    RESULT_VARIABLE EMBED_RTC_RC
-  )
-  if (EMBED_RTC_RC AND NOT EMBED_RTC_RC EQUAL 0)
-    message(FATAL_ERROR "Failed to create hiprtc shared lib")
-  endif()
-  install(FILES ${PROJECT_BINARY_DIR}/lib/${HIPRTC_LIB_NAME} DESTINATION lib)
 endif()
 
 #############################
diff --git a/hipamd/cmake/HIPRTC.cmake b/hipamd/cmake/HIPRTC.cmake
new file mode 100644
index 0000000000..2b16d08fed
--- /dev/null
+++ b/hipamd/cmake/HIPRTC.cmake
@@ -0,0 +1,86 @@
+# Copyright (C) 2021-2021 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}
+"#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 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__)"
+  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)
+
diff --git a/hipamd/src/hiprtc-builtins.def b/hipamd/src/hiprtc-builtins.def
new file mode 100644
index 0000000000..b878164125
--- /dev/null
+++ b/hipamd/src/hiprtc-builtins.def
@@ -0,0 +1,4 @@
+EXPORTS
+__hipRTC_header
+__hipRTC_header_size
+