diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc
index 4f56ffd875..d2614297c3 100755
--- a/hipamd/bin/hipcc
+++ b/hipamd/bin/hipcc
@@ -115,8 +115,13 @@ if ($HIP_PLATFORM eq "clang") {
     $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
     $HIPCC="$HIP_CLANG_PATH/clang++";
     $HIPCXXFLAGS .= "-std=c++11 -I$HIP_PATH/include";
-
     $HIPLDFLAGS = "--hip-link --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_PATH/lib -lhip_hcc";
+
+    # If $HIPCC clang++ is not compiled, use clang instead
+    if ( ! -e $HIPCC ) {
+        $HIPCC="$HIP_CLANG_PATH/clang";
+        $HIPLDFLAGS = "--driver-mode=g++ " . $HIPLDFLAGS;
+    }
 } elsif ($HIP_PLATFORM eq "hcc") {
     $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa";
 
@@ -375,15 +380,35 @@ foreach $arg (@ARGV)
         open my $out, ">", $new_file or die "$new_file: $!";
         while (my $line = <$in>) {
             chomp $line;
-            if ($line =~ m/\.a$/) {
+            if ($line =~ m/\.a$/ || $line =~ m/\.lo$/) {
+                my $libFile = $line;
                 my $path = abs_path($line);
                 my @objs = split ('\n', `cd $tmpdir; ar xv $path`);
+                ## Check if all files in .a are object files.
+                my $allIsObj = 1;
+                my $realObjs = "";
                 foreach my $obj (@objs) {
                     chomp $obj;
                     $obj =~ s/^x - //;
                     $obj = "$tmpdir/$obj";
-                    push (@inputs, $obj);
-                    $new_arg = "$new_arg $obj";
+                    my $fileType = `file $obj`;
+                    my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/);
+                    $allIsObj = ($allIsObj and $isObj);
+                    if ($isObj) {
+                        $realObjs = ($realObjs . " " . $obj);
+                    } else {
+                        push (@inputs, $obj);
+                        $new_arg = "$new_arg $obj";
+                    }
+                }
+                chomp $realObjs;
+                if ($allIsObj) {
+                    print $out "$line\n";
+                } elsif ($realObjs) {
+                    my($libBaseName, $libDir, $libExt) = fileparse($libFile);
+                    $libBaseName = mktemp($libBaseName . "XXXX") . $libExt;
+                    system("cd $tmpdir; ar c $libBaseName $realObjs");
+                    print $out "$tmpdir/$libBaseName\n";
                 }
             } else {
                 print $out "$line\n";
@@ -600,6 +625,9 @@ if ($needLDFLAGS and not $compileOnly) {
     $CMD .= " $HIPLDFLAGS";
 }
 $CMD .= " $toolArgs";
+if ($needLDFLAGS and not $compileOnly and $HIP_PLATFORM eq "clang") {
+    $CMD .= " -lgcc_s -lgcc";
+}
 
 if ($verbose & 0x1) {
     print "hipcc-cmd: ", $CMD, "\n";
diff --git a/hipamd/cmake/FindHIP.cmake b/hipamd/cmake/FindHIP.cmake
index a45fb1b44d..b761356c9a 100644
--- a/hipamd/cmake/FindHIP.cmake
+++ b/hipamd/cmake/FindHIP.cmake
@@ -20,6 +20,7 @@ foreach(config ${_hip_configuration_types})
     mark_as_advanced(HIP_HIPCC_FLAGS_${config_upper} HIP_HCC_FLAGS_${config_upper} HIP_NVCC_FLAGS_${config_upper})
 endforeach()
 option(HIP_HOST_COMPILATION_CPP "Host code compilation mode" ON)
+option(HIP_VERBOSE_BUILD "Print out the commands run while compiling the HIP source file.  With the Makefile generator this defaults to VERBOSE variable specified on the command line, but can be forced on with this option." OFF)
 mark_as_advanced(HIP_HOST_COMPILATION_CPP)
 
 ###############################################################################
@@ -471,7 +472,13 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files
                 INPUT "${custom_target_script_pregen}"
                 )
             set(main_dep DEPENDS ${source_file})
-            set(verbose_output "$(VERBOSE)")
+            if(CMAKE_GENERATOR MATCHES "Makefiles")
+                set(verbose_output "$(VERBOSE)")
+            elseif(HIP_VERBOSE_BUILD)
+                set(verbose_output ON)
+            else()
+                set(verbose_output OFF)
+            endif() 
 
             # Create up the comment string
             file(RELATIVE_PATH generated_file_relative_path "${CMAKE_BINARY_DIR}" "${generated_file}")
diff --git a/hipamd/include/hip/hcc_detail/hip_complex.h b/hipamd/include/hip/hcc_detail/hip_complex.h
index 973d5f564b..48a2852a5a 100644
--- a/hipamd/include/hip/hcc_detail/hip_complex.h
+++ b/hipamd/include/hip/hcc_detail/hip_complex.h
@@ -24,9 +24,35 @@ THE SOFTWARE.
 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
 
 #include "hip/hcc_detail/hip_vector_types.h"
-#include <math.h>
+
+// TODO: Clang has a bug which allows device functions to call std functions
+// when std functions are introduced into default namespace by using statement.
+// math.h may be included after this bug is fixed.
+#if __cplusplus
+#include <cmath>
+#else
+#include "math.h"
+#endif
 
 #if __cplusplus
+#define COMPLEX_NEG_OP_OVERLOAD(type)                                                              \
+    __device__ __host__ static inline type operator-(const type& op) {                             \
+        type ret;                                                                                  \
+        ret.x = -op.x;                                                                             \
+        ret.y = -op.y;                                                                             \
+        return ret;                                                                                \
+    }
+
+#define COMPLEX_EQ_OP_OVERLOAD(type)                                                               \
+    __device__ __host__ static inline bool operator==(const type& lhs, const type& rhs) {          \
+        return lhs.x == rhs.x && lhs.y == rhs.y;                                                   \
+    }
+
+#define COMPLEX_NE_OP_OVERLOAD(type)                                                               \
+    __device__ __host__ static inline bool operator!=(const type& lhs, const type& rhs) {          \
+        return !(lhs == rhs);                                                                      \
+    }
+
 #define COMPLEX_ADD_OP_OVERLOAD(type)                                                              \
     __device__ __host__ static inline type operator+(const type& lhs, const type& rhs) {           \
         type ret;                                                                                  \
@@ -94,12 +120,16 @@ THE SOFTWARE.
         ret.y = lhs.y * rhs;                                                                       \
         return ret;                                                                                \
     }
+#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T)                                      \
+    __device__ __host__ ComplexT(T val) : x(val), y(val) {}                                        \
+    __device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {}
 
 #endif
 
 struct hipFloatComplex {
 #ifdef __cplusplus
    public:
+   typedef float value_type;
     __device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {}
     __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {}
     __device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {}
@@ -119,6 +149,7 @@ struct hipFloatComplex {
 struct hipDoubleComplex {
 #ifdef __cplusplus
    public:
+   typedef double value_type;
     __device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {}
     __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {}
     __device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {}
@@ -137,6 +168,9 @@ struct hipDoubleComplex {
 
 #if __cplusplus
 
+COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex)
+COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex)
+COMPLEX_NE_OP_OVERLOAD(hipFloatComplex)
 COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex)
 COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex)
 COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex)
@@ -156,6 +190,9 @@ COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double)
 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long)
 COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long)
 
+COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex)
+COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex)
+COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex)
 COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex)
 COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex)
 COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex)
@@ -221,7 +258,6 @@ __device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hi
 
 __device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
 
-
 __device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
 
 __device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
@@ -301,4 +337,20 @@ __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, h
     return make_hipDoubleComplex(real, imag);
 }
 
+// Complex functions returning real numbers.
+#define __DEFINE_HIP_COMPLEX_REAL_FUN(func, hipFun) \
+__device__ __host__ inline float func(const hipFloatComplex& z) { return hipFun##f(z); } \
+__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipFun(z); }
+
+__DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs)
+__DEFINE_HIP_COMPLEX_REAL_FUN(real, hipCreal)
+__DEFINE_HIP_COMPLEX_REAL_FUN(imag, hipCimag)
+
+// Complex functions returning complex numbers.
+#define __DEFINE_HIP_COMPLEX_FUN(func, hipFun) \
+__device__ __host__ inline hipFloatComplex func(const hipFloatComplex& z) { return hipFun##f(z); } \
+__device__ __host__ inline hipDoubleComplex func(const hipDoubleComplex& z) { return hipFun(z); }
+
+__DEFINE_HIP_COMPLEX_FUN(conj, hipConj)
+
 #endif
diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index 8657bc30a3..777113a9de 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 #pragma once
-
+#include "hip/hcc_detail/host_defines.h"
 #include <assert.h>
 #if defined(__cplusplus)
     #include <algorithm>
diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h
index 8107f00a4e..68f2244014 100644
--- a/hipamd/include/hip/hcc_detail/hip_runtime.h
+++ b/hipamd/include/hip/hcc_detail/hip_runtime.h
@@ -98,6 +98,9 @@ struct Empty_launch_parm {};
 #include "grid_launch_GGL.hpp"
 #endif  // GENERIC_GRID_LAUNCH
 
+#endif // HCC
+
+#if __HCC_OR_HIP_CLANG__
 extern int HIP_TRACE_API;
 
 #ifdef __cplusplus
@@ -107,15 +110,16 @@ 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>
+#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.
 #if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__)
 #define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__
 #endif
 
-
 // TODO-HCC add a dummy implementation of assert, need to replace with a proper kernel exit call.
 #if __HIP_DEVICE_COMPILE__ == 1
 #undef assert
@@ -180,10 +184,6 @@ extern int HIP_TRACE_API;
 #define __HCC_C__
 #endif
 
-#endif  // defined __HCC__
-
-#if __HCC_OR_HIP_CLANG__
-
 // abort
 __device__ void abort();
 
diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h
index a9288dc204..8f7c885fc7 100644
--- a/hipamd/include/hip/hcc_detail/math_functions.h
+++ b/hipamd/include/hip/hcc_detail/math_functions.h
@@ -508,19 +508,19 @@ inline
 float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); }
 __device__
 inline
