From 03d3c6eaed5f12e723f0a7126ffcb506d887a2f2 Mon Sep 17 00:00:00 2001
From: Junjie Bai <bai@in.tum.de>
Date: Tue, 12 Jun 2018 23:53:18 -0700
Subject: [PATCH 1/4] Properly handle (empty) cmake generator expression

---
 cmake/FindHIP.cmake | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/cmake/FindHIP.cmake b/cmake/FindHIP.cmake
index 6668e60332..a45fb1b44d 100644
--- a/cmake/FindHIP.cmake
+++ b/cmake/FindHIP.cmake
@@ -265,7 +265,7 @@ endmacro()
 set(HIP_HIPCC_INCLUDE_ARGS_USER "")
 macro(HIP_INCLUDE_DIRECTORIES)
     foreach(dir ${ARGN})
-        list(APPEND HIP_HIPCC_INCLUDE_ARGS_USER -I${dir})
+        list(APPEND HIP_HIPCC_INCLUDE_ARGS_USER $<$<BOOL:${dir}>:-I${dir}>)
     endforeach()
 endmacro()
 
@@ -373,7 +373,7 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files
     list(REMOVE_DUPLICATES _hip_include_directories)
     if(_hip_include_directories)
         foreach(dir ${_hip_include_directories})
-            list(APPEND HIP_HIPCC_INCLUDE_ARGS -I${dir})
+            list(APPEND HIP_HIPCC_INCLUDE_ARGS $<$<BOOL:${dir}>:-I${dir}>)
         endforeach()
     endif()
 

From 069e2c34c9707a2a4045c823a064a0d35af8269c Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Thu, 14 Jun 2018 11:58:56 +0530
Subject: [PATCH 2/4] Fix stream resolution in memcpy2dasync

---
 src/hip_memory.cpp | 1 +
 1 file changed, 1 insertion(+)

diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 68966abaf1..d6c04ae98c 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -1668,6 +1668,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
     void *pinnedPtr=NULL;
     void *actualSrc = (void*)src;
     void *actualDest = dst;
+    stream = ihipSyncAndResolveStream(stream);
     if(kind == hipMemcpyHostToDevice ) {
         if(getLockedPointer((void*)src, spitch, &pinnedPtr) == hipSuccess ){
             isLocked = 1;

From 0b1e698e7409fead6fcabc38e84eb966bdcd894a Mon Sep 17 00:00:00 2001
From: Nico Trost <nico.trost@amd.com>
Date: Thu, 14 Jun 2018 21:49:54 +0200
Subject: [PATCH 3/4] added missing hipCmul() to nvcc_detail/hip_complex.h

---
 include/hip/nvcc_detail/hip_complex.h | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/include/hip/nvcc_detail/hip_complex.h b/include/hip/nvcc_detail/hip_complex.h
index ea2ddfdf4e..d0e45d26db 100644
--- a/include/hip/nvcc_detail/hip_complex.h
+++ b/include/hip/nvcc_detail/hip_complex.h
@@ -83,6 +83,10 @@ __device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, h
     return cuCsub(p, q);
 }
 
+__device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
+    return cuCmul(p, q);
+}
+
 __device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
     return cuCdiv(p, q);
 }

From cd239058976b1b73da76a8e58f6558f931ebd016 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Fri, 15 Jun 2018 16:08:29 +0530
Subject: [PATCH 4/4] TEMP- fix memcpy2dAsync for trsm issue

---
 src/hip_memory.cpp | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index d6c04ae98c..787e49683b 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -1680,9 +1680,12 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
             actualDest = pinnedPtr;
           }
     }
+#if 0
     if((width == dpitch) && (width == spitch)) {
             hip_internal::memcpyAsync(dst, src, width*height, kind, stream);
-    } else {
+    } else
+#endif
+   {
         try {
             if(!isLocked){
                 for (int i = 0; i < height; ++i)