diff --git a/CMakeLists.txt b/CMakeLists.txt
index 537764548f..a4da3b1920 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -193,9 +193,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
         src/program_state.cpp)
 
     set(SOURCE_FILES_DEVICE
-        src/device_util.cpp
-        src/hip_ldg.cpp
-        src/device_functions.cpp)
+        src/device_util.cpp)
 
     execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS)
     set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic")
@@ -238,7 +236,6 @@ endif()
 # Install hip_hcc if platform is hcc
 if(HIP_PLATFORM STREQUAL "hcc")
     install(TARGETS hip_hcc_static hip_hcc hip_device DESTINATION lib)
-    install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_hc.ll DESTINATION lib)
 
     # Install .hipInfo
     install(FILES ${PROJECT_BINARY_DIR}/.hipInfo DESTINATION lib)
diff --git a/bin/hipcc b/bin/hipcc
index 0abf09ed71..d2614297c3 100755
--- a/bin/hipcc
+++ b/bin/hipcc
@@ -259,8 +259,7 @@ if($HIP_PLATFORM eq "hcc"){
 }
 
 if(($HIP_PLATFORM eq "hcc")){
-    $ENV{HCC_EXTRA_LIBRARIES}="$HIP_PATH/lib/hip_hc.ll\n";
-    $ENV{HIP_HC_IR_FILE}="";
+    $ENV{HCC_EXTRA_LIBRARIES}="\n";
 }
 
 if($HIP_PLATFORM eq "nvcc"){
@@ -533,7 +532,7 @@ if($HIP_PLATFORM eq "hcc" or $HIP_PLATFORM eq "clang"){
         print "No valid AMD GPU target was either specified or found. Please specify a valid target using --amdgpu-target=" and die();
     }
 
-    $ENV{HCC_EXTRA_LIBRARIES}="$HIP_PATH/lib/hip_hc.ll\n";
+    $ENV{HCC_EXTRA_LIBRARIES}="\n";
 
     if($HIP_PLATFORM eq "hcc") {
         $GPU_ARCH_OPT = " --amdgpu-target=";
diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h
index ae1b96c979..fb453d02e9 100644
--- a/include/hip/hcc_detail/device_functions.h
+++ b/include/hip/hcc_detail/device_functions.h
@@ -24,46 +24,171 @@ THE SOFTWARE.
 #define HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
 
 #include "host_defines.h"
+#include "math_fwd.h"
 
+#include <hip/hip_runtime_api.h>
 #include <hip/hip_vector_types.h>
+#include <hip/hcc_detail/device_library_decls.h>
+#include <hip/hcc_detail/llvm_intrinsics.h>
 
-extern "C" __device__ unsigned int __hip_hc_ir_umul24_int(unsigned int, unsigned int);
-extern "C" __device__ signed int __hip_hc_ir_mul24_int(signed int, signed int);
-extern "C" __device__ signed int __hip_hc_ir_mulhi_int(signed int, signed int);
-extern "C" __device__ unsigned int __hip_hc_ir_umulhi_int(unsigned int, unsigned int);
-extern "C" __device__ unsigned int __hip_hc_ir_usad_int(unsigned int, unsigned int, unsigned int);
+typedef unsigned long ulong;
+typedef unsigned int uint;
+
+/*
+Integer Intrinsics
+*/
 
 // integer intrinsic function __poc __clz __ffs __brev
-__device__ unsigned int __brev(unsigned int x);
-__device__ unsigned long long int __brevll(unsigned long long int x);
-__device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s);
-__device__ unsigned int __clz(int x);
-__device__ unsigned int __clzll(long long int x);
-__device__ unsigned int __ffs(int x);
-__device__ unsigned int __ffsll(long long int x);
+__device__ static inline unsigned int __popc(unsigned int input) {
+    return __builtin_popcount(input);
+}
+__device__ static inline unsigned int __popcll(unsigned long long int input) {
+    return __builtin_popcountl(input);
+}
+
+__device__ static inline unsigned int __clz(unsigned int input) {
+#ifdef NVCC_COMPAT
+    return input == 0 ? 32 : __builtin_clz(input);
+#else
+    return input == 0 ? -1 : __builtin_clz(input);
+#endif
+}
+
+__device__ static inline unsigned int __clzll(unsigned long long int input) {
+#ifdef NVCC_COMPAT
+    return input == 0 ? 64 : ( input == 0 ? -1 : __builtin_clzl(input) );
+#else
+    return input == 0 ? -1 : __builtin_clzl(input);
+#endif
+}
+
+__device__ static inline unsigned int __clz(int input) {
+#ifdef NVCC_COMPAT
+    return input == 0 ? 32 : ( input > 0 ? __builtin_clz(input) : __builtin_clz(~input) );
+#else
+    if (input == 0) return -1;
+    return input > 0 ? __builtin_clz(input) : __builtin_clz(~input);
+#endif
+}
+
+__device__ static inline unsigned int __clzll(long long int input) {
+#ifdef NVCC_COMPAT
+    return input == 0 ? 64 : input > 0 ? __builtin_clzl(input) : __builtin_clzl(~input);
+#else
+    if (input == 0) return -1;
+    return input > 0 ? __builtin_clzl(input) : __builtin_clzl(~input);
+#endif
+}
+
+__device__ static inline unsigned int __ffs(unsigned int input) {
+#ifdef NVCC_COMPAT
+    return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1;
+#else
+    return input == 0 ? -1 : __builtin_ctz(input);
+#endif
+}
+
+__device__ static inline unsigned int __ffsll(unsigned long long int input) {
+#ifdef NVCC_COMPAT
+    return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1;
+#else
+    return input == 0 ? -1 : __builtin_ctzl(input);
+#endif
+}
+
+__device__ static inline unsigned int __ffs(int input) {
+#ifdef NVCC_COMPAT
+    return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1;
+#else
+    return input == 0 ? -1 : __builtin_ctz(input);
+#endif
+}
+
+__device__ static inline unsigned int __ffsll(long long int input) {
+#ifdef NVCC_COMPAT
+    return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1;
+#else
+    return input == 0 ? -1 : __builtin_ctzl(input);
+#endif
+}
+
+__device__ static inline unsigned int __brev(unsigned int input) { return __llvm_bitrev_b32(input); }
+
+__device__ static inline unsigned long long int __brevll(unsigned long long int input) {
+    return __llvm_bitrev_b64(input);
+}
+
+__device__ static unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s);
 __device__ static unsigned int __hadd(int x, int y);
 __device__ static int __mul24(int x, int y);
-__device__ long long int __mul64hi(long long int x, long long int y);
+__device__ static long long int __mul64hi(long long int x, long long int y);
 __device__ static int __mulhi(int x, int y);
-__device__ unsigned int __popc(unsigned int x);
-__device__ unsigned int __popcll(unsigned long long int x);
 __device__ static int __rhadd(int x, int y);
 __device__ static unsigned int __sad(int x, int y, int z);
 __device__ static unsigned int __uhadd(unsigned int x, unsigned int y);
 __device__ static int __umul24(unsigned int x, unsigned int y);
-__device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y);
+__device__ static unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y);
 __device__ static unsigned int __umulhi(unsigned int x, unsigned int y);
 __device__ static unsigned int __urhadd(unsigned int x, unsigned int y);
 __device__ static unsigned int __usad(unsigned int x, unsigned int y, unsigned int z);
 
+struct ucharHolder {
+    union {
+        unsigned char c[4];
+        unsigned int ui;
+    };
+} __attribute__((aligned(4)));
+
+struct uchar2Holder {
+    union {
+        unsigned int ui[2];
+        unsigned char c[8];
+    };
+} __attribute__((aligned(8)));
+
+__device__
+static inline unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s) {
+    struct uchar2Holder cHoldVal;
+    struct ucharHolder cHoldKey;
+    struct ucharHolder cHoldOut;
+    cHoldKey.ui = s;
+    cHoldVal.ui[0] = x;
+    cHoldVal.ui[1] = y;
+    cHoldOut.c[0] = cHoldVal.c[cHoldKey.c[0]];
+    cHoldOut.c[1] = cHoldVal.c[cHoldKey.c[1]];
+    cHoldOut.c[2] = cHoldVal.c[cHoldKey.c[2]];
+    cHoldOut.c[3] = cHoldVal.c[cHoldKey.c[3]];
+    return cHoldOut.ui;
+}
+
 __device__ static inline unsigned int __hadd(int x, int y) {
     int z = x + y;
     int sign = z & 0x8000000;
     int value = z & 0x7FFFFFFF;
     return ((value) >> 1 || sign);
 }
