Re-sync with upstream.

src/device_util.cpp

@@ -45,8 +45,8 @@ __device__ void *__hip_hc_malloc(size_t size)
     {
         return (void*)nullptr;
     }
-    uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z;
-    uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
+    uint32_t totalThreads = blockDim.x * gridDim.x * blockDim.y * gridDim.y * blockDim.z * gridDim.z;
+    uint32_t currentWorkItem = threadIdx.x + blockDim.x * blockIdx.x;
 
     uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads;
     uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads;
@@ -932,7 +932,7 @@ __device__  unsigned long long int atomicMax(unsigned long long int* address,
 template<typename T>
 __device__ T atomicCAS_impl(T* address, T compare, T val)
 {
-  // the implementation assumes the atomic is lock-free and 
+  // the implementation assumes the atomic is lock-free and
   // has the same size as the non-atmoic equivalent type
   static_assert(sizeof(T) == sizeof(std::atomic<T>)
                 , "size mismatch between atomic and non-atomic types");
@@ -945,7 +945,7 @@ __device__ T atomicCAS_impl(T* address, T compare, T val)
 
   T expected = compare;
 
-  // hcc should generate a system scope atomic CAS 
+  // hcc should generate a system scope atomic CAS
   std::atomic_compare_exchange_weak_explicit(u.atomic_address
                                             , &expected, val
                                             , std::memory_order_acq_rel
@@ -1110,8 +1110,8 @@ __device__ void* __get_dynamicgroupbaseptr() {
   return hc::get_dynamic_group_segment_base_pointer();
 }
 
-__host__ void* __get_dynamicgroupbaseptr() { 
-  return nullptr; 
+__host__ void* __get_dynamicgroupbaseptr() {
+  return nullptr;
 }
 
 // Precise Math Functions

src/grid_launch.cpp

@@ -91,6 +91,7 @@ namespace hip_impl
         const dim3& dimBlocks,
         uint32_t sharedMemBytes,
         hipStream_t stream,
+<<<<<<< HEAD
         void** kernarg)
     {
         const auto it0 = functions().find(function_address);
@@ -133,6 +134,22 @@ 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
         }
     }
 }

src/hip_context.cpp

@@ -269,7 +269,7 @@ hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig )
 hipError_t hipCtxSynchronize ( void )
 {
     HIP_INIT_API(1);
-    return ihipSynchronize(); //TODP Shall check validity of ctx?
+    return ihipLogStatus(ihipSynchronize()); //TODP Shall check validity of ctx?
 }
 
 hipError_t hipCtxGetFlags ( unsigned int* flags )

src/hip_hcc.cpp

@@ -1592,7 +1592,9 @@ void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launc
     tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n");
 
     stream->lockclose_postKernelCommand(kernelName, lp.av);
-    MARKER_END();
+    if(HIP_PROFILE_API) {
+        MARKER_END();
+    }
 }
 
 //=================================================================================================