diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt
index 8674ffdc10..7f5bdad978 100755
--- a/hipamd/CMakeLists.txt
+++ b/hipamd/CMakeLists.txt
@@ -8,10 +8,15 @@ set(BUILD_SHARED_LIBS ON  CACHE BOOL "Build shared library (.so) or static lib (
 
 set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
 
-if(NOT ${BUILD_SHARED_LIBS} AND NOT DEFINED ENABLE_HIP_PCH)
-  set(ENABLE_HIP_PCH ON CACHE BOOL "enable/disable pre-compiled hip headers")
+if(NOT DEFINED __HIP_ENABLE_PCH)
+  set(__HIP_ENABLE_PCH ON CACHE BOOL "enable/disable pre-compiled hip headers")
 endif()
 
+if(${__HIP_ENABLE_PCH})
+  set(_pchStatus 1)
+else()
+  set(_pchStatus 0)
+endif()
 #############################
 # Options
 #############################
@@ -455,6 +460,7 @@ set(_versionInfoHeader
 #define HIP_VERSION_MINOR ${HIP_VERSION_MINOR}
 #define HIP_VERSION_PATCH ${HIP_VERSION_GITDATE}
 #define HIP_VERSION       (HIP_VERSION_MAJOR * 100 + HIP_VERSION_MINOR)\n
+#define __HIP_HAS_GET_PCH ${_pchStatus}\n
 #endif\n
 ")
 file(WRITE "${PROJECT_BINARY_DIR}/include/hip/hip_version.h" ${_versionInfoHeader})
diff --git a/hipamd/bin/hip_embed_pch.sh b/hipamd/bin/hip_embed_pch.sh
index 8fe3c20f98..0f2cbabd84 100755
--- a/hipamd/bin/hip_embed_pch.sh
+++ b/hipamd/bin/hip_embed_pch.sh
@@ -1,8 +1,7 @@
 #!/bin/bash
 
 #set -x
-
-ROCM_PATH=${ROCM_PATH:-/opt/rocm}
+LLVM_DIR="$1/../../../"
 tmp=/tmp/hip_pch.$$
 mkdir -p $tmp
 
@@ -47,12 +46,12 @@ __hip_pch_size:
   .long __hip_pch_size - __hip_pch
 EOF
 
-$ROCM_PATH/llvm/bin/clang -O3 -c -std=c++17 -isystem /opt/rocm/llvm/lib/clang/11.0.0/include/.. -isystem /opt/rocm/include -nogpulib --cuda-device-only -x hip $tmp/hip_pch.h -E >$tmp/pch.cui
+$LLVM_DIR/bin/clang -O3 -c -std=c++17 -isystem $LLVM_DIR/lib/clang/11.0.0/include/.. -isystem /opt/rocm/include -nogpulib --cuda-device-only -x hip $tmp/hip_pch.h -E >$tmp/pch.cui
 
 cat $tmp/hip_macros.h >> $tmp/pch.cui
 
-$ROCM_PATH/llvm/bin/clang -cc1 -O3 -emit-pch -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -std=c++17 -fgnuc-version=4.2.1 -o $tmp/hip.pch -x hip-cpp-output - <$tmp/pch.cui
+$LLVM_DIR/bin/clang -cc1 -O3 -emit-pch -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -std=c++17 -fgnuc-version=4.2.1 -o $tmp/hip.pch -x hip-cpp-output - <$tmp/pch.cui
 
-$ROCM_PATH/llvm/bin/llvm-mc -o hip_pch.o $tmp/hip_pch.mcin --filetype=obj
+$LLVM_DIR/bin/llvm-mc -o hip_pch.o $tmp/hip_pch.mcin --filetype=obj
 
 rm -rf $tmp
diff --git a/hipamd/bin/hip_gen_pch.sh b/hipamd/bin/hip_gen_pch.sh
deleted file mode 100755
index b212177119..0000000000
--- a/hipamd/bin/hip_gen_pch.sh
+++ /dev/null
@@ -1,36 +0,0 @@
-#!/bin/bash
-
-#set -x
-
-cat >/tmp/hip_macros.h <<EOF
-#define __device__ __attribute__((device))
-#define __host__ __attribute__((host))
-#define __global__ __attribute__((global))
-#define __constant__ __attribute__((constant))
-#define __shared__ __attribute__((shared))
-
-#define launch_bounds_impl0(requiredMaxThreadsPerBlock)                                            \
-    __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock)))
-#define launch_bounds_impl1(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)                \
-    __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock),                     \
-                   amdgpu_waves_per_eu(minBlocksPerMultiprocessor)))
-#define select_impl_(_1, _2, impl_, ...) impl_
-#define __launch_bounds__(...)                                                                     \
-    select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)
-
-// Macro to replace extern __shared__ declarations
-// to local variable definitions
-#define HIP_DYNAMIC_SHARED(type, var) \
-    type* var = (type*)__amdgcn_get_dynamicgroupbaseptr();
-EOF
-
-cat >/tmp/hip_pch.h <<EOF
-#include "hip/hip_runtime.h"
-#include "hip/hip_fp16.h"
-EOF
-
-/opt/rocm/llvm/bin/clang -O3 -c -std=c++17 -isystem /opt/rocm/llvm/lib/clang/11.0.0/include/.. -isystem /opt/rocm/include -nogpulib --cuda-device-only -x hip /tmp/hip_pch.h -E >/tmp/pch.cui
-
-cat /tmp/hip_macros.h >> /tmp/pch.cui
-
-/opt/rocm/llvm/bin/clang -cc1 -O3 -emit-pch -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -std=c++17 -fgnuc-version=4.2.1 -o /tmp/hip.pch -x hip-cpp-output - </tmp/pch.cui
diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h
index 37fcccf192..74c0fb5f69 100755
--- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h
+++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h
@@ -345,13 +345,16 @@ typedef struct hipLaunchParams_t {
     hipStream_t stream;     ///< Stream identifier
 } hipLaunchParams;
 