-__device__ static inline int __mul24(int x, int y) { return __hip_hc_ir_mul24_int(x, y); }
-__device__ static inline int __mulhi(int x, int y) { return __hip_hc_ir_mulhi_int(x, y); }
+
+__device__ static inline int __mul24(int x, int y) {
+    return __ockl_mul24_i32(x, y);
+}
+
+__device__ static inline long long __mul64hi(long long int x, long long int y) {
+    ulong x0 = (ulong)x & 0xffffffffUL;
+    long x1 = x >> 32;
+    ulong y0 = (ulong)y & 0xffffffffUL;
+    long y1 = y >> 32;
+    ulong z0 = x0*y0;
+    long t = x1*y0 + (z0 >> 32);
+    long z1 = t & 0xffffffffL;
+    long z2 = t >> 32;
+    z1 = x0*y1 + z1;
+    return x1*y1 + z2 + (z1 >> 32);
+}
+
+__device__ static inline int __mulhi(int x, int y) {
+    return __ockl_mul_hi_i32(x, y);
+}
+
 __device__ static inline int __rhadd(int x, int y) {
     int z = x + y + 1;
     int sign = z & 0x8000000;
@@ -77,57 +202,307 @@ __device__ static inline unsigned int __uhadd(unsigned int x, unsigned int y) {
     return (x + y) >> 1;
 }
 __device__ static inline int __umul24(unsigned int x, unsigned int y) {
-    return __hip_hc_ir_umul24_int(x, y);
+    return __ockl_mul24_u32(x, y);
 }
+
+__device__
+static inline unsigned long long __umul64hi(unsigned long long int x, unsigned long long int y) {
+    ulong x0 = x & 0xffffffffUL;
+    ulong x1 = x >> 32;
+    ulong y0 = y & 0xffffffffUL;
+    ulong y1 = y >> 32;
+    ulong z0 = x0*y0;
+    ulong t = x1*y0 + (z0 >> 32);
+    ulong z1 = t & 0xffffffffUL;
+    ulong z2 = t >> 32;
+    z1 = x0*y1 + z1;
+    return x1*y1 + z2 + (z1 >> 32);
+}
+
 __device__ static inline unsigned int __umulhi(unsigned int x, unsigned int y) {
-    return __hip_hc_ir_umulhi_int(x, y);
+    return __ockl_mul_hi_u32(x, y);
 }
 __device__ static inline unsigned int __urhadd(unsigned int x, unsigned int y) {
     return (x + y + 1) >> 1;
 }
 __device__ static inline unsigned int __usad(unsigned int x, unsigned int y, unsigned int z) {
-    return __hip_hc_ir_usad_int(x, y, z);
+    return __ockl_sad_u32(x, y, z);
 }
 
-extern __device__ __attribute__((const)) unsigned int __mbcnt_lo(unsigned int x, unsigned int y) __asm("llvm.amdgcn.mbcnt.lo");
-extern __device__ __attribute__((const)) unsigned int __mbcnt_hi(unsigned int x, unsigned int y) __asm("llvm.amdgcn.mbcnt.hi");
-
 __device__ static inline unsigned int __lane_id() { return  __mbcnt_hi(-1, __mbcnt_lo(-1, 0)); }
 
 /*
-Rounding modes are not yet supported in HIP
+HIP specific device functions
 */
 
-__device__ float __double2float_rd(double x);
-__device__ float __double2float_rn(double x);
-__device__ float __double2float_ru(double x);
-__device__ float __double2float_rz(double x);
+// utility union type
+union __u {
+    int i;
+    unsigned int u;
+    float f;
+};
 
-__device__ int __double2hiint(double x);
+__device__ static inline unsigned __hip_ds_bpermute(int index, unsigned src) {
+    __u tmp; tmp.u = src;
+    tmp.i = __llvm_amdgcn_ds_bpermute(index, tmp.i);
+    return tmp.u;
+}
 
-__device__ int __double2int_rd(double x);
-__device__ int __double2int_rn(double x);
-__device__ int __double2int_ru(double x);
-__device__ int __double2int_rz(double x);
+__device__ static inline float __hip_ds_bpermutef(int index, float src) {
+    __u tmp; tmp.f = src;
+    tmp.i = __llvm_amdgcn_ds_bpermute(index, tmp.i);
+    return tmp.f;
+}
 
-__device__ long long int __double2ll_rd(double x);
-__device__ long long int __double2ll_rn(double x);
-__device__ long long int __double2ll_ru(double x);
-__device__ long long int __double2ll_rz(double x);
+__device__ static inline unsigned __hip_ds_permute(int index, unsigned src) {
+  __u tmp; tmp.u = src;
+  tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i);
+  return tmp.u;
+}
 
-__device__ int __double2loint(double x);
+__device__ static inline float __hip_ds_permutef(int index, float src) {
+  __u tmp; tmp.u = src;
+  tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i);
+  return tmp.u;
+}
 
-__device__ unsigned int __double2uint_rd(double x);
-__device__ unsigned int __double2uint_rn(double x);
-__device__ unsigned int __double2uint_ru(double x);
-__device__ unsigned int __double2uint_rz(double x);
+__device__ static inline unsigned __hip_ds_swizzle(unsigned int src, int pattern) {
+    __u tmp; tmp.u = src;
+    tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern);
+    return tmp.u;
+}
+__device__ static inline float __hip_ds_swizzlef(float src, int pattern) {
+    __u tmp; tmp.f = src;
+    tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern);
+    return tmp.f;
+}
 
-__device__ unsigned long long int __double2ull_rd(double x);
-__device__ unsigned long long int __double2ull_rn(double x);
-__device__ unsigned long long int __double2ull_ru(double x);
-__device__ unsigned long long int __double2ull_rz(double x);
+__device__ static inline int __hip_move_dpp(int src, int dpp_ctrl, int row_mask,
+                                            int bank_mask, bool bound_ctrl) {
+    return __llvm_amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
+}
+
+static constexpr int warpSize = 64;
+
+ __device__
+inline
+int __shfl(int var, int src_lane, int width = warpSize) {
+    int self = __lane_id();
+    int index = src_lane + (self & ~(width-1));
+    return __llvm_amdgcn_ds_bpermute(index<<2, var);
+}
+__device__
+inline
+unsigned int __shfl(unsigned int var, int src_lane, int width = warpSize) {
+     __u tmp; tmp.u = var;
+    tmp.i = __shfl(tmp.i, src_lane, width);
+    return tmp.u;
+}
+__device__
+inline
+float __shfl(float var, int src_lane, int width = warpSize) {
+    __u tmp; tmp.f = var;
+    tmp.i = __shfl(tmp.i, src_lane, width);
+    return tmp.f;
+}
+__device__
+inline
+double __shfl(double var, int src_lane, int width = warpSize) {
+    __u tmp; tmp.f = (float) var;
+    tmp.i = __shfl(tmp.i, src_lane, width);
+    return (double) tmp.f;
+}
+
+ __device__
+inline
+int __shfl_up(int var, unsigned int lane_delta, int width = warpSize) {
+    int self = __lane_id();
+    int index = self - lane_delta;
+    index = (index < (self & ~(width-1)))?self:index;
+    return __llvm_amdgcn_ds_bpermute(index<<2, var);
+}
+__device__
+inline
+unsigned int __shfl_up(unsigned int var, unsigned int lane_delta, int width = warpSize) {
+    __u tmp; tmp.u = var;
+    tmp.i = __shfl_up(tmp.i, lane_delta, width);
+    return tmp.u;
+}
+__device__
+inline
+float __shfl_up(float var, unsigned int lane_delta, int width = warpSize) {
+    __u tmp; tmp.f = var;
+    tmp.i = __shfl_up(tmp.i, lane_delta, width);
+    return tmp.f;
+}
+__device__
+inline
+double __shfl_up(double var, unsigned int lane_delta, int width = warpSize) {
+    __u tmp; tmp.f = (float) var;
+    tmp.i = __shfl_up(tmp.i, lane_delta, width);
+    return (double) tmp.f;
+}
+
+__device__
+inline
+int __shfl_down(int var, unsigned int lane_delta, int width = warpSize) {
+    int self = __lane_id();
+    int index = self + lane_delta;
+    index = (int)((self&(width-1))+lane_delta) >= width?self:index;
+    return __llvm_amdgcn_ds_bpermute(index<<2, var);
+}
+__device__
+inline
+unsigned int __shfl_down(unsigned int var, unsigned int lane_delta, int width = warpSize) {
+    __u tmp; tmp.u = var;
+    tmp.i = __shfl_down(tmp.i, lane_delta, width);
+    return tmp.u;
+}
+__device__
+inline
+float __shfl_down(float var, unsigned int lane_delta, int width = warpSize) {
+    __u tmp; tmp.f = var;
+    tmp.i = __shfl_down(tmp.i, lane_delta, width);
+    return tmp.f;
+}
+__device__
+inline
+double __shfl_down(double var, unsigned int lane_delta, int width = warpSize) {
+    __u tmp; tmp.f = (float) var;
+    tmp.i = __shfl_down(tmp.i, lane_delta, width);
+    return (double) tmp.f;
+}
+
+__device__
+inline
+int __shfl_xor(int var, int lane_mask, int width = warpSize) {
+    int self = __lane_id();
+    int index = self^lane_mask;
+    index = index >= ((self+width)&~(width-1))?self:index;
+    return __llvm_amdgcn_ds_bpermute(index<<2, var);
+}
+__device__
+inline
+unsigned int __shfl_xor(unsigned int var, int lane_mask, int width = warpSize) {
+    __u tmp; tmp.u = var;
+    tmp.i = __shfl_xor(tmp.i, lane_mask, width);
+    return tmp.u;
+}
+__device__
+inline
+float __shfl_xor(float var, int lane_mask, int width = warpSize) {
+    __u tmp; tmp.f = var;
+    tmp.i = __shfl_xor(tmp.i, lane_mask, width);
+    return tmp.f;
+}
+__device__
+inline
+double __shfl_xor(double var, int lane_mask, int width = warpSize) {
+    __u tmp; tmp.f = (float) var;
+    tmp.i = __shfl_xor(tmp.i, lane_mask, width);
+    return (double) tmp.f;
+}
+
+#define MASK1 0x00ff00ff
+#define MASK2 0xff00ff00
+
+__device__ static inline char4 __hip_hc_add8pk(char4 in1, char4 in2) {
+    char4 out;
+    unsigned one1 = in1.w & MASK1;
+    unsigned one2 = in2.w & MASK1;
+    out.w = (one1 + one2) & MASK1;
+    one1 = in1.w & MASK2;
+    one2 = in2.w & MASK2;
+    out.w = out.w | ((one1 + one2) & MASK2);
+    return out;
+}
+
+__device__ static inline char4 __hip_hc_sub8pk(char4 in1, char4 in2) {
+    char4 out;
+    unsigned one1 = in1.w & MASK1;
+    unsigned one2 = in2.w & MASK1;
+    out.w = (one1 - one2) & MASK1;
+    one1 = in1.w & MASK2;
+    one2 = in2.w & MASK2;
+    out.w = out.w | ((one1 - one2) & MASK2);
+    return out;
+}
+
+__device__ static inline char4 __hip_hc_mul8pk(char4 in1, char4 in2) {
+    char4 out;
+    unsigned one1 = in1.w & MASK1;
+    unsigned one2 = in2.w & MASK1;
+    out.w = (one1 * one2) & MASK1;
+    one1 = in1.w & MASK2;
+    one2 = in2.w & MASK2;
+    out.w = out.w | ((one1 * one2) & MASK2);
+    return out;
+}
+
+/*
+ * Rounding modes are not yet supported in HIP
+ * TODO: Conversion functions are not correct, need to fix when BE is ready
+*/
+
+__device__ static inline float __double2float_rd(double x) { return (double)x; }
+__device__ static inline float __double2float_rn(double x) { return (double)x; }
+__device__ static inline float __double2float_ru(double x) { return (double)x; }
+__device__ static inline float __double2float_rz(double x) { return (double)x; }
+
+__device__ static inline int __double2hiint(double x) {
+    static_assert(sizeof(double) == 2 * sizeof(int), "");
+
+    int tmp[2];
+    __builtin_memcpy(tmp, &x, sizeof(tmp));
+
+    return tmp[1];
+}
+__device__ static inline int __double2loint(double x) {
+    static_assert(sizeof(double) == 2 * sizeof(int), "");
+
+    int tmp[2];
+    __builtin_memcpy(tmp, &x, sizeof(tmp));
+
+    return tmp[0];
+}
+
+__device__ static inline int __double2int_rd(double x) { return (int)x; }
+__device__ static inline int __double2int_rn(double x) { return (int)x; }
+__device__ static inline int __double2int_ru(double x) { return (int)x; }
+__device__ static inline int __double2int_rz(double x) { return (int)x; }
+
+__device__ static inline long long int __double2ll_rd(double x) { return (long long int)x; }
+__device__ static inline long long int __double2ll_rn(double x) { return (long long int)x; }
+__device__ static inline long long int __double2ll_ru(double x) { return (long long int)x; }
+__device__ static inline long long int __double2ll_rz(double x) { return (long long int)x; }
+
+__device__ static inline unsigned int __double2uint_rd(double x) { return (unsigned int)x; }
+__device__ static inline unsigned int __double2uint_rn(double x) { return (unsigned int)x; }
+__device__ static inline unsigned int __double2uint_ru(double x) { return (unsigned int)x; }
+__device__ static inline unsigned int __double2uint_rz(double x) { return (unsigned int)x; }
+
+__device__ static inline unsigned long long int __double2ull_rd(double x) {
+    return (unsigned long long int)x;
+}
+__device__ static inline unsigned long long int __double2ull_rn(double x) {
+    return (unsigned long long int)x;
+}
+__device__ static inline unsigned long long int __double2ull_ru(double x) {
+    return (unsigned long long int)x;
+}
+__device__ static inline unsigned long long int __double2ull_rz(double x) {
+    return (unsigned long long int)x;
+}
+
+__device__ static inline long long int __double_as_longlong(double x) {
+    static_assert(sizeof(long long) == sizeof(double), "");
+
+    long long tmp;
+    __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+    return tmp;
+}
 
