From 7d5a45ac1a953f3a7958fbc138ec792834fcd64d Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Mon, 20 Nov 2017 16:33:52 +0000
Subject: [PATCH] Correct ill-formed merge in earlier commit and adjust for
 differences with the new CUDA natural indexing mechanism.

---
 include/hip/hcc_detail/hip_runtime.h | 15 ++++++++++++---
 src/grid_launch.cpp                  | 17 -----------------
 2 files changed, 12 insertions(+), 20 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index 924e774af0..c2302d4dc1 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -389,9 +389,18 @@ template<typename std::common_type<
 class Coordinates {
     using R = decltype(f(0));
 
-    struct X { __device__ operator R() const { return f(0); } };
-    struct Y { __device__ operator R() const { return f(1); } };
-    struct Z { __device__ operator R() const { return f(2); } };
+    struct X {
+        __device__ operator R() const { return f(0); }
+        __host__ operator R() const { return R{}; }
+    };
+    struct Y {
+        __device__ operator R() const { return f(1); }
+        __host__ operator R() const { return R{}; }
+    };
+    struct Z {
+        __device__ operator R() const { return f(2); }
+        __host__ operator R() const { return R{}; }
+    };
 public:
     static constexpr X x{};
     static constexpr Y y{};
diff --git a/src/grid_launch.cpp b/src/grid_launch.cpp
index 3fec390720..4a26f66c8c 100644
--- a/src/grid_launch.cpp
+++ b/src/grid_launch.cpp
@@ -91,7 +91,6 @@ namespace hip_impl
         const dim3& dimBlocks,
         uint32_t sharedMemBytes,
         hipStream_t stream,
-<<<<<<< HEAD
         void** kernarg)
     {
         const auto it0 = functions().find(function_address);
@@ -134,22 +133,6 @@ namespace hip_impl
                     nullptr,
                     kernarg);
             }
-=======
-        void* locked_stream,
-        const char* kernel_name,
-        hc::accelerator_view* acc_v)
-    {   // Precondition: acc_v is the accelerator_view associated with stream
-        //               which is guarded by locked_stream;
-        //               locked_stream is deletable.
-        using L = decltype(stream->lockopen_preKernelCommand());
-
-        stream->lockclose_postKernelCommand(kernel_name, acc_v);
-
-        delete static_cast<L*>(locked_stream);
-        locked_stream = nullptr;
-        if(HIP_PROFILE_API) {
-            MARKER_END();
->>>>>>> e8ede28ec4f5744185b171031e537237afb7affa
         }
     }
 }