-// Pre-Compiled header for online compilation
-#ifdef ENABLE_HIP_PCH
-extern const char* __hip_pch;
-extern unsigned __hip_pch_size;
-void __hipGetPCH(const char** pch, unsigned int*size);
+#if __HIP_HAS_GET_PCH
+/**
+ * Internal use only. This API may change in the future
+ * Pre-Compiled header for online compilation
+ *
+ */
+    void __hipGetPCH(const char** pch, unsigned int*size);
 #endif
 
+
 // Doxygen end group GlobalDefs
 /**  @} */
 
diff --git a/hipamd/include/hip/hip_runtime_api.h b/hipamd/include/hip/hip_runtime_api.h
index 4412bbd7da..3a26fb74f4 100644
--- a/hipamd/include/hip/hip_runtime_api.h
+++ b/hipamd/include/hip/hip_runtime_api.h
@@ -32,6 +32,7 @@ THE SOFTWARE.
 
 
 #include <string.h>  // for getDeviceProp
+#include <hip/hip_version.h>
 #include <hip/hip_common.h>
 
 enum {
diff --git a/hipamd/rocclr/CMakeLists.txt b/hipamd/rocclr/CMakeLists.txt
index 7f65666e2c..ec1dc50407 100755
--- a/hipamd/rocclr/CMakeLists.txt
+++ b/hipamd/rocclr/CMakeLists.txt
@@ -96,6 +96,14 @@ find_package(amd_comgr REQUIRED CONFIG
 
 message(STATUS "Code Object Manager found at ${amd_comgr_DIR}.")
 
+find_package(LLVM REQUIRED CONFIG
+   PATHS
+     /opt/rocm/llvm
+   PATH_SUFFIXES
+     lib/cmake/llvm)
+
+message(STATUS "llvm found at ${LLVM_DIR}.")
+
 add_library(hip64 OBJECT
  hip_context.cpp
  hip_code_object.cpp
@@ -148,10 +156,9 @@ endif()
 
 # Short-Term solution for pre-compiled headers for online compilation
 # Enable pre compiled header
-if(${ENABLE_HIP_PCH})
-    execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_gen_pch.sh")
-    execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh")
-    add_definitions(-DENABLE_HIP_PCH)
+if(${__HIP_ENABLE_PCH})
+    execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh ${LLVM_DIR}")
+    add_definitions(-D__HIP_ENABLE_PCH)
 endif()
 
 # Enable profiling API
@@ -216,7 +223,7 @@ add_library(device INTERFACE)
 target_link_libraries(device INTERFACE host)
 
 # Short-Term solution for pre-compiled headers for online compilation
-if(${ENABLE_HIP_PCH})
+if(${__HIP_ENABLE_PCH})
   target_link_libraries(amdhip64 PRIVATE ${CMAKE_BINARY_DIR}/hip_pch.o)
 endif()
 
diff --git a/hipamd/rocclr/hip_global.cpp b/hipamd/rocclr/hip_global.cpp
index bed2dcd850..46e6efcf52 100755
--- a/hipamd/rocclr/hip_global.cpp
+++ b/hipamd/rocclr/hip_global.cpp
@@ -5,7 +5,9 @@
 #include "hip_code_object.hpp"
 #include "platform/program.hpp"
 
-#ifdef ENABLE_HIP_PCH
+#ifdef __HIP_ENABLE_PCH
+extern const char __hip_pch[];
+extern unsigned __hip_pch_size;
 void __hipGetPCH(const char** pch, unsigned int *size) {
   *pch = __hip_pch;
   *size = __hip_pch_size;