From a0aef11da744037b1be7dd8c9c9bf8549a91bc2f Mon Sep 17 00:00:00 2001
From: Aditya Atluri <adityaavinash1@gmail.com>
Date: Fri, 15 Apr 2016 11:21:45 -0500
Subject: [PATCH 1/2] changed to guard from hc.hpp

---
 include/hcc_detail/hip_runtime.h | 23 ++++++++++-------------
 src/hip_hcc.cpp                  | 15 +++++++++------
 2 files changed, 19 insertions(+), 19 deletions(-)

diff --git a/include/hcc_detail/hip_runtime.h b/include/hcc_detail/hip_runtime.h
index aa420e992d..1a10059eba 100644
--- a/include/hcc_detail/hip_runtime.h
+++ b/include/hcc_detail/hip_runtime.h
@@ -45,10 +45,11 @@ THE SOFTWARE.
 //---
 // Remainder of this file only compiles with HCC
 #ifdef __HCC__
-#if __cplusplus
-#include <hc.hpp>
-#endif
+//#if __cplusplus
+//#include <hc.hpp>
+//#endif
 #include <grid_launch.h>
+#include <kalmar_defines.h>
 extern int HIP_TRACE_API;
 
 //TODO-HCC-GL - change this to typedef.
@@ -110,7 +111,6 @@ extern int HIP_TRACE_API;
 #define __HCC_C__
 #endif
 
-#if __cplusplus
 __device__ float acosf(float x);
 __device__ float acoshf(float x);
 __device__ float asinf(float x);
@@ -290,7 +290,6 @@ __device__ double trunc(double x);
 __device__ double y0(double x);
 __device__ double y1(double y);
 __device__ double yn(int n, double x);
-#endif
 
 // TODO - hipify-clang - change to use the function call.
 //#define warpSize hc::__wavesize()
@@ -491,10 +490,10 @@ __device__ float __dsqrt_rz(double x);
 #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)
 
 #define HIP_KERNEL_NAME(...) __VA_ARGS__
-
+#include<grid_launch.h>
 #ifdef __HCC_CPP__
-hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av);
-void ihipPostLaunchKernel(hipStream_t stream, hc::completion_future &cf);
+hipStream_t ihipPreLaunchKernel(hipStream_t stream, grid_launch_parm *lp);
+void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp);
 
 // TODO - move to common header file.
 #define KNRM  "\x1B[0m"
@@ -510,16 +509,14 @@ do {\
   lp.groupDim.x = _blockDim3D.x; \
   lp.groupDim.y = _blockDim3D.y; \
   lp.groupDim.z = _blockDim3D.z; \
-  lp.groupMemBytes = _groupMemBytes;\
-  hc::completion_future cf;\
-  lp.cf = &cf;  \
-  hipStream_t trueStream = (ihipPreLaunchKernel(_stream, &lp.av)); \
+  lp.groupMemBytes = _groupMemBytes; \
+  hipStream_t trueStream = (ihipPreLaunchKernel(_stream, &lp)); \
     if (HIP_TRACE_API) {\
         fprintf(stderr, KGRN "<<hip-api: hipLaunchKernel '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n" KNRM, \
                 #_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\
     }\
   _kernelName (lp, __VA_ARGS__);\
-  ihipPostLaunchKernel(trueStream, cf);\
+  ihipPostLaunchKernel(trueStream, lp);\
 } while(0)
 
 #else
diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index 41469c5ee1..48639cff33 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -1045,25 +1045,28 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
 // TODO - data-up to data-down:
 // Called just before a kernel is launched from hipLaunchKernel.
 // Allows runtime to track some information about the stream.
-hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av)
+hipStream_t ihipPreLaunchKernel(hipStream_t stream, grid_launch_parm *lp)
 {
 	std::call_once(hip_initialized, ihipInit);
     stream = ihipSyncAndResolveStream(stream);
 
 
     stream->lockopen_preKernelCommand();
-
-    *av = &stream->_av;
-
+//    *av = &stream->_av;
+    lp->av = &stream->_av;
+    lp->cf = new hc::completion_future;
+//    lp->av = static_cast<void*>(av);
+//    lp->cf = static_cast<void*>(malloc(sizeof(hc::completion_future)));
     return (stream);
 }
 
 
 //---
 //Called after kernel finishes execution.
-void ihipPostLaunchKernel(hipStream_t stream, hc::completion_future &kernelFuture)
+void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp)
 {
-    stream->lockclose_postKernelCommand(kernelFuture);
+//    stream->lockclose_postKernelCommand(cf);
+    stream->lockclose_postKernelCommand(*lp.cf);
     if (HIP_LAUNCH_BLOCKING) {
         tprintf(DB_SYNC, " stream:%p LAUNCH_BLOCKING for kernel completion\n", stream);
     }

From 07efbcf0d0c7728a939c8df5926ed287a11332a1 Mon Sep 17 00:00:00 2001
From: Jack Chung <whchung@gmail.com>
Date: Tue, 19 Apr 2016 23:17:09 +0800
Subject: [PATCH 2/2] Fix compilation error in hipSinglePrecisionMathHost test

---
 include/hcc_detail/hip_runtime.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/include/hcc_detail/hip_runtime.h b/include/hcc_detail/hip_runtime.h
index 1a10059eba..a637471995 100644
--- a/include/hcc_detail/hip_runtime.h
+++ b/include/hcc_detail/hip_runtime.h
@@ -33,7 +33,7 @@ THE SOFTWARE.
 
 
 //#include <cstring>
-//#include <cmath>
+#include <cmath>
 #include <string.h>
 #include <stddef.h>