-__device__ long long int __double_as_longlong(double x);
 /*
 __device__ unsigned short __float2half_rn(float x);
 __device__ float __half2float(unsigned short);
@@ -142,72 +517,457 @@ CUDA implements half as unsigned short whereas, HIP doesn't.
 
 */
 
-__device__ int __float2int_rd(float x);
-__device__ int __float2int_rn(float x);
-__device__ int __float2int_ru(float x);
-__device__ int __float2int_rz(float x);
+__device__ static inline int __float2int_rd(float x) { return (int)__ocml_floor_f32(x); }
+__device__ static inline int __float2int_rn(float x) { return (int)__ocml_rint_f32(x); }
+__device__ static inline int __float2int_ru(float x) { return (int)__ocml_ceil_f32(x); }
+__device__ static inline int __float2int_rz(float x) { return (int)__ocml_trunc_f32(x); }
 
-__device__ long long int __float2ll_rd(float x);
-__device__ long long int __float2ll_rn(float x);
-__device__ long long int __float2ll_ru(float x);
-__device__ long long int __float2ll_rz(float x);
+__device__ static inline long long int __float2ll_rd(float x) { return (long long int)x; }
+__device__ static inline long long int __float2ll_rn(float x) { return (long long int)x; }
+__device__ static inline long long int __float2ll_ru(float x) { return (long long int)x; }
+__device__ static inline long long int __float2ll_rz(float x) { return (long long int)x; }
 
-__device__ unsigned int __float2uint_rd(float x);
-__device__ unsigned int __float2uint_rn(float x);
-__device__ unsigned int __float2uint_ru(float x);
-__device__ unsigned int __float2uint_rz(float x);
+__device__ static inline unsigned int __float2uint_rd(float x) { return (unsigned int)x; }
+__device__ static inline unsigned int __float2uint_rn(float x) { return (unsigned int)x; }
+__device__ static inline unsigned int __float2uint_ru(float x) { return (unsigned int)x; }
+__device__ static inline unsigned int __float2uint_rz(float x) { return (unsigned int)x; }
 
-__device__ unsigned long long int __float2ull_rd(float x);
-__device__ unsigned long long int __float2ull_rn(float x);
-__device__ unsigned long long int __float2ull_ru(float x);
-__device__ unsigned long long int __float2ull_rz(float x);
+__device__ static inline unsigned long long int __float2ull_rd(float x) {
+    return (unsigned long long int)x;
+}
+__device__ static inline unsigned long long int __float2ull_rn(float x) {
+    return (unsigned long long int)x;
+}
+__device__ static inline unsigned long long int __float2ull_ru(float x) {
+    return (unsigned long long int)x;
+}
+__device__ static inline unsigned long long int __float2ull_rz(float x) {
+    return (unsigned long long int)x;
+}
 
-__device__ int __float_as_int(float x);
-__device__ unsigned int __float_as_uint(float x);
-__device__ double __hiloint2double(int hi, int lo);
-__device__ double __int2double_rn(int x);
+__device__ static inline int __float_as_int(float x) {
+    static_assert(sizeof(int) == sizeof(float), "");
 
-__device__ float __int2float_rd(int x);
-__device__ float __int2float_rn(int x);
-__device__ float __int2float_ru(int x);
-__device__ float __int2float_rz(int x);
+    int tmp;
+    __builtin_memcpy(&tmp, &x, sizeof(tmp));
 
-__device__ float __int_as_float(int x);
+    return tmp;
+}
 
-__device__ double __ll2double_rd(long long int x);
-__device__ double __ll2double_rn(long long int x);
-__device__ double __ll2double_ru(long long int x);
-__device__ double __ll2double_rz(long long int x);
+__device__ static inline unsigned int __float_as_uint(float x) {
+    static_assert(sizeof(unsigned int) == sizeof(float), "");
 
-__device__ float __ll2float_rd(long long int x);
-__device__ float __ll2float_rn(long long int x);
-__device__ float __ll2float_ru(long long int x);
-__device__ float __ll2float_rz(long long int x);
+    unsigned int tmp;
+    __builtin_memcpy(&tmp, &x, sizeof(tmp));
 
-__device__ double __longlong_as_double(long long int x);
+    return tmp;
+}
 
-__device__ double __uint2double_rn(int x);
+__device__ static inline double __hiloint2double(int hi, int lo) {
+    static_assert(sizeof(double) == sizeof(uint64_t), "");
 
-__device__ float __uint2float_rd(unsigned int x);
-__device__ float __uint2float_rn(unsigned int x);
-__device__ float __uint2float_ru(unsigned int x);
-__device__ float __uint2float_rz(unsigned int x);
+    uint64_t tmp0 = (static_cast<uint64_t>(hi) << 32ull) | static_cast<uint32_t>(lo);
+    double tmp1;
+    __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
 
-__device__ float __uint_as_float(unsigned int x);
+    return tmp1;
+}
 
-__device__ double __ull2double_rd(unsigned long long int x);
-__device__ double __ull2double_rn(unsigned long long int x);
-__device__ double __ull2double_ru(unsigned long long int x);
-__device__ double __ull2double_rz(unsigned long long int x);
+__device__ static inline double __int2double_rn(int x) { return (double)x; }
 
-__device__ float __ull2float_rd(unsigned long long int x);
-__device__ float __ull2float_rn(unsigned long long int x);
-__device__ float __ull2float_ru(unsigned long long int x);
-__device__ float __ull2float_rz(unsigned long long int x);
+__device__ static inline float __int2float_rd(int x) { return (float)x; }
+__device__ static inline float __int2float_rn(int x) { return (float)x; }
+__device__ static inline float __int2float_ru(int x) { return (float)x; }
+__device__ static inline float __int2float_rz(int x) { return (float)x; }
 
