diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc
index fadc74c137..30a4e34704 100755
--- a/projects/hip/bin/hipcc
+++ b/projects/hip/bin/hipcc
@@ -370,7 +370,6 @@ if($HIP_PLATFORM eq "nvcc"){
 
 my $toolArgs = "";  # arguments to pass to the hcc or nvcc tool
 my $optArg = ""; # -O args
-my $rdc = 0;
 
 foreach $arg (@ARGV)
 {
@@ -466,14 +465,6 @@ foreach $arg (@ARGV)
     {
         $optArg = $arg;
     }
-    if($arg =~ /-fgpu-rdc/)
-    {
-        $rdc = 1;
-    }
-    if($arg =~ /-fno-gpu-rdc/)
-    {
-        $rdc = 0;
-    }
 
     ## process linker response file for hip-clang
     ## extract object files from static library and pass them directly to
@@ -804,11 +795,7 @@ if ($HIP_PLATFORM eq "clang") {
         $HIPLDFLAGS .= " -O3";
     }
     $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH";
-    if ($rdc eq 0) {
-        $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS";
-    } else {
-        $HIPLDFLAGS .= " $HIP_DEVLIB_FLAGS";
-    }
+    $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS";
     if ($isWindows) {
         $HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility";
     } else {
@@ -848,7 +835,21 @@ if ($runCmd) {
         print ("HIP ($HIP_PATH) was built using hcc $hipConfig{'HCC_VERSION'}, but you are using $HCC_HOME/hcc with version $HCC_VERSION from hipcc. Please rebuild HIP including cmake or update HCC_HOME variable.\n") ;
         die unless $ENV{'HIP_IGNORE_HCC_VERSION'};
     }
-    system ("$CMD") or delete_temp_dirs () and die ();
+    system ("$CMD");
+    if ($? == -1) {
+        print "failed to execute: $!\n";
+        exit($?);
+    }
+    elsif ($? & 127) {
+        printf "child died with signal %d, %s coredump\n",
+        ($? & 127),  ($? & 128) ? 'with' : 'without';
+        exit($?);
+    }
+    else {
+         $CMD_EXIT_CODE = $? >> 8;
+    }
+    $? or delete_temp_dirs ();
+    exit($CMD_EXIT_CODE);
 }
 
 # vim: ts=4:sw=4:expandtab:smartindent
diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md
index aeb834c2ec..7c2d2c6631 100644
--- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md
+++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md
@@ -861,7 +861,7 @@
 | `cuMemAllocManaged`                                       |                               |
 | `cuMemAllocPitch`                                         |                               |
 | `cuMemcpy`                                                |                               |
-| `cuMemcpy2D`                                              |                               |
+| `cuMemcpy2D`                                              | `hipMemcpyParam2D`            |
 | `cuMemcpy2DAsync`                                         |                               |
 | `cuMemcpy2DUnaligned`                                     |                               |
 | `cuMemcpy3D`                                              |                               |
diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md
index d69f5a04a8..5479813675 100644
--- a/projects/hip/docs/markdown/hip_kernel_language.md
+++ b/projects/hip/docs/markdown/hip_kernel_language.md
@@ -35,6 +35,9 @@
 - [Warp Cross-Lane Functions](#warp-cross-lane-functions)
   * [Warp Vote and Ballot Functions](#warp-vote-and-ballot-functions)
   * [Warp Shuffle Functions](#warp-shuffle-functions)
+- [Cooperative Groups Functions](#cooperative-groups-functions)
+- [Warp Matrix Functions](#warp-matrix-functions)
+- [Independent Thread Scheduling](#independent-thread-scheduling)
 - [Profiler Counter Function](#profiler-counter-function)
 - [Assert](#assert)
 - [Printf](#printf)
@@ -599,6 +602,70 @@ float __shfl_xor  (float var, int laneMask, int width=warpSize);
 
 ```
 
+## Cooperative Groups Functions
+
+Cooperative groups is a mechanism for forming and communicating between groups of threads at
+a granularity different than the block.  This feature was introduced in Cuda 9.
+
+HIP does not support any of the kernel language cooperative groups
+types or functions.
+
+
+| **Function** | **Supported in HIP** | **Supported in CUDA** |
+| --- | --- | --- |
+| `void thread_group.sync()` | | ✓ |
+| `unsigned thread_group.size()` | | ✓ |
+| `unsigned thread_group.thread_rank()` | | ✓ |
+| `bool thread_group.is_valid()` | | ✓ |
+| `thread_group tiled_partition(thread_group, size)` | | ✓ |
+| `thread_block_tile<N> tiled_partition<N>(thread_group)` | | ✓ |
+| `thread_block this_thread_block()` | | ✓ |
+| `T thread_block_tile.shfl()` | | ✓ |
+| `T thread_block_tile.shfl_down()` | | ✓ |
+| `T thread_block_tile.shfl_up()` | | ✓ |
+| `T thread_block_tile.shfl_xor()` | | ✓ |
+| `T thread_block_tile.any()` | | ✓ |
+| `T thread_block_tile.all()` | | ✓ |
+| `T thread_block_tile.ballot()` | | ✓ |
+| `T thread_block_tile.match_any()` | | ✓ |
+| `T thread_block_tile.match_all()` | | ✓ |
+| `coalesced_group coalesced_threads()` | | ✓ |
+| `grid_group this_grid()` | | ✓ |
+| `void grid_group.sync()` | | ✓ |
+| `unsigned grid_group.size()` | | ✓ |
+| `unsigned grid_group.thread_rank()` | | ✓ |
+| `bool grid_group.is_valid()` | | ✓ |
+| `multi_grid_group this_multi_grid()` | | ✓ |
+| `void multi_grid_group.sync()` | | ✓ |
+| `unsigned multi_grid_group.size()` | | ✓ |
+| `unsigned multi_grid_group.thread_rank()` | | ✓ |
+| `bool multi_grid_group.is_valid()` | | ✓ |
+
+## Warp Matrix Functions
+
+Warp matrix functions allow a warp to cooperatively operate on small matrices
+whose elements are spread over the lanes in an unspecified manner.  This feature
+was introduced in Cuda 9.
+
+HIP does not support any of the kernel language warp matrix
+types or functions.
+
+| **Function** | **Supported in HIP** | **Supported in CUDA** |
+| --- | --- | --- |
+| `void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned lda)` | | ✓ |
+| `void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned lda, layout_t layout)` | | ✓ |
+| `void store_matrix_sync(T* mptr, fragment<...> &a,  unsigned lda, layout_t layout)` | | ✓ |
+| `void fill_fragment(fragment<...> &a, const T &value)` | | ✓ |
+| `void mma_sync(fragment<...> &d, const fragment<...> &a, const fragment<...> &b, const fragment<...> &c , bool sat)` | | ✓ |
+
+## Independent Thread Scheduling
+
+The hardware support for independent thread scheduling introduced in certain architectures
+supporting Cuda allows threads to progress independently of each other and enables
+intra-warp synchronizations that were previously not allowed.
+
+HIP does not support this type of scheduling.
+
 ## Profiler Counter Function
 
 The Cuda `__prof_trigger()` instruction is not supported.
diff --git a/projects/hip/hipify-clang/src/CUDA2HIP.cpp b/projects/hip/hipify-clang/src/CUDA2HIP.cpp
index 9a26a051ca..1e530745e6 100644
--- a/projects/hip/hipify-clang/src/CUDA2HIP.cpp
+++ b/projects/hip/hipify-clang/src/CUDA2HIP.cpp
@@ -33,6 +33,7 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
   {"driver_types.h",            {"hip/driver_types.h",       "", CONV_INCLUDE,             API_RUNTIME}},
   {"cuda_fp16.h",               {"hip/hip_fp16.h",           "", CONV_INCLUDE,             API_RUNTIME}},
   {"cuda_texture_types.h",      {"hip/hip_texture_types.h",  "", CONV_INCLUDE,             API_RUNTIME}},
+  {"texture_fetch_functions.h", {"",                         "", CONV_INCLUDE,             API_RUNTIME}},
   {"vector_types.h",            {"hip/hip_vector_types.h",   "", CONV_INCLUDE,             API_RUNTIME}},
   {"cuda_profiler_api.h",       {"hip/hip_profile.h",        "", CONV_INCLUDE,             API_RUNTIME}},
   // cuComplex includes
diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp
index b71e49710d..8be20774ea 100644
--- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp
+++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp
@@ -175,8 +175,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_FUNCTION_MAP{
   {"cuMemcpy",                                             {"hipMemcpy_",                                              "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
   // no analogue
   // NOTE: Not equal to cudaMemcpy2D due to different signatures
-  {"cuMemcpy2D",                                           {"hipMemcpy2D_",                                            "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
-  {"cuMemcpy2D_v2",                                        {"hipMemcpy2D_",                                            "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
+  {"cuMemcpy2D",                                           {"hipMemcpyParam2D",                                        "", CONV_MEMORY, API_DRIVER}},
+  {"cuMemcpy2D_v2",                                        {"hipMemcpyParam2D",                                        "", CONV_MEMORY, API_DRIVER}},
   // no analogue
   // NOTE: Not equal to cudaMemcpy2DAsync due to different signatures
   {"cuMemcpy2DAsync",                                      {"hipMemcpy2DAsync_",                                       "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp
index 5ea83c8376..241ca7ecae 100644
--- a/projects/hip/hipify-clang/src/HipifyAction.cpp
+++ b/projects/hip/hipify-clang/src/HipifyAction.cpp
@@ -189,6 +189,9 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) {
       }
       return false;
     case CONV_INCLUDE:
+      if (hipToken.hipName.empty()) {
+        return true;
+      }
       switch (hipToken.apiType) {
         case API_RAND:
           if (hipToken.hipName == "hiprand_kernel.h") {
diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h
index e1f81e4547..808ed216e4 100644
--- a/projects/hip/include/hip/hcc_detail/device_functions.h
+++ b/projects/hip/include/hip/hcc_detail/device_functions.h
@@ -1009,14 +1009,20 @@ void __syncthreads()
 
 #define GETREG_IMMED(SZ,OFF,REG) (SZ << 11) | (OFF << 6) | REG
 
+/*
+  __smid returns the wave's assigned Compute Unit and Shader Engine.
+  The Compute Unit, CU_ID returned in bits 3:0, and Shader Engine, SE_ID in bits 5:4.
+  Note: the results vary over time.
+  SZ minus 1 since SIZE is 1-based.
+*/
 __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));
+            GETREG_IMMED(HW_ID_CU_ID_SIZE-1, 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));
+            GETREG_IMMED(HW_ID_SE_ID_SIZE-1, HW_ID_SE_ID_OFFSET, HW_ID));
 
     /* Each shader engine has 16 CU */
     return (se_id << HW_ID_CU_ID_SIZE) + cu_id;
diff --git a/projects/hip/include/hip/hcc_detail/driver_types.h b/projects/hip/include/hip/hcc_detail/driver_types.h
index 8e1fec11fa..5b2297114f 100644
--- a/projects/hip/include/hip/hcc_detail/driver_types.h
+++ b/projects/hip/include/hip/hcc_detail/driver_types.h
@@ -80,22 +80,22 @@ typedef struct hipArray {
 }hipArray;
 
 typedef struct hip_Memcpy2D {
-    size_t height;
-    size_t widthInBytes;
-    hipArray* dstArray;
-    hipDeviceptr_t dstDevice;
-    void* dstHost;
-    hipMemoryType dstMemoryType;
-    size_t dstPitch;
-    size_t dstXInBytes;
-    size_t dstY;
-    hipArray* srcArray;
-    hipDeviceptr_t srcDevice;
-    const void* srcHost;
-    hipMemoryType srcMemoryType;
-    size_t srcPitch;
     size_t srcXInBytes;
     size_t srcY;
+    hipMemoryType srcMemoryType;
+    const void* srcHost;
+    hipDeviceptr_t srcDevice;
+    hipArray* srcArray;
+    size_t srcPitch;
+    size_t dstXInBytes;
+    size_t dstY;
+    hipMemoryType dstMemoryType;
+    void* dstHost;
+    hipDeviceptr_t dstDevice;
+    hipArray* dstArray;
+    size_t dstPitch;
+    size_t WidthInBytes;
+    size_t Height;
 } hip_Memcpy2D;
 
 
diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
index 5b598b54a8..d870963101 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
@@ -1858,6 +1858,16 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
  */
 hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
                        size_t height, hipMemcpyKind kind);
+
+/**
+*  @brief Copies memory for 2D arrays.
+*  @param[in]   pCopy Parameters for the memory copy
+ *  @return      #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
+ * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
+ *
+ *  @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray,
+ * hipMemcpyToSymbol, hipMemcpyAsync
+*/
 hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy);
 
 /**
diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
index 2ebd4f8b7d..c1846c1b1e 100644
--- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
@@ -162,6 +162,7 @@ typedef CUdeviceptr hipDeviceptr_t;
 typedef struct cudaArray hipArray;
 typedef struct cudaArray* hipArray_const_t;
 typedef cudaFuncAttributes hipFuncAttributes;
+#define hip_Memcpy2D CUDA_MEMCPY2D
 #define hipMemcpy3DParms cudaMemcpy3DParms
 #define hipArrayDefault cudaArrayDefault
 #define hipArrayLayered cudaArrayLayered
@@ -578,6 +579,10 @@ inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src,
         cudaMemcpy2D(dst, dpitch, src, spitch, width, height, hipMemcpyKindToCudaMemcpyKind(kind)));
 }
 
+inline static hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
+  return hipCUResultTohipError(cuMemcpy2D(pCopy));
+}
+
 inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p)
 {
     return hipCUDAErrorTohipError(cudaMemcpy3D(p));
diff --git a/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
index a6ba44696a..36d37a4fad 100644
--- a/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
+++ b/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
@@ -71,8 +71,8 @@ bool runTest(int argc, char** argv) {
     copyParam.srcMemoryType = hipMemoryTypeHost;
     copyParam.srcHost = hData;
     copyParam.srcPitch = width * sizeof(float);
-    copyParam.widthInBytes = copyParam.srcPitch;
-    copyParam.height = height;
+    copyParam.WidthInBytes = copyParam.srcPitch;
+    copyParam.Height = height;
     hipMemcpyParam2D(&copyParam);
 
     textureReference* texref;
diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp
index 0fad8ab890..36edcdb338 100644
--- a/projects/hip/src/hip_memory.cpp
+++ b/projects/hip/src/hip_memory.cpp
@@ -1715,8 +1715,8 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
     if (pCopy == nullptr) {
         e = hipErrorInvalidValue;
     }
-    e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch,
-                     pCopy->widthInBytes, pCopy->height, hipMemcpyDefault);
+    e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch,
+                     pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault);
     return ihipLogStatus(e);
 }
 
diff --git a/projects/hip/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/projects/hip/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
new file mode 100644
index 0000000000..d5dffd0b09
--- /dev/null
+++ b/projects/hip/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
@@ -0,0 +1,36 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+/*
+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.
+*/
+
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda.h>
+// CHECK-NOT: #include <texture_fetch_functions.h>
+#include <texture_fetch_functions.h>
+
+// CHECK: extern texture<float, 2, hipReadModeElementType> tex;
+extern texture<float, 2, cudaReadModeElementType> tex;
+
+extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) {
+  int x = blockDim.x * blockIdx.x + threadIdx.x;
+  int y = blockDim.y * blockIdx.y + threadIdx.y;
+  outputData[y * width + x] = tex2D(tex, x, y);
+}
diff --git a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code
index 737ee1ce9b..4246151be1 100755
Binary files a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code and b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code differ
diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp
index c365dce70a..8160ac4bc3 100644
--- a/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp
+++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp
@@ -178,7 +178,9 @@ void runTests(int64_t numElements) {
     {
         test(0x01, C_d, C_h, numElements, syncNone, true /*expectMismatch*/);
         test(0x02, C_d, C_h, numElements, syncNullStream, false /*expectMismatch*/);
+#ifndef __HIP_CLANG_ONLY__
         test(0x04, C_d, C_h, numElements, syncOtherStream, true /*expectMismatch*/);
+#endif
         test(0x08, C_d, C_h, numElements, syncDevice, false /*expectMismatch*/);
 
         // Sending a marker to to null stream may synchronize the otherStream