diff --git a/hipamd/Jenkinsfile b/hipamd/Jenkinsfile
index 540e7e6689..6d37f10e3c 100644
--- a/hipamd/Jenkinsfile
+++ b/hipamd/Jenkinsfile
@@ -296,8 +296,53 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_
 String build_config = 'Release'
 String job_name = env.JOB_NAME.toLowerCase( )
 
-// The following launches 2 builds in parallel: rocm-head and cuda-9.x
-parallel rocm_head:
+// The following launches 3 builds in parallel: rocm-head, rocm-1.9.x and cuda-9.x
+parallel rocm_1_9:
+{
+  node('hip-rocm')
+  {
+    String hcc_ver = 'rocm-1.9.x'
+    String from_image = 'ci_test_nodes/rocm-1.9.x/ubuntu-16.04:latest'
+    String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video'
+
+    // Checkout source code, dependencies and version files
+    String source_hip_rel = checkout_and_version( hcc_ver )
+
+    // Create/reuse a docker image that represents the hip build environment
+    def hip_build_image = docker_build_image( hcc_ver, 'hip', '', source_hip_rel, from_image )
+
+    // Print system information for the log
+    hip_build_image.inside( inside_args )
+    {
+      sh  """#!/usr/bin/env bash
+          set -x
+          /opt/rocm/bin/rocm_agent_enumerator -t ALL
+          /opt/rocm/bin/hcc --version
+        """
+    }
+
+    // Conctruct a binary directory path based on build config
+    String build_hip_rel = build_directory_rel( build_config );
+
+    // Build hip inside of the build environment
+    docker_build_inside_image( hip_build_image, inside_args, hcc_ver, '', build_config, source_hip_rel, build_hip_rel )
+
+    // Clean docker build image
+    docker_clean_images( 'hip', docker_build_image_name( ) )
+
+    // After a successful build, upload a docker image of the results
+    /*
+    String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel )
+    if( params.push_image_to_docker_hub )
+    {
+      docker_upload_dockerhub( job_name, hip_image_name, 'rocm' )
+      docker_clean_images( 'rocm', hip_image_name )
+    }
+    docker_clean_images( job_name, hip_image_name )
+    */
+  }
+},
+rocm_head:
 {
   node('hip-rocm')
   {
diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc
index f22884ea1b..5942b20b8d 100755
--- a/hipamd/bin/hipcc
+++ b/hipamd/bin/hipcc
@@ -687,7 +687,11 @@ if($HIP_PLATFORM eq "hcc" or $HIP_PLATFORM eq "clang"){
         $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX900__=1 ";
     }
     if ($target_gfx906 eq 1) {
-        $HIPLDFLAGS .= " --amdgpu-target=gfx906";
+        $GPU_ARCH_ARG = $GPU_ARCH_OPT . "gfx906";
+        $HIPLDFLAGS .= $GPU_ARCH_ARG;
+        if ($HIP_PLATFORM eq 'clang') {
+            $HIPCXXFLAGS .= $GPU_ARCH_ARG;;
+        }
         $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX906__=1 ";
     }
 }
diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt
index a3952ba418..56de7c7904 100644
--- a/hipamd/hipify-clang/CMakeLists.txt
+++ b/hipamd/hipify-clang/CMakeLists.txt
@@ -23,7 +23,7 @@ set(CMAKE_CXX_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang++)
 set(CMAKE_C_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang)
 
 # Link against LLVM and CLANG libraries
-target_link_libraries(hipify-clang
+target_link_libraries(hipify-clang PRIVATE
     clangASTMatchers
     clangFrontend
     clangTooling
@@ -48,11 +48,11 @@ target_link_libraries(hipify-clang
     LLVMCore)
 
 if(WIN32)
-    target_link_libraries(hipify-clang version)
+    target_link_libraries(hipify-clang PRIVATE version)
 endif()
 
 if ((LLVM_PACKAGE_VERSION VERSION_EQUAL "7") OR (LLVM_PACKAGE_VERSION VERSION_GREATER "7"))
-    target_link_libraries(hipify-clang clangToolingInclusions)
+    target_link_libraries(hipify-clang PRIVATE clangToolingInclusions)
 endif()
 
 set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}")
diff --git a/hipamd/hipify-clang/README.md b/hipamd/hipify-clang/README.md
index 02bc9cfed0..ec46c7cb2f 100644
--- a/hipamd/hipify-clang/README.md
+++ b/hipamd/hipify-clang/README.md
@@ -30,11 +30,7 @@
 ## <a name="dependencies"></a> Dependencies
 
 `hipify-clang` requires:
-1. LLVM+CLANG of at least version 3.8.0, latest stable and recommended releases:
-   
-   6.0.1 (linux and windows),
-
-   7.0.0 (linux only).
+1. LLVM+CLANG of at least version 3.8.0, latest stable and recommended release: 6.0.1 (linux and windows).
 
 2. CUDA at least version 7.5, latest supported release is 9.2.
 
@@ -50,8 +46,8 @@
 | 5.0.1                    | 8.0                               |
 | 5.0.2                    | 8.0                               |
 | 6.0.0                    | 9.0                               |
-| 6.0.1                    | 9.0                               |
-| 7.0.0                    | 9.2                               | linux only |
+| **6.0.1**                | **9.0**                           | **LATEST STABLE RELEASE** |
+| 7.0.0                    | 9.2                               | windows is not supported, on linux there is a clang bug: https://bugs.llvm.org/show_bug.cgi?id=36384  |
 |                          | 10.0                              | not yet supported |
 
 In most cases, you can get a suitable version of LLVM+CLANG with your package manager.
@@ -77,10 +73,10 @@ cmake \
 
 make -j install
 ```
-On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"` and after `cmake` the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make`.
+On Windows, the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"`; the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make.`
 
-Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested, `LLVM+CLANG` should be built in `Debug` mode as well.
-64 bit build mode `-Thost=x64` is supported as well, `LLVM+CLANG` should be built (installed) in 64bit mode as well.
+Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested; `LLVM+CLANG` should be built in `Debug` mode as well.
+64 bit build mode `-Thost=x64` is supported as well; `LLVM+CLANG` should be built in 64bit mode as well.
 
 The binary can then be found at `./dist/bin/hipify-clang`.
 
@@ -108,11 +104,11 @@ To run it:
 
     make -j install
     ```
-    On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"` and after `cmake` the generated `LLVM.sln` should be built by `Visual Studio 15 2017` instead of `make`.
+    On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"`; the generated `LLVM.sln` should be built by `Visual Studio 15 2017` instead of `make`.
 
 3. Ensure [`CUDA`](https://developer.nvidia.com/cuda-toolkit-archive) of minimum version 7.5 is installed.
 
-    * Having multiple CUDA installations in order to choose a concrete version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified:
+    * Having multiple CUDA installations, in order to choose a particular version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified:
 
       `-DCUDA_TOOLKIT_ROOT_DIR="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"`
 
@@ -127,7 +123,9 @@ To run it:
       `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1`
 
 5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed.
+
 6. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM.
+
     * installing `lit` into `python` might be required:
 
       `python f:/LLVM/6.0.1/llvm/utils/lit/setup.py install`,
@@ -139,7 +137,9 @@ To run it:
       `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py`,
 
       where `f:/LLVM/6.0.1/build/Release` is LLVM build directory.
+
 7. Build with the `HIPIFY_CLANG_TESTS` option turned on: -DHIPIFY_CLANG_TESTS=1.
+
 8. `make test-hipify`
 
     On Windows after `cmake` the project `test-hipify` in the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make test-hipify`.
@@ -178,7 +178,7 @@ A corresponding successful output:
 --    - CMake module path: F:/LLVM/6.0.1/dist/lib/cmake/llvm
 --    - Include path     : F:/LLVM/6.0.1/dist/include
 --    - Binary path      : F:/LLVM/6.0.1/dist/bin
--- Found PythonInterp: C:/Program Files/Python36/python.exe (found suitable version "3.6.4", minimum required is "2.7")
+-- Found PythonInterp: C:/Program Files/Python36/python.exe (found suitable version "3.6.4", minimum required is "3.6")
 -- Found lit: C:/Program Files/Python36/Scripts/lit.exe
 -- Found FileCheck: F:/LLVM/6.0.1/dist/bin/FileCheck.exe
 -- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0 (found version "9.0")
diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h
index 42927e3246..7b0a13e83a 100644
--- a/hipamd/include/hip/hcc_detail/device_functions.h
+++ b/hipamd/include/hip/hcc_detail/device_functions.h
@@ -297,9 +297,16 @@ float __shfl(float var, int src_lane, int width = warpSize) {
 __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;
+    static_assert(sizeof(double) == 2 * sizeof(int), "");
+    static_assert(sizeof(double) == sizeof(uint64_t), "");
+
+    int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp));
+    tmp[0] = __shfl(tmp[0], src_lane, width);
+    tmp[1] = __shfl(tmp[1], src_lane, width);
+
+    uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
+    double tmp1;  __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
+    return tmp1;
 }
 
  __device__
@@ -327,9 +334,16 @@ float __shfl_up(float var, unsigned int lane_delta, int width = warpSize) {
 __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;
+    static_assert(sizeof(double) == 2 * sizeof(int), "");
+    static_assert(sizeof(double) == sizeof(uint64_t), "");
+
+    int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp));
+    tmp[0] = __shfl_up(tmp[0], lane_delta, width);
+    tmp[1] = __shfl_up(tmp[1], lane_delta, width);
+
+    uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
+    double tmp1;  __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
+    return tmp1;
 }
 
 __device__
@@ -357,9 +371,16 @@ float __shfl_down(float var, unsigned int lane_delta, int width = warpSize) {
 __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;
+    static_assert(sizeof(double) == 2 * sizeof(int), "");
+    static_assert(sizeof(double) == sizeof(uint64_t), "");
+
+    int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp));
+    tmp[0] = __shfl_down(tmp[0], lane_delta, width);
+    tmp[1] = __shfl_down(tmp[1], lane_delta, width);
+
+    uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
+    double tmp1;  __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
+    return tmp1;
 }
 
 __device__
@@ -387,9 +408,16 @@ float __shfl_xor(float var, int lane_mask, int width = warpSize) {
 __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;
+    static_assert(sizeof(double) == 2 * sizeof(int), "");
+    static_assert(sizeof(double) == sizeof(uint64_t), "");
+
+    int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp));
+    tmp[0] = __shfl_xor(tmp[0], lane_mask, width);
+    tmp[1] = __shfl_xor(tmp[1], lane_mask, width);
+
+    uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
+    double tmp1;  __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
+    return tmp1;
 }
 
 #define MASK1 0x00ff00ff
@@ -752,6 +780,80 @@ void *__amdgcn_get_dynamicgroupbaseptr() {
     return __get_dynamicgroupbaseptr();
 }
 
+#if defined(__HCC__) && (__hcc_minor__ < 3)
+// 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;
+        }
+    }
+}
+#endif
+
 // Memory Fence Functions
 __device__
 inline
diff --git a/hipamd/include/hip/hcc_detail/device_library_decls.h b/hipamd/include/hip/hcc_detail/device_library_decls.h
index 2a0b968fe3..2bf3c8cc51 100644
--- a/hipamd/include/hip/hcc_detail/device_library_decls.h
+++ b/hipamd/include/hip/hcc_detail/device_library_decls.h
@@ -65,6 +65,29 @@ extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float)
 __device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; }
 #endif //__HIP_DEVICE_COMPILE__
 
+#if defined(__HCC__) && (__hcc_minor__ < 3)
+// __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);
+#else
 // Using hip.amdgcn.bc - sync threads
 #define __CLK_LOCAL_MEM_FENCE    0x01
 typedef unsigned __cl_mem_fence_flags;
@@ -90,5 +113,6 @@ typedef enum __memory_order
 // Linked from hip.amdgcn.bc
 extern "C" __device__ void
 __atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope);
+#endif
 
 #endif
diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp
index 1e239cb2fc..2268a203dd 100644
--- a/hipamd/src/hip_stream.cpp
+++ b/hipamd/src/hip_stream.cpp
@@ -31,12 +31,21 @@ THE SOFTWARE.
 //-------------------------------------------------------------------------------------------------
 // Stream
 //
+#if defined(__HCC__) && (__hcc_minor__ < 3)
+enum queue_priority
+{
+    priority_high = 0,
+    priority_normal = 0,
+    priority_low = 0
+};
+#else
 enum queue_priority
 {
     priority_high = Kalmar::priority_high,
     priority_normal = Kalmar::priority_normal,
     priority_low = Kalmar::priority_low
 };
+#endif
 
 //---
 hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priority) {
@@ -59,7 +68,11 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
                 // Obtain mutex access to the device critical data, release by destructor
                 LockedAccessor_CtxCrit_t ctxCrit(ctx->criticalData());
 
+#if defined(__HCC__) && (__hcc_minor__ < 3)
+                auto istream = new ihipStream_t(ctx, acc.create_view(), flags);
+#else
                 auto istream = new ihipStream_t(ctx, acc.create_view(Kalmar::execute_in_order, Kalmar::queuing_mode_automatic, (Kalmar::queue_priority)priority), flags);
+#endif
 
                 ctxCrit->addStream(istream);
                 *stream = istream;
@@ -223,8 +236,12 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
     } else if (stream == hipStreamNull) {
         return ihipLogStatus(hipErrorInvalidResourceHandle);
     } else {
+#if defined(__HCC__) && (__hcc_minor__ < 3)
+        *priority = 0;
+#else
         LockedAccessor_StreamCrit_t crit(stream->_criticalData);
         *priority = crit->_av.get_queue_priority();
+#endif
         return ihipLogStatus(hipSuccess);
     }
 }