-__device__ char4 __hip_hc_add8pk(char4, char4);
-__device__ char4 __hip_hc_sub8pk(char4, char4);
-__device__ char4 __hip_hc_mul8pk(char4, char4);
+__device__ static inline float __int_as_float(int x) {
+    static_assert(sizeof(float) == sizeof(int), "");
+
+    float tmp;
+    __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+    return tmp;
+}
+
+__device__ static inline double __ll2double_rd(long long int x) { return (double)x; }
+__device__ static inline double __ll2double_rn(long long int x) { return (double)x; }
+__device__ static inline double __ll2double_ru(long long int x) { return (double)x; }
+__device__ static inline double __ll2double_rz(long long int x) { return (double)x; }
+
+__device__ static inline float __ll2float_rd(long long int x) { return (float)x; }
+__device__ static inline float __ll2float_rn(long long int x) { return (float)x; }
+__device__ static inline float __ll2float_ru(long long int x) { return (float)x; }
+__device__ static inline float __ll2float_rz(long long int x) { return (float)x; }
+
+__device__ static inline double __longlong_as_double(long long int x) {
+    static_assert(sizeof(double) == sizeof(long long), "");
+
+    double tmp;
+    __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+    return x;
+}
+
+__device__ static inline double __uint2double_rn(int x) { return (double)x; }
+
+__device__ static inline float __uint2float_rd(unsigned int x) { return (float)x; }
+__device__ static inline float __uint2float_rn(unsigned int x) { return (float)x; }
+__device__ static inline float __uint2float_ru(unsigned int x) { return (float)x; }
+__device__ static inline float __uint2float_rz(unsigned int x) { return (float)x; }
+
+__device__ static inline float __uint_as_float(unsigned int x) {
+   static_assert(sizeof(float) == sizeof(unsigned int), "");
+
+    float tmp;
+    __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+    return tmp;
+}
+
+__device__ static inline double __ull2double_rd(unsigned long long int x) { return (double)x; }
+__device__ static inline double __ull2double_rn(unsigned long long int x) { return (double)x; }
+__device__ static inline double __ull2double_ru(unsigned long long int x) { return (double)x; }
+__device__ static inline double __ull2double_rz(unsigned long long int x) { return (double)x; }
+
+__device__ static inline float __ull2float_rd(unsigned long long int x) { return (float)x; }
+__device__ static inline float __ull2float_rn(unsigned long long int x) { return (float)x; }
+__device__ static inline float __ull2float_ru(unsigned long long int x) { return (float)x; }
+__device__ static inline float __ull2float_rz(unsigned long long int x) { return (float)x; }
+
+#if defined(__HCC__)
+#define __HCC_OR_HIP_CLANG__ 1
+#elif defined(__clang__) && defined(__HIP__)
+#define __HCC_OR_HIP_CLANG__ 1
+#else
+#define __HCC_OR_HIP_CLANG__ 0
+#endif
+
+#ifdef __HCC_OR_HIP_CLANG__
+
+#ifdef __HIP_DEVICE_COMPILE__
+
+// Clock functions
+__device__
+inline
+long long int __clock64() { return (long long int) __builtin_amdgcn_s_memrealtime(); }
+
+__device__
+inline
+long long int  __clock() { return (long long int) __builtin_amdgcn_s_memrealtime(); }
+
+// hip.amdgcn.bc - named sync
+__device__
+inline
+void __named_sync(int a, int b) { __builtin_amdgcn_s_barrier(); }
+
+#endif // __HIP_DEVICE_COMPILE__
+
+// warp vote function __all __any __ballot
+__device__
+inline
+int __all(int predicate) {
+    return __ockl_wfall_i32(predicate);
+}
+
+__device__
+inline
+int __any(int predicate) {
+#ifdef NVCC_COMPAT
+    if (__ockl_wfany_i32(predicate) != 0)
+        return 1;
+    else
+        return 0;
+#else
+    return __ockl_wfany_i32(predicate);
+#endif
+}
+
+// XXX from llvm/include/llvm/IR/InstrTypes.h
+#define ICMP_NE 33
+
+__device__
+inline
+unsigned long long int __ballot(int predicate) {
+    return __llvm_amdgcn_icmp_i32(predicate, 0, ICMP_NE);
+}
+
+__device__
+inline
+unsigned long long int __ballot64(int predicate) {
+    return __llvm_amdgcn_icmp_i32(predicate, 0, ICMP_NE);
+}
+
+// hip.amdgcn.bc - lanemask
+__device__
+inline
+int64_t  __lanemask_gt()
+{
+    int32_t activelane = __ockl_activelane_u32();
+    int64_t ballot = __ballot64(1);
+    if (activelane != 63) {
+        int64_t tmp = (~0UL) << (activelane + 1);
+        return tmp & ballot;
+    }
+    return 0;
+}
+
+__device__
+inline
+int64_t __lanemask_lt()
+{
+    int32_t activelane = __ockl_activelane_u32();
+    int64_t ballot = __ballot64(1);
+    if (activelane == 0)
+        return 0;
+    return ballot;
+}
+
+__device__ inline void* __local_to_generic(void* p) { return p; }
+
+#ifdef __HIP_DEVICE_COMPILE__
+__device__
+inline
+void* __get_dynamicgroupbaseptr()
+{
+    // Get group segment base pointer.
+    return (char*)__local_to_generic((void*)__to_local(__llvm_amdgcn_groupstaticsize()));
+}
+#else
+__device__
+void* __get_dynamicgroupbaseptr();
+#endif // __HIP_DEVICE_COMPILE__
+
+__device__
+inline
+void *__amdgcn_get_dynamicgroupbaseptr() {
+    return __get_dynamicgroupbaseptr();
+}
+
+
+
+// hip.amdgcn.bc - sync threads
+#define __CLK_LOCAL_MEM_FENCE    0x01
+typedef unsigned __cl_mem_fence_flags;
+
+typedef enum __memory_scope {
+  __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
+  __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
+  __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
+  __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
+  __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
+} __memory_scope;
+
+// enum values aligned with what clang uses in EmitAtomicExpr()
+typedef enum __memory_order
+{
+  __memory_order_relaxed = __ATOMIC_RELAXED,
+  __memory_order_acquire = __ATOMIC_ACQUIRE,
+  __memory_order_release = __ATOMIC_RELEASE,
+  __memory_order_acq_rel = __ATOMIC_ACQ_REL,
+  __memory_order_seq_cst = __ATOMIC_SEQ_CST
+} __memory_order;
+
+__device__
+inline
+static void
+__atomic_work_item_fence(__cl_mem_fence_flags flags, __memory_order order, __memory_scope scope)
+{
+    // We're tying global-happens-before and local-happens-before together as does HSA
+    if (order != __memory_order_relaxed) {
+        switch (scope) {
+        case __memory_scope_work_item:
+            break;
+        case __memory_scope_sub_group:
+            switch (order) {
+            case __memory_order_relaxed: break;
+            case __memory_order_acquire: __llvm_fence_acq_sg(); break;
+            case __memory_order_release: __llvm_fence_rel_sg(); break;
+            case __memory_order_acq_rel: __llvm_fence_ar_sg(); break;
+            case __memory_order_seq_cst: __llvm_fence_sc_sg(); break;
+            }
+            break;
+        case __memory_scope_work_group:
+            switch (order) {
+            case __memory_order_relaxed: break;
+            case __memory_order_acquire: __llvm_fence_acq_wg(); break;
+            case __memory_order_release: __llvm_fence_rel_wg(); break;
+            case __memory_order_acq_rel: __llvm_fence_ar_wg(); break;
+            case __memory_order_seq_cst: __llvm_fence_sc_wg(); break;
+            }
+            break;
+        case __memory_scope_device:
+            switch (order) {
+            case __memory_order_relaxed: break;
+            case __memory_order_acquire: __llvm_fence_acq_dev(); break;
+            case __memory_order_release: __llvm_fence_rel_dev(); break;
+            case __memory_order_acq_rel: __llvm_fence_ar_dev(); break;
+            case __memory_order_seq_cst: __llvm_fence_sc_dev(); break;
+            }
+            break;
+        case __memory_scope_all_svm_devices:
+            switch (order) {
+            case __memory_order_relaxed: break;
+            case __memory_order_acquire: __llvm_fence_acq_sys(); break;
+            case __memory_order_release: __llvm_fence_rel_sys(); break;
+            case __memory_order_acq_rel: __llvm_fence_ar_sys(); break;
+            case __memory_order_seq_cst: __llvm_fence_sc_sys(); break;
+            }
+            break;
+        }
+    }
+}
+
+// Memory Fence Functions
+__device__
+inline
+static void __threadfence()
+{
+  __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_device);
+}
+
+__device__
+inline
+static void __threadfence_block()
+{
+  __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_work_group);
+}
+
+__device__
+inline
+static void __threadfence_system()
+{
+  __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_all_svm_devices);
+}
+
+#endif // __HCC_OR_HIP_CLANG__
+
+#ifdef __HCC__
+
+/**
+ * extern __shared__
+ */
+
+// Macro to replace extern __shared__ declarations
+// to local variable definitions
+#define HIP_DYNAMIC_SHARED(type, var) type* var = (type*)__get_dynamicgroupbaseptr();
+
+#define HIP_DYNAMIC_SHARED_ATTRIBUTE
+
+
+#elif defined(__clang__) && defined(__HIP__)
+
+#pragma push_macro("__DEVICE__")
+#define __DEVICE__ extern "C" __device__ __attribute__((always_inline)) \
+  __attribute__((weak))
+
+__DEVICE__
+inline
+void __assert_fail(const char * __assertion,
+                                     const char *__file,
+                                     unsigned int __line,
+                                     const char *__function)
+{
+    // Ignore all the args for now.
+    __builtin_trap();
+}
+
+__DEVICE__
+inline
+void __assertfail(const char * __assertion,
+                  const char *__file,
+                  unsigned int __line,
+                  const char *__function,
+                  size_t charsize)
+{
+    // ignore all the args for now.
+    __builtin_trap();
+}
+
+__device__
+inline
+static void __work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope)
+{
+    if (flags) {
+        __atomic_work_item_fence(flags, __memory_order_release, scope);
+        __builtin_amdgcn_s_barrier();
+        __atomic_work_item_fence(flags, __memory_order_acquire, scope);
+    } else {
+        __builtin_amdgcn_s_barrier();
+    }
+}
+
+__device__
+inline
+static void __barrier(int n)
+{
+  __work_group_barrier((__cl_mem_fence_flags)n, __memory_scope_work_group);
+}
+
+__device__
+inline
+__attribute__((noduplicate))
+void __syncthreads()
+{
+  __barrier(__CLK_LOCAL_MEM_FENCE);
+}
+
+// hip.amdgcn.bc - device routine
+/*
+   HW_ID Register bit structure
+   WAVE_ID     3:0     Wave buffer slot number. 0-9.
+   SIMD_ID     5:4     SIMD which the wave is assigned to within the CU.
+   PIPE_ID     7:6     Pipeline from which the wave was dispatched.
+   CU_ID       11:8    Compute Unit the wave is assigned to.
+   SH_ID       12      Shader Array (within an SE) the wave is assigned to.
+   SE_ID       14:13   Shader Engine the wave is assigned to.
+   TG_ID       19:16   Thread-group ID
+   VM_ID       23:20   Virtual Memory ID
+   QUEUE_ID    26:24   Queue from which this wave was dispatched.
+   STATE_ID    29:27   State ID (graphics only, not compute).
+   ME_ID       31:30   Micro-engine ID.
+ */
+
+#define HW_ID               4
+
+#define HW_ID_CU_ID_SIZE    4
+#define HW_ID_CU_ID_OFFSET  8
+
+#define HW_ID_SE_ID_SIZE    2
+#define HW_ID_SE_ID_OFFSET  13
+
+/*
+   Encoding of parameter bitmask
+   HW_ID        5:0     HW_ID
+   OFFSET       10:6    Range: 0..31
+   SIZE         15:11   Range: 1..32
+ */
+
+#define GETREG_IMMED(SZ,OFF,REG) (SZ << 11) | (OFF << 6) | REG
+
+__device__
+inline
+unsigned __smid(void)
+{
+    unsigned cu_id = __builtin_amdgcn_s_getreg(
+            GETREG_IMMED(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID));
+    unsigned se_id = __builtin_amdgcn_s_getreg(
+            GETREG_IMMED(HW_ID_SE_ID_SIZE, HW_ID_SE_ID_OFFSET, HW_ID));
+
+    /* Each shader engine has 16 CU */
+    return (se_id << HW_ID_CU_ID_SIZE) + cu_id;
+}
+
+#pragma push_macro("__DEVICE__")
+
+// Macro to replace extern __shared__ declarations
+// to local variable definitions
+#define HIP_DYNAMIC_SHARED(type, var) \
+    type* var = (type*)__amdgcn_get_dynamicgroupbaseptr();
+
+#define HIP_DYNAMIC_SHARED_ATTRIBUTE
+
+
+#endif //defined(__clang__) && defined(__HIP__)
 
 
 #endif
diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h
new file mode 100644
index 0000000000..2a14b0b814
--- /dev/null
+++ b/include/hip/hcc_detail/device_library_decls.h
@@ -0,0 +1,80 @@
+/*
+Copyright (c) 2015 - present 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.
+*/
+
+/**
+ *  @file  hcc_detail/device_library_decls.h
+ *  @brief Contains declarations for types and functions in device library.
+ */
+
+#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_LIBRARY_DECLS_H
+#define HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_LIBRARY_DECLS_H
+
+#include "hip/hcc_detail/host_defines.h"
+
+extern "C" __device__ __attribute__((const)) bool __ockl_wfany_i32(int);
+extern "C" __device__ __attribute__((const)) bool __ockl_wfall_i32(int);
+extern "C" __device__ uint __ockl_activelane_u32(void);
+
+extern "C" __device__ __attribute__((const)) uint __ockl_mul24_u32(uint, uint);
+extern "C" __device__ __attribute__((const)) int __ockl_mul24_i32(int, int);
+extern "C" __device__ __attribute__((const)) uint __ockl_mul_hi_u32(uint, uint);
+extern "C" __device__ __attribute__((const)) int __ockl_mul_hi_i32(int, int);
+extern "C" __device__ __attribute__((const)) uint __ockl_sad_u32(uint, uint, uint);
+
+extern "C" __device__ __attribute__((const)) float __ocml_floor_f32(float);
+extern "C" __device__ __attribute__((const)) float __ocml_rint_f32(float);
+extern "C" __device__ __attribute__((const)) float __ocml_ceil_f32(float);
+extern "C" __device__ __attribute__((const)) float __ocml_trunc_f32(float);
+
+extern "C" __device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
+extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);
+
+// Introduce local address space
+#define __local __attribute__((address_space(3)))
+
+#ifdef __HIP_DEVICE_COMPILE__
+__device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; }
+#endif //__HIP_DEVICE_COMPILE__
+
+// __llvm_fence* functions from device-libs/irif/src/fence.ll
+extern "C" __device__ void __llvm_fence_acq_sg(void);
+extern "C" __device__ void __llvm_fence_acq_wg(void);
+extern "C" __device__ void __llvm_fence_acq_dev(void);
+extern "C" __device__ void __llvm_fence_acq_sys(void);
+
+extern "C" __device__ void __llvm_fence_rel_sg(void);
+extern "C" __device__ void __llvm_fence_rel_wg(void);
+extern "C" __device__ void __llvm_fence_rel_dev(void);
+extern "C" __device__ void __llvm_fence_rel_sys(void);
+
+extern "C" __device__ void __llvm_fence_ar_sg(void);
+extern "C" __device__ void __llvm_fence_ar_wg(void);
+extern "C" __device__ void __llvm_fence_ar_dev(void);
+extern "C" __device__ void __llvm_fence_ar_sys(void);
+
+
+extern "C" __device__ void __llvm_fence_sc_sg(void);
+extern "C" __device__ void __llvm_fence_sc_wg(void);
+extern "C" __device__ void __llvm_fence_sc_dev(void);
+extern "C" __device__ void __llvm_fence_sc_sys(void);
+
+#endif
diff --git a/include/hip/hcc_detail/hip_ldg.h b/include/hip/hcc_detail/hip_ldg.h
index 281550cd4c..a5b80b0a1b 100644
--- a/include/hip/hcc_detail/hip_ldg.h
+++ b/include/hip/hcc_detail/hip_ldg.h
@@ -23,54 +23,81 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H
 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H
 
-#if defined __HCC__
-#if __hcc_workweek__ >= 16164
+#if defined(__HCC_OR_HIP_CLANG__)
+#if __hcc_workweek__ >= 16164 || defined(__HIP_CLANG_ONLY__)
 #include "hip_vector_types.h"
 #include "host_defines.h"
 
-__device__ char __ldg(const char*);
-__device__ char2 __ldg(const char2*);
-__device__ char4 __ldg(const char4*);
-__device__ signed char __ldg(const signed char*);
-__device__ unsigned char __ldg(const unsigned char*);
+__device__ inline static char __ldg(const char* ptr) { return *ptr; }
 
-__device__ short __ldg(const short*);
-__device__ short2 __ldg(const short2*);
-__device__ short4 __ldg(const short4*);
-__device__ unsigned short __ldg(const unsigned short*);
+__device__ inline static char2 __ldg(const char2* ptr) { return *ptr; }
 
-__device__ int __ldg(const int*);
-__device__ int2 __ldg(const int2*);
-__device__ int4 __ldg(const int4*);
-__device__ unsigned int __ldg(const unsigned int*);
+__device__ inline static char4 __ldg(const char4* ptr) { return *ptr; }
+
+__device__ inline static signed char __ldg(const signed char* ptr) { return ptr[0]; }
+
+__device__ inline static unsigned char __ldg(const unsigned char* ptr) { return ptr[0]; }
 
 
-__device__ long __ldg(const long*);
-__device__ unsigned long __ldg(const unsigned long*);
+__device__ inline static short __ldg(const short* ptr) { return ptr[0]; }
 
-__device__ long long __ldg(const long long*);
-__device__ longlong2 __ldg(const longlong2*);
-__device__ unsigned long long __ldg(const unsigned long long*);
+__device__ inline static short2 __ldg(const short2* ptr) { return ptr[0]; }
 
-__device__ uchar2 __ldg(const uchar2*);
-__device__ uchar4 __ldg(const uchar4*);
+__device__ inline static short4 __ldg(const short4* ptr) { return ptr[0]; }
 
-__device__ ushort2 __ldg(const ushort2*);
+__device__ inline static unsigned short __ldg(const unsigned short* ptr) { return ptr[0]; }
 
-__device__ uint2 __ldg(const uint2*);
-__device__ uint4 __ldg(const uint4*);
 
-__device__ ulonglong2 __ldg(const ulonglong2*);
+__device__ inline static int __ldg(const int* ptr) { return ptr[0]; }
 
-__device__ float __ldg(const float*);
-__device__ float2 __ldg(const float2*);
-__device__ float4 __ldg(const float4*);
+__device__ inline static int2 __ldg(const int2* ptr) { return ptr[0]; }
 
-__device__ double __ldg(const double*);
-__device__ double2 __ldg(const double2*);
+__device__ inline static int4 __ldg(const int4* ptr) { return ptr[0]; }
 
-#endif  // __hcc_workweek__
+__device__ inline static unsigned int __ldg(const unsigned int* ptr) { return ptr[0]; }
 
-#endif  // __HCC__
+
+__device__ inline static long __ldg(const long* ptr) { return ptr[0]; }
+
+__device__ inline static unsigned long __ldg(const unsigned long* ptr) { return ptr[0]; }
+
+
+__device__ inline static long long __ldg(const long long* ptr) { return ptr[0]; }
+
+__device__ inline static longlong2 __ldg(const longlong2* ptr) { return ptr[0]; }
+
+__device__ inline static unsigned long long __ldg(const unsigned long long* ptr) { return ptr[0]; }
+
+
+__device__ inline static uchar2 __ldg(const uchar2* ptr) { return ptr[0]; }
+
+__device__ inline static uchar4 __ldg(const uchar4* ptr) { return ptr[0]; }
+
+
+__device__ inline static ushort2 __ldg(const ushort2* ptr) { return ptr[0]; }
+
+
+__device__ inline static uint2 __ldg(const uint2* ptr) { return ptr[0]; }
+
+__device__ inline static uint4 __ldg(const uint4* ptr) { return ptr[0]; }
+
+
+__device__ inline static ulonglong2 __ldg(const ulonglong2* ptr) { return ptr[0]; }
+
+
+__device__ inline static float __ldg(const float* ptr) { return ptr[0]; }
+
+__device__ inline static float2 __ldg(const float2* ptr) { return ptr[0]; }
+
+__device__ inline static float4 __ldg(const float4* ptr) { return ptr[0]; }
+
+
+__device__ inline static double __ldg(const double* ptr) { return ptr[0]; }
+
+__device__ inline static double2 __ldg(const double2* ptr) { return ptr[0]; }
+
+#endif  // __hcc_workweek__ || defined(__HIP_CLANG_ONLY__)
+
+#endif  // defined(__HCC_OR_HIP_CLANG__)
 
 #endif  // HIP_LDG_H
diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index f67da8d51e..13eaf4fda4 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -110,9 +110,9 @@ extern int HIP_TRACE_API;
 #include <hip/hcc_detail/host_defines.h>
 #include <hip/hcc_detail/math_functions.h>
 #include <hip/hcc_detail/device_functions.h>
+#include <hip/hcc_detail/surface_functions.h>
 #if __HCC__
 #include <hip/hcc_detail/texture_functions.h>
-#include <hip/hcc_detail/surface_functions.h>
 #endif // __HCC__
 
 // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
@@ -184,45 +184,11 @@ extern int HIP_TRACE_API;
 #define __HCC_C__
 #endif
 
-// TODO - hipify-clang - change to use the function call.
-//#define warpSize hc::__wavesize()
-static constexpr int warpSize = 64;
-
-#define clock_t long long int
-__device__ long long int clock64();
-__device__ clock_t clock();
-
 // abort
 __device__ void abort();
 
-// warp vote function __all __any __ballot
-__device__ int __all(int input);
-__device__ int __any(int input);
-__device__ unsigned long long int __ballot(int input);
-
 #if __HIP_ARCH_GFX701__ == 0
 
-// warp shuffle functions
-#ifdef __cplusplus
-__device__ int __shfl(int input, int lane, int width = warpSize);
-__device__ int __shfl_up(int input, unsigned int lane_delta, int width = warpSize);
-__device__ int __shfl_down(int input, unsigned int lane_delta, int width = warpSize);
-__device__ int __shfl_xor(int input, int lane_mask, int width = warpSize);
-__device__ float __shfl(float input, int lane, int width = warpSize);
-__device__ float __shfl_up(float input, unsigned int lane_delta, int width = warpSize);
-__device__ float __shfl_down(float input, unsigned int lane_delta, int width = warpSize);
-__device__ float __shfl_xor(float input, int lane_mask, int width = warpSize);
-#else
-__device__ int __shfl(int input, int lane, int width);
-__device__ int __shfl_up(int input, unsigned int lane_delta, int width);
-__device__ int __shfl_down(int input, unsigned int lane_delta, int width);
-__device__ int __shfl_xor(int input, int lane_mask, int width);
-__device__ float __shfl(float input, int lane, int width);
-__device__ float __shfl_up(float input, unsigned int lane_delta, int width);
-__device__ float __shfl_down(float input, unsigned int lane_delta, int width);
-__device__ float __shfl_xor(float input, int lane_mask, int width);
-#endif  //__cplusplus
-
 __device__ unsigned __hip_ds_bpermute(int index, unsigned src);
 __device__ float __hip_ds_bpermutef(int index, float src);
 __device__ unsigned __hip_ds_permute(int index, unsigned src);
@@ -235,85 +201,15 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
 
 #endif  //__HIP_ARCH_GFX803__ == 1
 