-float __fdiv_rd(float x, float y) { return __ocml_div_rtp_f32(x, y); }
+float __fdiv_rd(float x, float y) { return x / y; }
 __device__
 inline
-float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); }
+float __fdiv_rn(float x, float y) { return x / y; }
 __device__
 inline
-float __fdiv_ru(float x, float y) { return __ocml_div_rtn_f32(x, y); }
+float __fdiv_ru(float x, float y) { return x / y; }
 __device__
 inline
-float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); }
+float __fdiv_rz(float x, float y) { return x / y; }
 __device__
 inline
-float __fdividef(float x, float y) { return __ocml_div_rte_f32(x, y); }
+float __fdividef(float x, float y) { return x / y; }
 __device__
 inline
 float __fmaf_rd(float x, float y, float z)
@@ -1028,16 +1028,16 @@ inline
 double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); }
 __device__
 inline
-double __ddiv_rd(double x, double y) { return __ocml_div_rtp_f64(x, y); }
+double __ddiv_rd(double x, double y) { return x / y; }
 __device__
 inline
-double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); }
+double __ddiv_rn(double x, double y) { return x / y; }
 __device__
 inline
-double __ddiv_ru(double x, double y) { return __ocml_div_rtn_f64(x, y); }
+double __ddiv_ru(double x, double y) { return x / y; }
 __device__
 inline
-double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); }
+double __ddiv_rz(double x, double y) { return x / y; }
 __device__
 inline
 double __dmul_rd(double x, double y) { return __ocml_mul_rtp_f64(x, y); }
diff --git a/hipamd/include/hip/hcc_detail/math_fwd.h b/hipamd/include/hip/hcc_detail/math_fwd.h
index ac5d6e7dc6..102714e117 100644
--- a/hipamd/include/hip/hcc_detail/math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/math_fwd.h
@@ -69,8 +69,6 @@ float __ocml_cosh_f32(float);
 __device__
 float __ocml_cospi_f32(float);
 __device__
-float __ocml_div_rtz_f32(float, float);
-__device__
 float __ocml_i0_f32(float);
 __device__
 float __ocml_i1_f32(float);
@@ -290,18 +288,6 @@ __attribute__((const))
 float __ocml_mul_rtz_f32(float, float);
 __device__
 __attribute__((const))
