diff --git a/projects/hip/docs/markdown/hip-math-api.md b/projects/hip/docs/markdown/hip-math-api.md
index 37efafbbbf..9b8a3f2f11 100644
--- a/projects/hip/docs/markdown/hip-math-api.md
+++ b/projects/hip/docs/markdown/hip-math-api.md
@@ -1433,7 +1433,7 @@ __device__  float __expf(float x);
 __device__ static  float __fadd_rd(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fadd_rn
@@ -1441,7 +1441,7 @@ __device__ static  float __fadd_rd(float x, float y);
 __device__ static  float __fadd_rn(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fadd_ru
@@ -1449,7 +1449,7 @@ __device__ static  float __fadd_rn(float x, float y);
 __device__ static  float __fadd_ru(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fadd_rz
@@ -1457,7 +1457,7 @@ __device__ static  float __fadd_ru(float x, float y);
 __device__ static  float __fadd_rz(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fdiv_rd
@@ -1465,7 +1465,7 @@ __device__ static  float __fadd_rz(float x, float y);
 __device__ static  float __fdiv_rd(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fdiv_rn
@@ -1473,7 +1473,7 @@ __device__ static  float __fdiv_rd(float x, float y);
 __device__ static  float __fdiv_rn(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fdiv_ru
@@ -1481,7 +1481,7 @@ __device__ static  float __fdiv_rn(float x, float y);
 __device__ static  float __fdiv_ru(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fdiv_rz
@@ -1489,7 +1489,7 @@ __device__ static  float __fdiv_ru(float x, float y);
 __device__ static  float __fdiv_rz(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fdividef
@@ -1505,7 +1505,7 @@ __device__ static  float __fdividef(float x, float y);
 __device__  float __fmaf_rd(float x, float y, float z);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fmaf_rn
@@ -1513,7 +1513,7 @@ __device__  float __fmaf_rd(float x, float y, float z);
 __device__  float __fmaf_rn(float x, float y, float z);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fmaf_ru
@@ -1521,7 +1521,7 @@ __device__  float __fmaf_rn(float x, float y, float z);
 __device__  float __fmaf_ru(float x, float y, float z);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fmaf_rz
@@ -1529,7 +1529,7 @@ __device__  float __fmaf_ru(float x, float y, float z);
 __device__  float __fmaf_rz(float x, float y, float z);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fmul_rd
@@ -1537,7 +1537,7 @@ __device__  float __fmaf_rz(float x, float y, float z);
 __device__ static  float __fmul_rd(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fmul_rn
@@ -1545,7 +1545,7 @@ __device__ static  float __fmul_rd(float x, float y);
 __device__ static  float __fmul_rn(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fmul_ru
@@ -1553,7 +1553,7 @@ __device__ static  float __fmul_rn(float x, float y);
 __device__ static  float __fmul_ru(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fmul_rz
@@ -1561,7 +1561,7 @@ __device__ static  float __fmul_ru(float x, float y);
 __device__ static  float __fmul_rz(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __frcp_rd
@@ -1569,7 +1569,7 @@ __device__ static  float __fmul_rz(float x, float y);
 __device__  float __frcp_rd(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __frcp_rn
@@ -1577,7 +1577,7 @@ __device__  float __frcp_rd(float x);
 __device__  float __frcp_rn(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __frcp_ru
@@ -1585,7 +1585,7 @@ __device__  float __frcp_rn(float x);
 __device__  float __frcp_ru(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __frcp_rz
@@ -1593,7 +1593,7 @@ __device__  float __frcp_ru(float x);
 __device__  float __frcp_rz(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __frsqrt_rn
@@ -1601,7 +1601,7 @@ __device__  float __frcp_rz(float x);
 __device__  float __frsqrt_rn(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fsqrt_rd
@@ -1609,7 +1609,7 @@ __device__  float __frsqrt_rn(float x);
 __device__  float __fsqrt_rd(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fsqrt_rn
@@ -1617,7 +1617,7 @@ __device__  float __fsqrt_rd(float x);
 __device__  float __fsqrt_rn(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fsqrt_ru
@@ -1625,7 +1625,7 @@ __device__  float __fsqrt_rn(float x);
 __device__  float __fsqrt_ru(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fsqrt_rz
@@ -1633,7 +1633,7 @@ __device__  float __fsqrt_ru(float x);
 __device__  float __fsqrt_rz(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fsub_rd
@@ -1641,7 +1641,7 @@ __device__  float __fsqrt_rz(float x);
 __device__ static  float __fsub_rd(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fsub_rn
@@ -1649,7 +1649,7 @@ __device__ static  float __fsub_rd(float x, float y);
 __device__ static  float __fsub_rn(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fsub_ru
@@ -1657,7 +1657,15 @@ __device__ static  float __fsub_rn(float x, float y);
 __device__ static  float __fsub_ru(float x, float y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
+
+
+### __fsub_rz
+```cpp 
+__device__ static  float __fsub_rz(float x, float y);
+
+```
+**Description:**  Unsupported
 
 
 ### __log10f
@@ -1729,7 +1737,7 @@ __device__  float __tanf(float x);
 __device__ static  double __dadd_rd(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dadd_rn
@@ -1737,7 +1745,7 @@ __device__ static  double __dadd_rd(double x, double y);
 __device__ static  double __dadd_rn(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dadd_ru
@@ -1745,7 +1753,7 @@ __device__ static  double __dadd_rn(double x, double y);
 __device__ static  double __dadd_ru(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dadd_rz
@@ -1753,7 +1761,7 @@ __device__ static  double __dadd_ru(double x, double y);
 __device__ static  double __dadd_rz(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __ddiv_rd
@@ -1761,7 +1769,7 @@ __device__ static  double __dadd_rz(double x, double y);
 __device__ static  double __ddiv_rd(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __ddiv_rn
@@ -1769,7 +1777,7 @@ __device__ static  double __ddiv_rd(double x, double y);
 __device__ static  double __ddiv_rn(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __ddiv_ru
@@ -1777,7 +1785,7 @@ __device__ static  double __ddiv_rn(double x, double y);
 __device__ static  double __ddiv_ru(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __ddiv_rz
@@ -1785,7 +1793,7 @@ __device__ static  double __ddiv_ru(double x, double y);
 __device__ static  double __ddiv_rz(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dmul_rd
@@ -1793,7 +1801,7 @@ __device__ static  double __ddiv_rz(double x, double y);
 __device__ static  double __dmul_rd(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dmul_rn
@@ -1801,7 +1809,7 @@ __device__ static  double __dmul_rd(double x, double y);
 __device__ static  double __dmul_rn(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dmul_ru
@@ -1809,7 +1817,7 @@ __device__ static  double __dmul_rn(double x, double y);
 __device__ static  double __dmul_ru(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dmul_rz
@@ -1817,7 +1825,7 @@ __device__ static  double __dmul_ru(double x, double y);
 __device__ static  double __dmul_rz(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __drcp_rd
@@ -1825,7 +1833,7 @@ __device__ static  double __dmul_rz(double x, double y);
 __device__  double __drcp_rd(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __drcp_rn
@@ -1833,7 +1841,7 @@ __device__  double __drcp_rd(double x);
 __device__  double __drcp_rn(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __drcp_ru
@@ -1841,7 +1849,7 @@ __device__  double __drcp_rn(double x);
 __device__  double __drcp_ru(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __drcp_rz
@@ -1849,7 +1857,7 @@ __device__  double __drcp_ru(double x);
 __device__  double __drcp_rz(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dsqrt_rd
@@ -1857,7 +1865,7 @@ __device__  double __drcp_rz(double x);
 __device__  double __dsqrt_rd(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dsqrt_rn
@@ -1865,7 +1873,7 @@ __device__  double __dsqrt_rd(double x);
 __device__  double __dsqrt_rn(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dsqrt_ru
@@ -1873,7 +1881,7 @@ __device__  double __dsqrt_rn(double x);
 __device__  double __dsqrt_ru(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dsqrt_rz
@@ -1881,7 +1889,7 @@ __device__  double __dsqrt_ru(double x);
 __device__  double __dsqrt_rz(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dsub_rd
@@ -1889,7 +1897,7 @@ __device__  double __dsqrt_rz(double x);
 __device__ static  double __dsub_rd(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dsub_rn
@@ -1897,7 +1905,7 @@ __device__ static  double __dsub_rd(double x, double y);
 __device__ static  double __dsub_rn(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dsub_ru
@@ -1905,7 +1913,7 @@ __device__ static  double __dsub_rn(double x, double y);
 __device__ static  double __dsub_ru(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dsub_rz
@@ -1913,7 +1921,7 @@ __device__ static  double __dsub_ru(double x, double y);
 __device__ static  double __dsub_rz(double x, double y);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fma_rd
@@ -1921,7 +1929,7 @@ __device__ static  double __dsub_rz(double x, double y);
 __device__  double __fma_rd(double x, double y, double z);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fma_rn
@@ -1929,7 +1937,7 @@ __device__  double __fma_rd(double x, double y, double z);
 __device__  double __fma_rn(double x, double y, double z);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fma_ru
@@ -1937,7 +1945,7 @@ __device__  double __fma_rn(double x, double y, double z);
 __device__  double __fma_ru(double x, double y, double z);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fma_rz
@@ -1945,7 +1953,7 @@ __device__  double __fma_ru(double x, double y, double z);
 __device__  double __fma_rz(double x, double y, double z);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __brev
diff --git a/projects/hip/include/hip/hcc_detail/hip_memory.h b/projects/hip/include/hip/hcc_detail/hip_memory.h
index 739e488ca3..2c9ec1b7c3 100644
--- a/projects/hip/include/hip/hcc_detail/hip_memory.h
+++ b/projects/hip/include/hip/hcc_detail/hip_memory.h
@@ -27,6 +27,8 @@ THE SOFTWARE.
 // HIP heap is implemented as a global array with fixed size. Users may define
 // __HIP_SIZE_OF_PAGE and __HIP_NUM_PAGES to have a larger heap.
 
+#if __HCC__ || __HIP__
+
 // Size of page in bytes.
 #ifndef __HIP_SIZE_OF_PAGE
 #define __HIP_SIZE_OF_PAGE 64
@@ -39,11 +41,8 @@ THE SOFTWARE.
 
 #define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE)
 
-#if __HCC__ || __HIP__
-
-__attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
-__attribute__((weak)) __device__
-    uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
+extern __device__ char __hip_device_heap[];
+extern __device__ uint32_t __hip_device_page_flag[];
 
 extern "C" inline __device__ void* __hip_malloc(size_t size) {
     char* heap = (char*)__hip_device_heap;
diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h
index b12e7aca89..557257b2b0 100644
--- a/projects/hip/include/hip/hcc_detail/math_functions.h
+++ b/projects/hip/include/hip/hcc_detail/math_functions.h
@@ -514,38 +514,41 @@ float __exp10f(float x) { return __ocml_exp10_f32(x); }
 __DEVICE__
 inline
 float __expf(float x) { return __ocml_exp_f32(x); }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
-float __fadd_rd(float x, float y) { return __ocml_add_rtp_f32(x, y); }
+float __fadd_rd(float x, float y) { return __ocml_add_rtn_f32(x, y); }
 __DEVICE__
 inline
 float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); }
 __DEVICE__
 inline
-float __fadd_ru(float x, float y) { return __ocml_add_rtn_f32(x, y); }
+float __fadd_ru(float x, float y) { return __ocml_add_rtp_f32(x, y); }
 __DEVICE__
 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 x / y; }
+float __fdiv_rd(float x, float y) { return __ocml_div_rtn_f32(x, y); }
 __DEVICE__
 inline
-float __fdiv_rn(float x, float y) { return x / y; }
+float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); }
 __DEVICE__
 inline
-float __fdiv_ru(float x, float y) { return x / y; }
+float __fdiv_ru(float x, float y) { return __ocml_div_rtp_f32(x, y); }
 __DEVICE__
 inline
-float __fdiv_rz(float x, float y) { return x / y; }
+float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); }
+#endif
 __DEVICE__
 inline
 float __fdividef(float x, float y) { return x / y; }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 float __fmaf_rd(float x, float y, float z)
 {
-    return __ocml_fma_rtp_f32(x, y, z);
+    return __ocml_fma_rtn_f32(x, y, z);
 }
 __DEVICE__
 inline
@@ -557,7 +560,7 @@ __DEVICE__
 inline
 float __fmaf_ru(float x, float y, float z)
 {
-    return __ocml_fma_rtn_f32(x, y, z);
+    return __ocml_fma_rtp_f32(x, y, z);
 }
 __DEVICE__
 inline
@@ -567,13 +570,13 @@ float __fmaf_rz(float x, float y, float z)
 }
 __DEVICE__
 inline
-float __fmul_rd(float x, float y) { return __ocml_mul_rtp_f32(x, y); }
+float __fmul_rd(float x, float y) { return __ocml_mul_rtn_f32(x, y); }
 __DEVICE__
 inline
 float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); }
 __DEVICE__
 inline
-float __fmul_ru(float x, float y)  { return __ocml_mul_rtn_f32(x, y); }
+float __fmul_ru(float x, float y)  { return __ocml_mul_rtp_f32(x, y); }
 __DEVICE__
 inline
 float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); }
@@ -594,28 +597,29 @@ inline
 float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); }
 __DEVICE__
 inline
-float __fsqrt_rd(float x) { return __ocml_sqrt_f32(x); }
+float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); }
 __DEVICE__
 inline
-float __fsqrt_rn(float x) { return __ocml_sqrt_f32(x); }
+float __fsqrt_rn(float x) { return __ocml_sqrt_rte_f32(x); }
 __DEVICE__
 inline
-float __fsqrt_ru(float x) { return __ocml_sqrt_f32(x); }
+float __fsqrt_ru(float x) { return __ocml_sqrt_rtp_f32(x); }
 __DEVICE__
 inline
-float __fsqrt_rz(float x) { return __ocml_sqrt_f32(x); }
+float __fsqrt_rz(float x) { return __ocml_sqrt_rtz_f32(x); }
 __DEVICE__
 inline
-float __fsub_rd(float x, float y) { return __ocml_sub_rtp_f32(x, y); }
+float __fsub_rd(float x, float y) { return __ocml_sub_rtn_f32(x, y); }
 __DEVICE__
 inline
 float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); }
 __DEVICE__
 inline
-float __fsub_ru(float x, float y) { return __ocml_sub_rtn_f32(x, y); }
+float __fsub_ru(float x, float y) { return __ocml_sub_rtp_f32(x, y); }
 __DEVICE__
 inline
 float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); }
+#endif
 __DEVICE__
 inline
 float __log10f(float x) { return __ocml_log10_f32(x); }
@@ -1034,39 +1038,40 @@ double yn(int n, double x)
 }
 
 // BEGIN INTRINSICS
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
-double __dadd_rd(double x, double y) { return __ocml_add_rtp_f64(x, y); }
+double __dadd_rd(double x, double y) { return __ocml_add_rtn_f64(x, y); }
 __DEVICE__
 inline
 double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); }
 __DEVICE__
 inline
-double __dadd_ru(double x, double y) { return __ocml_add_rtn_f64(x, y); }
+double __dadd_ru(double x, double y) { return __ocml_add_rtp_f64(x, y); }
 __DEVICE__
 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 x / y; }
+double __ddiv_rd(double x, double y) { return __ocml_div_rtn_f64(x, y); }
 __DEVICE__
 inline
-double __ddiv_rn(double x, double y) { return x / y; }
+double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); }
 __DEVICE__
 inline
-double __ddiv_ru(double x, double y) { return x / y; }
+double __ddiv_ru(double x, double y) { return __ocml_div_rtp_f64(x, y); }
 __DEVICE__
 inline
-double __ddiv_rz(double x, double y) { return x / y; }
+double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); }
 __DEVICE__
 inline
-double __dmul_rd(double x, double y) { return __ocml_mul_rtp_f64(x, y); }
+double __dmul_rd(double x, double y) { return __ocml_mul_rtn_f64(x, y); }
 __DEVICE__
 inline
 double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); }
 __DEVICE__
 inline
-double __dmul_ru(double x, double y) { return __ocml_mul_rtn_f64(x, y); }
+double __dmul_ru(double x, double y) { return __ocml_mul_rtp_f64(x, y); }
 __DEVICE__
 inline
 double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); }
@@ -1084,25 +1089,25 @@ inline
 double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); }
 __DEVICE__
 inline
-double __dsqrt_rd(double x) { return __ocml_sqrt_f64(x); }
+double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); }
 __DEVICE__
 inline
-double __dsqrt_rn(double x) { return __ocml_sqrt_f64(x); }
+double __dsqrt_rn(double x) { return __ocml_sqrt_rte_f64(x); }
 __DEVICE__
 inline
-double __dsqrt_ru(double x) { return __ocml_sqrt_f64(x); }
+double __dsqrt_ru(double x) { return __ocml_sqrt_rtp_f64(x); }
 __DEVICE__
 inline
-double __dsqrt_rz(double x) { return __ocml_sqrt_f64(x); }
+double __dsqrt_rz(double x) { return __ocml_sqrt_rtz_f64(x); }
 __DEVICE__
 inline
-double __dsub_rd(double x, double y) { return __ocml_sub_rtp_f64(x, y); }
+double __dsub_rd(double x, double y) { return __ocml_sub_rtn_f64(x, y); }
 __DEVICE__
 inline
 double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); }
 __DEVICE__
 inline
-double __dsub_ru(double x, double y) { return __ocml_sub_rtn_f64(x, y); }
+double __dsub_ru(double x, double y) { return __ocml_sub_rtp_f64(x, y); }
 __DEVICE__
 inline
 double __dsub_rz(double x, double y) { return __ocml_sub_rtz_f64(x, y); }
@@ -1110,7 +1115,7 @@ __DEVICE__
 inline
 double __fma_rd(double x, double y, double z)
 {
-    return __ocml_fma_rtp_f64(x, y, z);
+    return __ocml_fma_rtn_f64(x, y, z);
 }
 __DEVICE__
 inline
@@ -1122,7 +1127,7 @@ __DEVICE__
 inline
 double __fma_ru(double x, double y, double z)
 {
-    return __ocml_fma_rtn_f64(x, y, z);
+    return __ocml_fma_rtp_f64(x, y, z);
 }
 __DEVICE__
 inline
@@ -1130,6 +1135,7 @@ double __fma_rz(double x, double y, double z)
 {
     return __ocml_fma_rtz_f64(x, y, z);
 }
+#endif
 // END INTRINSICS
 // END DOUBLE
 
diff --git a/projects/hip/include/hip/hcc_detail/math_fwd.h b/projects/hip/include/hip/hcc_detail/math_fwd.h
index 404c2f81d5..e5594924ba 100644
--- a/projects/hip/include/hip/hcc_detail/math_fwd.h
+++ b/projects/hip/include/hip/hcc_detail/math_fwd.h
@@ -288,6 +288,30 @@ __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, float);
+__device__
+__attribute__((const))
+float __ocml_sqrt_rtn_f32(float, float);
+__device__
+__attribute__((const))
+float __ocml_sqrt_rtp_f32(float, float);
+__device__
+__attribute__((const))
+float __ocml_sqrt_rtz_f32(float, float);
+__device__
+__attribute__((const))
 float __ocml_fma_rte_f32(float, float, float);
 __device__
 __attribute__((const))
@@ -572,6 +596,30 @@ __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, double);
+__device__
+__attribute__((const))
+double __ocml_sqrt_rtn_f64(double, double);
+__device__
+__attribute__((const))
+double __ocml_sqrt_rtp_f64(double, double);
+__device__
+__attribute__((const))
+double __ocml_sqrt_rtz_f64(double, double);
+__device__
+__attribute__((const))
 double __ocml_fma_rte_f64(double, double, double);
 __device__
 __attribute__((const))
@@ -594,4 +642,4 @@ double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64");
 
 #if defined(__cplusplus)
     } // extern "C"
-#endif
\ No newline at end of file
+#endif
diff --git a/projects/hip/src/hip_clang.cpp b/projects/hip/src/hip_clang.cpp
index 15a96d298a..cfd75df562 100644
--- a/projects/hip/src/hip_clang.cpp
+++ b/projects/hip/src/hip_clang.cpp
@@ -22,6 +22,7 @@ THE SOFTWARE.
 
 #include <unordered_map>
 #include <string>
+#include <fstream>
 
 #include "hip/hip_runtime.h"
 #include "hip_hcc_internal.h"
@@ -86,6 +87,7 @@ __hipRegisterFatBinary(const void* data)
 
     std::string target{&desc->triple[sizeof(AMDGCN_AMDHSA_TRIPLE)],
       desc->tripleSize - sizeof(AMDGCN_AMDHSA_TRIPLE)};
+    tprintf(DB_FB, "Found bundle for %s\n", target.c_str());
 
     for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
       hsa_agent_t agent = g_allAgents[deviceId + 1];
@@ -110,10 +112,35 @@ __hipRegisterFatBinary(const void* data)
 
       if (module->executable.handle) {
         modules->at(deviceId) = module;
+        tprintf(DB_FB, "Loaded code object for %s\n", name);
+        if (HIP_DUMP_CODE_OBJECT) {
+          char fname[30];
+          static std::atomic<int> index;
+          sprintf(fname, "__hip_dump_code_object%04d.o", index++);
+          tprintf(DB_FB, "Dump code object %s\n", fname);
+          std::ofstream ofs;
+          ofs.open(fname, std::ios::binary);
+          ofs << image;
+          ofs.close();
+        }
+      } else {
+        fprintf(stderr, "Failed to load code object for %s\n", name);
+        abort();
       }
     }
   }
 
+  for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
+    hsa_agent_t agent = g_allAgents[deviceId + 1];
+
+    char name[64] = {};
+    hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, name);
+    if (!(*modules)[deviceId]) {
+      fprintf(stderr, "No device code bundle for %s\n", name);
+      abort();
+    }
+  }
+
   tprintf(DB_FB, "__hipRegisterFatBinary succeeds and returns %p\n", modules);
   return modules;
 }
@@ -132,13 +159,18 @@ extern "C" void __hipRegisterFunction(
   dim3*        gridDim,
   int*         wSize)
 {
+  HIP_INIT_API(modules, hostFunction, deviceFunction, deviceName);
   std::vector<hipFunction_t> functions{g_deviceCnt};
 
+  assert(modules && modules->size() >= g_deviceCnt);
   for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
     hipFunction_t function;
     if (hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName)) {
       functions[deviceId] = function;
     }
+    else {
+      tprintf(DB_FB, "missing kernel %s for device %d\n", deviceName, deviceId);
+    }
   }
 
   g_functions.insert(std::make_pair(hostFunction, std::move(functions)));
@@ -180,6 +212,7 @@ hipError_t hipSetupArgument(
   size_t size,
   size_t offset)
 {
+  HIP_INIT_API(arg, size, offset);
   auto ctx = ihipGetTlsDefaultCtx();
   LockedAccessor_CtxCrit_t crit(ctx->criticalData());
   auto& arguments = crit->_execStack.top()._arguments;
@@ -194,6 +227,7 @@ hipError_t hipSetupArgument(
 
 hipError_t hipLaunchByPtr(const void *hostFunction)
 {
+  HIP_INIT_API(hostFunction);
   ihipExec_t exec;
   {
     auto ctx = ihipGetTlsDefaultCtx();
@@ -213,20 +247,26 @@ hipError_t hipLaunchByPtr(const void *hostFunction)
     deviceId = 0;
   }
 
+  hipError_t e = hipSuccess;
   decltype(g_functions)::iterator it;
-  if ((it = g_functions.find(hostFunction)) == g_functions.end())
-    return hipErrorUnknown;
+  if ((it = g_functions.find(hostFunction)) == g_functions.end()) {
+    e = hipErrorUnknown;
+    fprintf(stderr, "kernel %p not found!\n", hostFunction);
+    abort();
+  } else {
+    size_t size = exec._arguments.size();
+    void *extra[] = {
+        HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec._arguments[0],
+        HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
+        HIP_LAUNCH_PARAM_END
+      };
 
-  size_t size = exec._arguments.size();
-  void *extra[] = {
-      HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec._arguments[0],
-      HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
-      HIP_LAUNCH_PARAM_END
-    };
+    e = hipModuleLaunchKernel(it->second[deviceId],
+      exec._gridDim.x, exec._gridDim.y, exec._gridDim.z,
+      exec._blockDim.x, exec._blockDim.y, exec._blockDim.z,
+      exec._sharedMem, exec._hStream, nullptr, extra);
+  }
 
-  return hipModuleLaunchKernel(it->second[deviceId],
-    exec._gridDim.x, exec._gridDim.y, exec._gridDim.z,
-    exec._blockDim.x, exec._blockDim.y, exec._blockDim.z,
-    exec._sharedMem, exec._hStream, nullptr, extra);
+  return ihipLogStatus(e);
 }
 
diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index eff93da847..e152e7ba69 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -97,6 +97,8 @@ int HIP_INIT_ALLOC = -1;
 int HIP_SYNC_STREAM_WAIT = 0;
 int HIP_FORCE_NULL_STREAM = 0;
 
+int HIP_DUMP_CODE_OBJECT = 0;
+
 
 #if (__hcc_workweek__ >= 17300)
 // Make sure we have required bug fix in HCC
@@ -1294,6 +1296,10 @@ void HipReadEnv() {
                "overridden by specifying hipEventReleaseToSystem or hipEventReleaseToDevice flag "
                "when creating the event.");
 
+    READ_ENV_I(release, HIP_DUMP_CODE_OBJECT, 0,
+               "If set, dump code object as __hip_dump_code_object[nnnn].o in the current directory,"
+               "where nnnn is the index number.");
+
     // Some flags have both compile-time and runtime flags - generate a warning if user enables the
     // runtime flag but the compile-time flag is disabled.
     if (HIP_DB && !COMPILE_HIP_DB) {
diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h
index d64a4a4cbe..8102f066de 100644
--- a/projects/hip/src/hip_hcc_internal.h
+++ b/projects/hip/src/hip_hcc_internal.h
@@ -83,11 +83,11 @@ extern int HIP_SYNC_NULL_STREAM;
 extern int HIP_INIT_ALLOC;
 extern int HIP_FORCE_NULL_STREAM;
 
+extern int HIP_DUMP_CODE_OBJECT;
 
 // TODO - remove when this is standard behavior.
 extern int HCC_OPT_FLUSH;
 
-
 // Class to assign a short TID to each new thread, for HIP debugging purposes.
 class TidInfo {
    public:
diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp
index 93ac527826..4ea5b24f43 100644
--- a/projects/hip/src/hip_memory.cpp
+++ b/projects/hip/src/hip_memory.cpp
@@ -28,6 +28,9 @@ THE SOFTWARE.
 #include "hip_hcc_internal.h"
 #include "trace_helper.h"
 
+__device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
+__device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
+
 // Internal HIP APIS:
 namespace hip_internal {
 
diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp
index a6d486b6de..1dc6701fe6 100644
--- a/projects/hip/src/hip_module.cpp
+++ b/projects/hip/src/hip_module.cpp
@@ -267,7 +267,12 @@ inline void track(const Agent_global& x) {
     hc::AmPointerInfo ptr_info(nullptr, x.address, x.address, x.byte_cnt, device->_acc, true,
                                false);
     hc::am_memtracker_add(x.address, ptr_info);
+#if USE_APP_PTR_FOR_CTX
+    hc::am_memtracker_update(x.address, device->_deviceId, 0u, ihipGetTlsDefaultCtx());
+#else
     hc::am_memtracker_update(x.address, device->_deviceId, 0u);
+#endif
+
 }
 
 template <typename Container = vector<Agent_global>>
@@ -342,7 +347,7 @@ hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, hi
 
     tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(), it0->second.cend(), name);
 
-    return dptr ? hipSuccess : hipErrorNotFound;
+    return *dptr ? hipSuccess : hipErrorNotFound;
 }
 
 hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, const char* name) {
@@ -367,7 +372,7 @@ hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, c
 
     tie(*dptr, *bytes) = read_global_description(it->second.cbegin(), it->second.cend(), name);
 
-    return dptr ? hipSuccess : hipErrorNotFound;
+    return *dptr ? hipSuccess : hipErrorNotFound;
 }
 
 hsa_executable_symbol_t find_kernel_by_name(hsa_executable_t executable, const char* kname) {
diff --git a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
index 17cd82c9ab..939bdae743 100644
--- a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
+++ b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
@@ -34,6 +34,7 @@ THE SOFTWARE.
 #pragma clang diagnostic ignored "-Wunused-variable"
 
 __device__ void double_precision_intrinsics() {
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dadd_rd(0.0, 1.0);
     __dadd_rn(0.0, 1.0);
     __dadd_ru(0.0, 1.0);
@@ -62,6 +63,7 @@ __device__ void double_precision_intrinsics() {
     __fma_rn(1.0, 2.0, 3.0);
     __fma_ru(1.0, 2.0, 3.0);
     __fma_rz(1.0, 2.0, 3.0);
+#endif
 }
 
 __global__ void compileDoublePrecisionIntrinsics(int ignored) {
diff --git a/projects/hip/tests/src/deviceLib/hipFloatMath.cpp b/projects/hip/tests/src/deviceLib/hipFloatMath.cpp
index ee83309f28..c6a07e26a9 100644
--- a/projects/hip/tests/src/deviceLib/hipFloatMath.cpp
+++ b/projects/hip/tests/src/deviceLib/hipFloatMath.cpp
@@ -38,11 +38,13 @@ __global__ void floatMath(float* In, float* Out) {
     Out[tid] = __cosf(In[tid]);
     Out[tid] = __exp10f(Out[tid]);
     Out[tid] = __expf(Out[tid]);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     Out[tid] = __frsqrt_rn(Out[tid]);
-    //Out[tid] = __fsqrt_rd(Out[tid]);
-    //Out[tid] = __fsqrt_rn(Out[tid]);
-    //Out[tid] = __fsqrt_ru(Out[tid]);
-    //Out[tid] = __fsqrt_rz(Out[tid]);
+    Out[tid] = __fsqrt_rd(Out[tid]);
+    Out[tid] = __fsqrt_rn(Out[tid]);
+    Out[tid] = __fsqrt_ru(Out[tid]);
+    Out[tid] = __fsqrt_rz(Out[tid]);
+#endif
     Out[tid] = __log10f(Out[tid]);
     Out[tid] = __log2f(Out[tid]);
     Out[tid] = __logf(Out[tid]);
diff --git a/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp
index f3d2a36931..b216b3cb54 100644
--- a/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp
+++ b/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp
@@ -39,6 +39,7 @@ __device__ void single_precision_intrinsics() {
     __cosf(0.0f);
     __exp10f(0.0f);
     __expf(0.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fadd_rd(0.0f, 1.0f);
     __fadd_rn(0.0f, 1.0f);
     __fadd_ru(0.0f, 1.0f);
@@ -47,7 +48,9 @@ __device__ void single_precision_intrinsics() {
     __fdiv_rn(4.0f, 2.0f);
     __fdiv_ru(4.0f, 2.0f);
     __fdiv_rz(4.0f, 2.0f);
+#endif
     __fdividef(4.0f, 2.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fmaf_rd(1.0f, 2.0f, 3.0f);
     __fmaf_rn(1.0f, 2.0f, 3.0f);
     __fmaf_ru(1.0f, 2.0f, 3.0f);
@@ -69,6 +72,7 @@ __device__ void single_precision_intrinsics() {
     __fsub_rn(2.0f, 1.0f);
     __fsub_ru(2.0f, 1.0f);
     __fsub_rz(2.0f, 1.0f);
+#endif
     __log10f(1.0f);
     __log2f(1.0f);
     __logf(1.0f);