-__host__ __device__ int min(int arg1, int arg2);
-__host__ __device__ int max(int arg1, int arg2);
+__device__ inline static int min(int arg1, int arg2) {
+    return (arg1 < arg2) ? arg1 : arg2;
+}
+__device__ inline static int max(int arg1, int arg2) {
+    return (arg1 > arg2) ? arg1 : arg2;
+}
 
-__device__ void* __get_dynamicgroupbaseptr();
-
-
-/**
- * CUDA 8 device function features
-
- */
-
-
-/**
- * Kernel launching
- */
-
-/**
- *-------------------------------------------------------------------------------------------------
- *-------------------------------------------------------------------------------------------------
- *  @defgroup Fence Fence Functions
- *  @{
- *
- *
- *  @warning The HIP memory fence functions are currently not supported yet.
- *  If any of those threadfence stubs are reached by the application, you should set "export
- *HSA_DISABLE_CACHE=1" to disable L1 and L2 caches.
- *
- *
- *  On AMD platforms, the threadfence* routines are currently empty stubs.
- */
-
-extern __attribute__((const)) __device__ void __hip_hc_threadfence() __asm("__llvm_fence_sc_dev");
-extern __attribute__((const)) __device__ void __hip_hc_threadfence_block() __asm(
-    "__llvm_fence_sc_wg");
-
-
-/**
- * @brief threadfence_block makes writes visible to threads running in same block.
- *
- * @Returns void
- *
- * @param void
- *
- * @warning __threadfence_block is a stub and map to no-op.
- */
-// __device__ void  __threadfence_block(void);
-__device__ static inline void __threadfence_block(void) { return __hip_hc_threadfence_block(); }
-
-/**
- * @brief threadfence makes wirtes visible to other threads running on same GPU.
- *
- * @Returns void
- *
- * @param void
- *
- * @warning __threadfence is a stub and map to no-op, application should set "export
- * HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches.
- */
-// __device__ void  __threadfence(void) __attribute__((deprecated("Provided for compile-time
-// compatibility, not yet functional")));
-__device__ static inline void __threadfence(void) { return __hip_hc_threadfence(); }
-
-/**
- * @brief threadfence_system makes writes to pinned system memory visible on host CPU.
- *
- * @Returns void
- *
- * @param void
- *
- * @warning __threadfence_system is a stub and map to no-op.
- */
-//__device__ void  __threadfence_system(void) __attribute__((deprecated("Provided with workaround
-//configuration, see hip_kernel_language.md for details")));
-__device__ void __threadfence_system(void);
-
-// doxygen end Fence Fence
-/**
- * @}
- */
+__host__ inline static int min(int arg1, int arg2) { return std::min(arg1, arg2); }
+__host__ inline static int max(int arg1, int arg2) { return std::max(arg1, arg2); }
 
 #endif  // __HCC_OR_HIP_CLANG__
 
@@ -437,17 +333,6 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
 
 #endif  //__HCC_CPP__
 
-/**
- * extern __shared__
- */
-
-// Macro to replace extern __shared__ declarations
-// to local variable definitions
-#define HIP_DYNAMIC_SHARED(type, var) type* var = (type*)__get_dynamicgroupbaseptr();
-
-#define HIP_DYNAMIC_SHARED_ATTRIBUTE
-
-
 /**
  * @defgroup HIP-ENV HIP Environment Variables
  * @{
@@ -566,33 +451,6 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
 #define hipGridDim_y gridDim.y
 #define hipGridDim_z gridDim.z
 
-#pragma push_macro("__DEVICE__")
-#define __DEVICE__ extern "C" __device__ __attribute__((always_inline)) \
-  __attribute__((weak))
-
-__DEVICE__ void __device_trap() __asm("llvm.trap");
-
-__DEVICE__ void inline __assert_fail(const char * __assertion,
-                                     const char *__file,
-                                     unsigned int __line,
-                                     const char *__function)
-{
-    // Ignore all the args for now.
-    __device_trap();
-}
-
-extern "C" __device__ __attribute__((noduplicate)) void __syncthreads();
-extern "C" __device__ void *__amdgcn_get_dynamicgroupbaseptr();
-
-// Macro to replace extern __shared__ declarations
-// to local variable definitions
-#define HIP_DYNAMIC_SHARED(type, var) \
-    type* var = (type*)__amdgcn_get_dynamicgroupbaseptr();
-
-#define HIP_DYNAMIC_SHARED_ATTRIBUTE
-
-#pragma push_macro("__DEVICE__")
-
 #include <hip/hcc_detail/math_functions.h>
 
 #endif
diff --git a/include/hip/hcc_detail/llvm_intrinsics.h b/include/hip/hcc_detail/llvm_intrinsics.h
new file mode 100644
index 0000000000..dc6fd05c52
--- /dev/null
+++ b/include/hip/hcc_detail/llvm_intrinsics.h
@@ -0,0 +1,70 @@
+/*
+Copyright (c) 2015 - present 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.
+*/
+
+/**
+ *  @file  hcc_detail/llvm_intrinsics.h
+ *  @brief Contains declarations for wrapper functions for llvm intrinsics
+ *         like llvm.amdgcn.s.barrier.
+ */
+
+#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_LLVM_INTRINSICS_H
+#define HIP_INCLUDE_HIP_HCC_DETAIL_LLVM_INTRINSICS_H
+
+#include "hip/hcc_detail/host_defines.h"
+
+__device__
+__attribute__((convergent))
+ulong __llvm_amdgcn_icmp_i32(uint x, uint y, uint z) __asm("llvm.amdgcn.icmp.i32");
+
+__device__
+unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize");
+
+__device__
+unsigned int __llvm_bitrev_b32(unsigned int src0) __asm("llvm.bitreverse.i32");
+
+__device__
+uint64_t __llvm_bitrev_b64(uint64_t src0) __asm("llvm.bitreverse.i64");
+
+extern
+__device__
+__attribute__((const))
+unsigned int __mbcnt_lo(unsigned int x, unsigned int y) __asm("llvm.amdgcn.mbcnt.lo");
+
+extern
+__device__
+__attribute__((const))
+unsigned int __mbcnt_hi(unsigned int x, unsigned int y) __asm("llvm.amdgcn.mbcnt.hi");
+
+__device__
+int __llvm_amdgcn_ds_bpermute(int index, int src) __asm("llvm.amdgcn.ds.bpermute");
+
+__device__
+int __llvm_amdgcn_ds_permute(int index, int src) __asm("llvm.amdgcn.ds.permute");
+
+__device__
+int __llvm_amdgcn_ds_swizzle(int index, int pattern) __asm("llvm.amdgcn.ds.swizzle");
+
+__device__
+int __llvm_amdgcn_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask,
+                           bool bound_ctrl) __asm("llvm.amdgcn.mov.dpp.i32");
+
+#endif
diff --git a/include/hip/hcc_detail/surface_functions.h b/include/hip/hcc_detail/surface_functions.h
index 562cc440ed..607f221901 100644
--- a/include/hip/hcc_detail/surface_functions.h
+++ b/include/hip/hcc_detail/surface_functions.h
@@ -23,8 +23,6 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_SURFACE_FUNCTIONS_H
 #define HIP_INCLUDE_HIP_HCC_DETAIL_SURFACE_FUNCTIONS_H
 
-#include <hc.hpp>
-#include <hc_short_vector.hpp>
 #include <hip/hcc_detail/hip_surface_types.h>
 
 #define __SURFACE_FUNCTIONS_DECL__ static __inline__ __device__
diff --git a/packaging/hip_hcc.txt b/packaging/hip_hcc.txt
index 04293f2044..9d4b96761d 100644
--- a/packaging/hip_hcc.txt
+++ b/packaging/hip_hcc.txt
@@ -5,7 +5,6 @@ install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
 install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib)
 install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib)
 install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
-install(FILES @hip_SOURCE_DIR@/src/hip_hc.ll DESTINATION lib)
 install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip)
 install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/packaging/hip-targets-release.cmake DESTINATION lib/cmake/hip)
 