-float __ocml_div_rte_f32(float, float);
-__device__
-__attribute__((const))
-float __ocml_div_rtn_f32(float, float);
-__device__
-__attribute__((const))
-float __ocml_div_rtp_f32(float, float);
-__device__
-__attribute__((const))
-float __ocml_div_rtz_f32(float, float);
-__device__
-__attribute__((const))
 float __ocml_sqrt_rte_f32(float);
 __device__
 __attribute__((const))
@@ -598,18 +584,6 @@ __attribute__((const))
 double __ocml_mul_rtz_f64(double, double);
 __device__
 __attribute__((const))
-double __ocml_div_rte_f64(double, double);
-__device__
-__attribute__((const))
-double __ocml_div_rtn_f64(double, double);
-__device__
-__attribute__((const))
-double __ocml_div_rtp_f64(double, double);
-__device__
-__attribute__((const))
-double __ocml_div_rtz_f64(double, double);
-__device__
-__attribute__((const))
 double __ocml_sqrt_rte_f64(double);
 __device__
 __attribute__((const))
diff --git a/hipamd/include/hip/hip_runtime_api.h b/hipamd/include/hip/hip_runtime_api.h
index e18c17e07b..2ff562cc01 100644
--- a/hipamd/include/hip/hip_runtime_api.h
+++ b/hipamd/include/hip/hip_runtime_api.h
@@ -293,6 +293,12 @@ typedef enum hipDeviceAttribute_t {
     hipDeviceAttributeIntegrated,                        ///< iGPU
 } hipDeviceAttribute_t;
 
+enum hipComputeMode {
+    hipComputeModeDefault = 0,
+    hipComputeModeExclusive = 1,
+    hipComputeModeProhibited = 2,
+    hipComputeModeExcusiveProcess = 3
+};
 
 /**
  *     @}
diff --git a/hipamd/packaging/hip-targets.cmake b/hipamd/packaging/hip-targets.cmake
index 65370eec9e..5aff2ca1ed 100644
--- a/hipamd/packaging/hip-targets.cmake
+++ b/hipamd/packaging/hip-targets.cmake
@@ -16,7 +16,7 @@ set(CMAKE_IMPORT_FILE_VERSION 1)
 set(_targetsDefined)
 set(_targetsNotDefined)
 set(_expectedTargets)
-foreach(_expectedTarget hip::hip_hcc_static hip::hip_hcc hip::hip_device)
+foreach(_expectedTarget hip::hip_hcc_static hip::hip_hcc hip::hip_device hip::host hip::device)
   list(APPEND _expectedTargets ${_expectedTarget})
   if(NOT TARGET ${_expectedTarget})
     list(APPEND _targetsNotDefined ${_expectedTarget})
@@ -65,6 +65,24 @@ set_target_properties(hip::hip_device PROPERTIES
   INTERFACE_SYSTEM_INCLUDE_DIRECTORIES "${_IMPORT_PREFIX}/include;/opt/rocm/hsa/include"
 )
 
+# Create imported target hip::host
+add_library(hip::host INTERFACE IMPORTED)
+
+set_target_properties(hip::host PROPERTIES
+  INTERFACE_LINK_LIBRARIES "hip::hip_hcc"
+)
+
+# Create imported target hip::device
+add_library(hip::device INTERFACE IMPORTED)
+
+set_target_properties(hip::device PROPERTIES
+  INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device;hcc::hccrt;hcc::hc_am"
+)
+
+if(CMAKE_VERSION VERSION_LESS 3.0.0)
+  message(FATAL_ERROR "This file relies on consumers using CMake 3.0.0 or greater.")
+endif()
+
 # Load information for each installed configuration.
 get_filename_component(_DIR "${CMAKE_CURRENT_LIST_FILE}" PATH)
 file(GLOB CONFIG_FILES "${_DIR}/hip-targets-*.cmake")
diff --git a/hipamd/src/hip_hcc_internal.h b/hipamd/src/hip_hcc_internal.h
index 4008df1574..be257aff4f 100644
--- a/hipamd/src/hip_hcc_internal.h
+++ b/hipamd/src/hip_hcc_internal.h
@@ -26,6 +26,7 @@ THE SOFTWARE.
 #include <hc.hpp>
 #include <hsa/hsa.h>
 #include <unordered_map>
+#include <stack>
 
 #include "hsa/hsa_ext_amd.h"
 #include "hip/hip_runtime.h"