diff --git a/src/device_functions.cpp b/src/device_functions.cpp
deleted file mode 100644
index 8dfd5a07c4..0000000000
--- a/src/device_functions.cpp
+++ /dev/null
@@ -1,387 +0,0 @@
-/*
-Copyright (c) 2015 - present 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/device_functions.h>
-#include <hc.hpp>
-#include <grid_launch.h>
-#include <hc_math.hpp>
-#include "device_util.h"
-
-__device__ float __double2float_rd(double x) { return (double)x; }
-__device__ float __double2float_rn(double x) { return (double)x; }
-__device__ float __double2float_ru(double x) { return (double)x; }
-__device__ float __double2float_rz(double x) { return (double)x; }
-
-
-__device__ int __double2hiint(double x) {
-    static_assert(sizeof(double) == 2 * sizeof(int), "");
-
-    int tmp[2];
-    __builtin_memcpy(tmp, &x, sizeof(tmp));
-
-    return tmp[1];
-}
-__device__ int __double2loint(double x) {
-    static_assert(sizeof(double) == 2 * sizeof(int), "");
-
-    int tmp[2];
-    __builtin_memcpy(tmp, &x, sizeof(tmp));
-
-    return tmp[0];
-}
-
-
-__device__ int __double2int_rd(double x) { return (int)x; }
-__device__ int __double2int_rn(double x) { return (int)x; }
-__device__ int __double2int_ru(double x) { return (int)x; }
-__device__ int __double2int_rz(double x) { return (int)x; }
-
-__device__ long long int __double2ll_rd(double x) { return (long long int)x; }
-__device__ long long int __double2ll_rn(double x) { return (long long int)x; }
-__device__ long long int __double2ll_ru(double x) { return (long long int)x; }
-__device__ long long int __double2ll_rz(double x) { return (long long int)x; }
-
-
-__device__ unsigned int __double2uint_rd(double x) { return (unsigned int)x; }
-__device__ unsigned int __double2uint_rn(double x) { return (unsigned int)x; }
-__device__ unsigned int __double2uint_ru(double x) { return (unsigned int)x; }
-__device__ unsigned int __double2uint_rz(double x) { return (unsigned int)x; }
-
-__device__ unsigned long long int __double2ull_rd(double x) { return (unsigned long long int)x; }
-__device__ unsigned long long int __double2ull_rn(double x) { return (unsigned long long int)x; }
-__device__ unsigned long long int __double2ull_ru(double x) { return (unsigned long long int)x; }
-__device__ unsigned long long int __double2ull_rz(double x) { return (unsigned long long int)x; }
-
-__device__ long long int __double_as_longlong(double x) {
-    static_assert(sizeof(long long) == sizeof(double), "");
-
-    long long tmp;
-    __builtin_memcpy(&tmp, &x, sizeof(tmp));
-
-    return tmp;
-}
-
-__device__ int __float2int_rd(float x) { return (int)__ocml_floor_f32(x); }
-__device__ int __float2int_rn(float x) { return (int)__ocml_rint_f32(x); }
-__device__ int __float2int_ru(float x) { return (int)__ocml_ceil_f32(x); }
-__device__ int __float2int_rz(float x) { return (int)__ocml_trunc_f32(x); }
-
-__device__ long long int __float2ll_rd(float x) { return (long long int)x; }
-__device__ long long int __float2ll_rn(float x) { return (long long int)x; }
-__device__ long long int __float2ll_ru(float x) { return (long long int)x; }
-__device__ long long int __float2ll_rz(float x) { return (long long int)x; }
-
-__device__ unsigned int __float2uint_rd(float x) { return (unsigned int)x; }
-__device__ unsigned int __float2uint_rn(float x) { return (unsigned int)x; }
-__device__ unsigned int __float2uint_ru(float x) { return (unsigned int)x; }
-__device__ unsigned int __float2uint_rz(float x) { return (unsigned int)x; }
-
-__device__ unsigned long long int __float2ull_rd(float x) { return (unsigned long long int)x; }
-__device__ unsigned long long int __float2ull_rn(float x) { return (unsigned long long int)x; }
-__device__ unsigned long long int __float2ull_ru(float x) { return (unsigned long long int)x; }
-__device__ unsigned long long int __float2ull_rz(float x) { return (unsigned long long int)x; }
-
-__device__ int __float_as_int(float x) {
-    static_assert(sizeof(int) == sizeof(float), "");
-
-    int tmp;
-    __builtin_memcpy(&tmp, &x, sizeof(tmp));
-
-    return tmp;
-}
-__device__ unsigned int __float_as_uint(float x) {
-    static_assert(sizeof(unsigned int) == sizeof(float), "");
-
-    unsigned int tmp;
-    __builtin_memcpy(&tmp, &x, sizeof(tmp));
-
-    return tmp;
-}
-__device__ double __hiloint2double(int32_t hi, int32_t lo) {
-    static_assert(sizeof(double) == sizeof(uint64_t), "");
-
-    uint64_t tmp0 = (static_cast<uint64_t>(hi) << 32ull) | static_cast<uint32_t>(lo);
-    double tmp1;
-    __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
-
-    return tmp1;
-}
-__device__ double __int2double_rn(int x) { return (double)x; }
-
-__device__ float __int2float_rd(int x) { return (float)x; }
-__device__ float __int2float_rn(int x) { return (float)x; }
-__device__ float __int2float_ru(int x) { return (float)x; }
-__device__ float __int2float_rz(int x) { return (float)x; }
-
-__device__ float __int_as_float(int x) {
-    static_assert(sizeof(float) == sizeof(int), "");
-
-    float tmp;
-    __builtin_memcpy(&tmp, &x, sizeof(tmp));
-
-    return tmp;
-}
-
-__device__ double __ll2double_rd(long long int x) { return (double)x; }
-__device__ double __ll2double_rn(long long int x) { return (double)x; }
-__device__ double __ll2double_ru(long long int x) { return (double)x; }
-__device__ double __ll2double_rz(long long int x) { return (double)x; }
-
-__device__ float __ll2float_rd(long long int x) { return (float)x; }
-__device__ float __ll2float_rn(long long int x) { return (float)x; }
-__device__ float __ll2float_ru(long long int x) { return (float)x; }
-__device__ float __ll2float_rz(long long int x) { return (float)x; }
-
-__device__ double __longlong_as_double(long long int x) {
-    static_assert(sizeof(double) == sizeof(long long), "");
-
-    double tmp;
-    __builtin_memcpy(&tmp, &x, sizeof(tmp));
-
-    return x;
-}
-
-__device__ double __uint2double_rn(int x) { return (double)x; }
-
-__device__ float __uint2float_rd(unsigned int x) { return (float)x; }
-__device__ float __uint2float_rn(unsigned int x) { return (float)x; }
-__device__ float __uint2float_ru(unsigned int x) { return (float)x; }
-__device__ float __uint2float_rz(unsigned int x) { return (float)x; }
-
-__device__ float __uint_as_float(unsigned int x) {
-    static_assert(sizeof(float) == sizeof(unsigned int), "");
-
-    float tmp;
-    __builtin_memcpy(&tmp, &x, sizeof(tmp));
-
-    return tmp;
-}
-
-__device__ double __ull2double_rd(unsigned long long int x) { return (double)x; }
-__device__ double __ull2double_rn(unsigned long long int x) { return (double)x; }
-__device__ double __ull2double_ru(unsigned long long int x) { return (double)x; }
-__device__ double __ull2double_rz(unsigned long long int x) { return (double)x; }
-
-__device__ float __ull2float_rd(unsigned long long int x) { return (float)x; }
-__device__ float __ull2float_rn(unsigned long long int x) { return (float)x; }
-__device__ float __ull2float_ru(unsigned long long int x) { return (float)x; }
-__device__ float __ull2float_rz(unsigned long long int x) { return (float)x; }
-
-/*
-Integer Intrinsics
-*/
-
-// integer intrinsic function __poc __clz __ffs __brev
-__device__ unsigned int __popc(unsigned int input) { return hc::__popcount_u32_b32(input); }
-
-__device__ unsigned int __popcll(unsigned long long int input) {
-    return hc::__popcount_u32_b64(input);
-}
-
-__device__ unsigned int __clz(unsigned int input) {
-#ifdef NVCC_COMPAT
-    return input == 0 ? 32 : hc::__firstbit_u32_u32(input);
-#else
-    return hc::__firstbit_u32_u32(input);
-#endif
-}
-
-__device__ unsigned int __clzll(unsigned long long int input) {
-#ifdef NVCC_COMPAT
-    return input == 0 ? 64 : hc::__firstbit_u32_u64(input);
-#else
-    return hc::__firstbit_u32_u64(input);
-#endif
-}
-
-__device__ unsigned int __clz(int input) {
-#ifdef NVCC_COMPAT
-    return input == 0 ? 32 : hc::__firstbit_u32_s32(input);
-#else
-    return hc::__firstbit_u32_s32(input);
-#endif
-}
-
-__device__ unsigned int __clzll(long long int input) {
-#ifdef NVCC_COMPAT
-    return input == 0 ? 64 : hc::__firstbit_u32_s64(input);
-#else
-    return hc::__firstbit_u32_s64(input);
-#endif
-}
-
-__device__ unsigned int __ffs(unsigned int input) {
-#ifdef NVCC_COMPAT
-    return hc::__lastbit_u32_u32(input) + 1;
-#else
-    return hc::__lastbit_u32_u32(input);
-#endif
-}
-
-__device__ unsigned int __ffsll(unsigned long long int input) {
-#ifdef NVCC_COMPAT
-    return hc::__lastbit_u32_u64(input) + 1;
-#else
-    return hc::__lastbit_u32_u64(input);
-#endif
-}
-
-__device__ unsigned int __ffs(int input) {
-#ifdef NVCC_COMPAT
-    return hc::__lastbit_u32_s32(input) + 1;
-#else
-    return hc::__lastbit_u32_s32(input);
-#endif
-}
-
-__device__ unsigned int __ffsll(long long int input) {
-#ifdef NVCC_COMPAT
-    return hc::__lastbit_u32_s64(input) + 1;
-#else
-    return hc::__lastbit_u32_s64(input);
-#endif
-}
-
-__device__ unsigned int __brev(unsigned int input) { return hc::__bitrev_b32(input); }
-
-__device__ unsigned long long int __brevll(unsigned long long int input) {
-    return hc::__bitrev_b64(input);
-}
-
-struct ucharHolder {
-    union {
-        unsigned char c[4];
-        unsigned int ui;
-    };
-} __attribute__((aligned(4)));
-
-struct uchar2Holder {
-    union {
-        unsigned int ui[2];
-        unsigned char c[8];
-    };
-} __attribute__((aligned(8)));
-
-__device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s) {
-    struct uchar2Holder cHoldVal;
-    struct ucharHolder cHoldKey;
-    struct ucharHolder cHoldOut;
-    cHoldKey.ui = s;
-    cHoldVal.ui[0] = x;
-    cHoldVal.ui[1] = y;
-    cHoldOut.c[0] = cHoldVal.c[cHoldKey.c[0]];
-    cHoldOut.c[1] = cHoldVal.c[cHoldKey.c[1]];
-    cHoldOut.c[2] = cHoldVal.c[cHoldKey.c[2]];
-    cHoldOut.c[3] = cHoldVal.c[cHoldKey.c[3]];
-    return cHoldOut.ui;
-}
-
-__device__ long long __mul64hi(long long int x, long long int y) {
-    ulong x0 = (ulong)x & 0xffffffffUL;
-    long x1 = x >> 32;
-    ulong y0 = (ulong)y & 0xffffffffUL;
-    long y1 = y >> 32;
-    ulong z0 = x0*y0;
-    long t = x1*y0 + (z0 >> 32);
-    long z1 = t & 0xffffffffL;
-    long z2 = t >> 32;
-    z1 = x0*y1 + z1;
-    return x1*y1 + z2 + (z1 >> 32);
-}
-
-__device__ unsigned long long __umul64hi(unsigned long long int x, unsigned long long int y) {
-    ulong x0 = x & 0xffffffffUL;
-    ulong x1 = x >> 32;
-    ulong y0 = y & 0xffffffffUL;
-    ulong y1 = y >> 32;
-    ulong z0 = x0*y0;
-    ulong t = x1*y0 + (z0 >> 32);
-    ulong z1 = t & 0xffffffffUL;
-    ulong z2 = t >> 32;
-    z1 = x0*y1 + z1;
-    return x1*y1 + z2 + (z1 >> 32);
-}
-
-/*
-HIP specific device functions
-*/
-
-__device__ unsigned __hip_ds_bpermute(int index, unsigned src) {
-    return hc::__amdgcn_ds_bpermute(index, src);
-}
-
-__device__ float __hip_ds_bpermutef(int index, float src) {
-    return hc::__amdgcn_ds_bpermute(index, src);
-}
-
-__device__ unsigned __hip_ds_permute(int index, unsigned src) {
-    return hc::__amdgcn_ds_permute(index, src);
-}
-
-__device__ float __hip_ds_permutef(int index, float src) {
-    return hc::__amdgcn_ds_permute(index, src);
-}
-
-__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern) {
-    return hc::__amdgcn_ds_swizzle(src, pattern);
-}
-
-__device__ float __hip_ds_swizzlef(float src, int pattern) {
-    return hc::__amdgcn_ds_swizzle(src, pattern);
-}
-
-__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) {
-    return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
-}
-
-#define MASK1 0x00ff00ff
-#define MASK2 0xff00ff00
-
-__device__ char4 __hip_hc_add8pk(char4 in1, char4 in2) {
-    char4 out;
-    unsigned one1 = in1.w & MASK1;
-    unsigned one2 = in2.w & MASK1;
-    out.w = (one1 + one2) & MASK1;
-    one1 = in1.w & MASK2;
-    one2 = in2.w & MASK2;
-    out.w = out.w | ((one1 + one2) & MASK2);
-    return out;
-}
-
-__device__ char4 __hip_hc_sub8pk(char4 in1, char4 in2) {
-    char4 out;
-    unsigned one1 = in1.w & MASK1;
-    unsigned one2 = in2.w & MASK1;
-    out.w = (one1 - one2) & MASK1;
-    one1 = in1.w & MASK2;
-    one2 = in2.w & MASK2;
-    out.w = out.w | ((one1 - one2) & MASK2);
-    return out;
-}
-
-__device__ char4 __hip_hc_mul8pk(char4 in1, char4 in2) {
-    char4 out;
-    unsigned one1 = in1.w & MASK1;
-    unsigned one2 = in2.w & MASK1;
-    out.w = (one1 * one2) & MASK1;
-    one1 = in1.w & MASK2;
-    one2 = in2.w & MASK2;
-    out.w = out.w | ((one1 * one2) & MASK2);
-    return out;
-}
diff --git a/src/device_util.cpp b/src/device_util.cpp
index 613e35f0cc..34e198c61d 100644
--- a/src/device_util.cpp
+++ b/src/device_util.cpp
@@ -144,70 +144,7 @@ __device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) {
     return dst;
 }
 
-__device__ long long int clock64() { return (long long int)hc::__cycle_u64(); };
-__device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); };
-
 // abort
 __device__ void abort() { return hc::abort(); }
 
-// warp vote function __all __any __ballot
-__device__ int __all(int input) { return hc::__all(input); }
-
-
-__device__ int __any(int input) {
-#ifdef NVCC_COMPAT
-    if (hc::__any(input) != 0)
-        return 1;
-    else
-        return 0;
-#else
-    return hc::__any(input);
-#endif
-}
-
-__device__ unsigned long long int __ballot(int input) { return hc::__ballot(input); }
-
-// warp shuffle functions
-__device__ int __shfl(int input, int lane, int width) { return hc::__shfl(input, lane, width); }
-
-__device__ int __shfl_up(int input, unsigned int lane_delta, int width) {
-    return hc::__shfl_up(input, lane_delta, width);
-}
-
-__device__ int __shfl_down(int input, unsigned int lane_delta, int width) {
-    return hc::__shfl_down(input, lane_delta, width);
-}
-
-__device__ int __shfl_xor(int input, int lane_mask, int width) {
-    return hc::__shfl_xor(input, lane_mask, width);
-}
-
-__device__ float __shfl(float input, int lane, int width) { return hc::__shfl(input, lane, width); }
-
-__device__ float __shfl_up(float input, unsigned int lane_delta, int width) {
-    return hc::__shfl_up(input, lane_delta, width);
-}
-
-__device__ float __shfl_down(float input, unsigned int lane_delta, int width) {
-    return hc::__shfl_down(input, lane_delta, width);
-}
-
-__device__ float __shfl_xor(float input, int lane_mask, int width) {
-    return hc::__shfl_xor(input, lane_mask, width);
-}
-
-__host__ __device__ int min(int arg1, int arg2) {
-    return (int)(hc::precise_math::fmin((float)arg1, (float)arg2));
-}
-__host__ __device__ int max(int arg1, int arg2) {
-    return (int)(hc::precise_math::fmax((float)arg1, (float)arg2));
-}
-
-__device__ void* __get_dynamicgroupbaseptr() {
-    return hc::get_dynamic_group_segment_base_pointer();
-}
-
 __host__ void* __get_dynamicgroupbaseptr() { return nullptr; }
-
-
-__device__ void __threadfence_system(void) { std::atomic_thread_fence(std::memory_order_seq_cst); }
\ No newline at end of file
diff --git a/src/device_util.h b/src/device_util.h
index 6603689d82..8fa96da9d9 100644
--- a/src/device_util.h
+++ b/src/device_util.h
@@ -125,7 +125,6 @@ __device__ double __hip_fast_dsqrt_rd(double x);
 __device__ double __hip_fast_dsqrt_rn(double x);
 __device__ double __hip_fast_dsqrt_ru(double x);
 __device__ double __hip_fast_dsqrt_rz(double x);
-__device__ void __threadfence_system(void);
 
 float __hip_host_j0f(float x);
 double __hip_host_j0(double x);
diff --git a/src/hip_hc.ll b/src/hip_hc.ll
deleted file mode 100644
index aba9205912..0000000000
--- a/src/hip_hc.ll
+++ /dev/null
@@ -1,30 +0,0 @@
-target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
-target triple = "amdgcn--amdhsa"
-
-define i32 @__hip_hc_ir_mul24_int(i32 %a, i32 %b) #1 {
-  %1 = tail call i32 asm sideeffect "v_mul_i32_i24 $0, $1, $2","=v,v,v"(i32 %a, i32 %b)
-  ret i32 %1
-}
-
-define i32 @__hip_hc_ir_umul24_int(i32 %a, i32 %b) #1 {
-  %1 = tail call i32 asm sideeffect "v_mul_u32_u24 $0, $1, $2","=v,v,v"(i32 %a, i32 %b)
-  ret i32 %1
-}
-
-define i32 @__hip_hc_ir_mulhi_int(i32 %a, i32 %b) #1 {
-  %1 = tail call i32 asm sideeffect "v_mul_hi_i32 $0, $1, $2","=v,v,v"(i32 %a, i32 %b)
-  ret i32 %1
-}
-
-define i32 @__hip_hc_ir_umulhi_int(i32 %a, i32 %b) #1 {
-  %1 = tail call i32 asm sideeffect "v_mul_hi_u32 $0, $1, $2","=v,v,v"(i32 %a, i32 %b)
-  ret i32 %1
-}
-
-define i32 @__hip_hc_ir_usad_int(i32 %a, i32 %b, i32 %c) #1 {
-  %1 = tail call i32 asm sideeffect "v_sad_u32 $0, $1, $2, $3","=v,v,v,v"(i32 %a, i32 %b, i32 %c)
-  ret i32 %1
-}
-
-attributes #1 = { alwaysinline nounwind }
-
diff --git a/src/hip_ldg.cpp b/src/hip_ldg.cpp
deleted file mode 100644
index bf94c05571..0000000000
--- a/src/hip_ldg.cpp
+++ /dev/null
@@ -1,83 +0,0 @@
-/*
-Copyright (c) 2015 - present 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/hcc_detail/hip_ldg.h"
-#include "hip/hcc_detail/hip_vector_types.h"
-
-__device__ char __ldg(const char* ptr) { return *ptr; }
-
-__device__ char2 __ldg(const char2* ptr) { return *ptr; }
-
-__device__ char4 __ldg(const char4* ptr) { return *ptr; }
-
-__device__ signed char __ldg(const signed char* ptr) { return ptr[0]; }
-
-__device__ unsigned char __ldg(const unsigned char* ptr) { return ptr[0]; }
-
-__device__ short __ldg(const short* ptr) { return ptr[0]; }
-
-__device__ short2 __ldg(const short2* ptr) { return ptr[0]; }
-
-__device__ short4 __ldg(const short4* ptr) { return ptr[0]; }
-
-__device__ unsigned short __ldg(const unsigned short* ptr) { return ptr[0]; }
-
-__device__ int __ldg(const int* ptr) { return ptr[0]; }
-
-__device__ int2 __ldg(const int2* ptr) { return ptr[0]; }
-
-__device__ int4 __ldg(const int4* ptr) { return ptr[0]; }
-
-__device__ unsigned int __ldg(const unsigned int* ptr) { return ptr[0]; }
-
-
-__device__ long __ldg(const long* ptr) { return ptr[0]; }
-
-__device__ unsigned long __ldg(const unsigned long* ptr) { return ptr[0]; }
-
-__device__ long long __ldg(const long long* ptr) { return ptr[0]; }
-
-__device__ longlong2 __ldg(const longlong2* ptr) { return ptr[0]; }
-
-__device__ unsigned long long __ldg(const unsigned long long* ptr) { return ptr[0]; }
-
-__device__ uchar2 __ldg(const uchar2* ptr) { return ptr[0]; }
-
-__device__ uchar4 __ldg(const uchar4* ptr) { return ptr[0]; }
-
-__device__ ushort2 __ldg(const ushort2* ptr) { return ptr[0]; }
-
-__device__ uint2 __ldg(const uint2* ptr) { return ptr[0]; }
-
-__device__ uint4 __ldg(const uint4* ptr) { return ptr[0]; }
-
-__device__ ulonglong2 __ldg(const ulonglong2* ptr) { return ptr[0]; }
-
-__device__ float __ldg(const float* ptr) { return ptr[0]; }
-
-__device__ float2 __ldg(const float2* ptr) { return ptr[0]; }
-
-__device__ float4 __ldg(const float4* ptr) { return ptr[0]; }
-
-__device__ double __ldg(const double* ptr) { return ptr[0]; }
-
-__device__ double2 __ldg(const double2* ptr) { return ptr[0]; }
diff --git a/tests/src/deviceLib/hip_test_ldg.cpp b/tests/src/deviceLib/hip_test_ldg.cpp
index 4dea81d9b9..7274baa92c 100644
--- a/tests/src/deviceLib/hip_test_ldg.cpp
+++ b/tests/src/deviceLib/hip_test_ldg.cpp
@@ -35,7 +35,7 @@ THE SOFTWARE.
 #include "hip/hip_vector_types.h"
 #include "test_common.h"
 
-#if (__hcc_workweek__ >= 16164) || defined(__HIP_PLATFORM_NVCC__)
+#if (__hcc_workweek__ >= 16164) || defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_CLANG_ONLY__)
 
 #define HIP_ASSERT(x) (assert((x) == hipSuccess))