From d24ba1ecbb704091e15e16e882d5b0ebab370b78 Mon Sep 17 00:00:00 2001
From: bwicakso <besarwicaksono@gmail.com>
Date: Wed, 20 Apr 2016 15:51:39 -0500
Subject: [PATCH 1/2] Fix for kernel synchronization

The completion future of a particular kernel is lost if there are
multiple kernels in the stream. This can cause a racing condition where
the signal associated with the unreferenced completion_future might get
released by hcc runtime.


[ROCm/hip commit: df98fd8531e7b099b88bb0f0f9bdbad33abd64e9]
---
 projects/hip/include/hcc_detail/hip_hcc.h | 6 +++++-
 projects/hip/src/hip_hcc.cpp              | 5 +++++
 2 files changed, 10 insertions(+), 1 deletion(-)

diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h
index 58b58e2c99..5fa9be0fbc 100644
--- a/projects/hip/include/hcc_detail/hip_hcc.h
+++ b/projects/hip/include/hcc_detail/hip_hcc.h
@@ -440,9 +440,13 @@ public:
     hc::accelerator_view        _av;
     unsigned                    _flags;
 
-private: // Critical Data.  THis MUST be accessed through LockedAccessor_StreamCrit_t
+private:
+    // Critical Data.  THis MUST be accessed through LockedAccessor_StreamCrit_t
     ihipStreamCritical_t        _criticalData;
 
+    // Array of dependency completion_future.
+    std::vector<hc::completion_future> _depFutures;
+
 private:
     void                        enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal);
     void                        waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal);
diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 44cccf14fa..64d529f875 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -181,6 +181,8 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty
     // Reset the stream to "empty" - next command will not set up an inpute dependency on any older signal.
     crit->_last_command_type = ihipCommandCopyH2D;
     crit->_last_copy_signal = NULL;
+
+    _depFutures.clear();
 }
 
 
@@ -428,6 +430,9 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t
             needSync = 1;
             hsa_signal_t *hsaSignal = (static_cast<hsa_signal_t*> (crit->_last_kernel_future.get_native_handle()));
             if (hsaSignal) {
+                // Keep reference to the kernel future in order to keep the
+                // dependent signal alive.
+                _depFutures.push_back(crit->_last_kernel_future);
                 *waitSignal = * hsaSignal;
             } else {
                 assert(0); // if NULL signal, and we return 1, hsa_amd_memory_copy_async will fail.  Confirm this never happens.

From dcd09cebcaf0036289ccdf112b2092255bc92050 Mon Sep 17 00:00:00 2001
From: bwicakso <besarwicaksono@gmail.com>
Date: Mon, 25 Apr 2016 14:42:35 -0500
Subject: [PATCH 2/2] Update with original

[ROCm/hip commit: c976e6bcd330b183ed4e544805aa4097ccc2eef5]
---
 projects/hip/include/hcc_detail/hip_ldg.h |   1 +
 projects/hip/tests/src/hip_test_ldg.cpp   | 154 +++++++++++++++++++++-
 2 files changed, 149 insertions(+), 6 deletions(-)

diff --git a/projects/hip/include/hcc_detail/hip_ldg.h b/projects/hip/include/hcc_detail/hip_ldg.h
index e6fe45cd77..4dab90b4e8 100644
--- a/projects/hip/include/hcc_detail/hip_ldg.h
+++ b/projects/hip/include/hcc_detail/hip_ldg.h
@@ -102,3 +102,4 @@ __device__ double4              __ldg(const double4* );
 #endif  // __HCC__
 
 #endif  // HIP_LDG_H
+
diff --git a/projects/hip/tests/src/hip_test_ldg.cpp b/projects/hip/tests/src/hip_test_ldg.cpp
index a58652d240..31c7d4865a 100644
--- a/projects/hip/tests/src/hip_test_ldg.cpp
+++ b/projects/hip/tests/src/hip_test_ldg.cpp
@@ -26,7 +26,7 @@ THE SOFTWARE.
 #include<iostream>
 #include "hip_runtime.h"
 #include "test_common.h"
-
+#if __hcc_workweek__ >= 16164
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 
 
@@ -149,14 +149,156 @@ int main() {
   cout << " System major " << devProp.major << endl;
   cout << " agent prop name " << devProp.name << endl;
 
-   int errors = dataTypesRun<char>() &
-           dataTypesRun<signed char>() &
-           dataTypesRun<short>() &
-           dataTypesRun<int>();
-  //hipResetDefaultAccelerator();
+    int errors;
+
+    errors =    dataTypesRun<char>() &
+                dataTypesRun<char1>() &
+                dataTypesRun<char2>() &
+                dataTypesRun<char3>() &
+                dataTypesRun<char4>() &
+                dataTypesRun<signed char>() &
+                dataTypesRun<unsigned char>();
+
+    if(errors == 1){
+        errors = 0;
+    }else{
+        std::cout<<"Failed Char"<<std::endl;
+        return -1;
+    }
+
+    errors =    dataTypesRun<short>() &
+                dataTypesRun<short1>() &
+                dataTypesRun<short2>() &
+                dataTypesRun<short3>() &
+                dataTypesRun<short4>() &
+                dataTypesRun<unsigned short>();
+
+    if(errors == 1){
+        errors = 0;
+    }else{
+        std::cout<<"Failed Short"<<std::endl;
+        return -1;
+    }
+
+    errors =    dataTypesRun<int>() &
+                dataTypesRun<int1>() &
+                dataTypesRun<int2>() &
+                dataTypesRun<int3>() &
+                dataTypesRun<int4>() &
+                dataTypesRun<unsigned int>();
+
+    if(errors == 1){
+        errors = 0;
+    }else{
+        std::cout<<"Failed Int"<<std::endl;
+        return -1;
+    }
+
+    errors =    dataTypesRun<long>() &
+                dataTypesRun<long1>() &
+                dataTypesRun<long2>() &
+                dataTypesRun<long3>() &
+                dataTypesRun<long4>() &
+                dataTypesRun<unsigned long>();
+
+    if(errors == 1){
+        errors = 0;
+    }else{
+        std::cout<<"Failed Long"<<std::endl;
+        return -1;
+    }
+
+    errors =    dataTypesRun<long long>() &
+                dataTypesRun<longlong1>() &
+                dataTypesRun<longlong2>() &
+                dataTypesRun<longlong3>() &
+                dataTypesRun<longlong4>() &
+                dataTypesRun<unsigned long long>();
+
+    if(errors == 1){
+        errors = 0;
+    }else{
+        std::cout<<"Failed Long Long"<<std::endl;
+        return -1;
+    }
+
+    errors =    dataTypesRun<uchar1>() &
+                dataTypesRun<uchar2>() &
+                dataTypesRun<uchar3>() &
+                dataTypesRun<uchar4>();
+
+    if(errors == 1){
+        errors = 0;
+    }else{
+        std::cout<<"Failed Unsigned Char"<<std::endl;
+        return -1;
+    }
+
+    errors =    dataTypesRun<ushort1>() &
+                dataTypesRun<ushort2>() &
+                dataTypesRun<ushort3>() &
+                dataTypesRun<ushort4>();
+
+    if(errors == 1){
+        errors = 0;
+    }else{
+        std::cout<<"Failed Unsigned Short"<<std::endl;
+        return -1;
+    }
+
+    errors =    dataTypesRun<uint1>() &
+                dataTypesRun<uint2>() &
+                dataTypesRun<uint3>() &
+                dataTypesRun<uint4>();
+
+    if(errors == 1){
+        errors = 0;
+    }else{
+        std::cout<<"Failed Unsigned Int"<<std::endl;
+        return -1;
+    }
+
+    errors =    dataTypesRun<ulonglong1>() &
+                dataTypesRun<ulonglong2>() &
+                dataTypesRun<ulonglong3>() &
+                dataTypesRun<ulonglong4>();
+
+    if(errors == 1){
+        errors = 0;
+    }else{
+        std::cout<<"Failed Unsigned Long Long"<<std::endl;
+        return -1;
+    }
+
+    errors =    dataTypesRun<float>() &
+                dataTypesRun<float1>() &
+                dataTypesRun<float2>() &
+                dataTypesRun<float3>() &
+                dataTypesRun<float4>();
+
+    if(errors == 1){
+        errors = 0;
+    }else{
+        std::cout<<"Failed Float"<<std::endl;
+        return -1;
+    }
+
+    errors =    dataTypesRun<double>() &
+                dataTypesRun<double1>() &
+                dataTypesRun<double2>() &
+                dataTypesRun<double3>() &
+                dataTypesRun<double4>();
+
+
+ //hipResetDefaultAccelerator();
     if(errors == 1){
         passed();
         return 0;
+    }else{
+        std::cout<<"Failed Float"<<std::endl;
+        return -1;
     }
 
 }
+
+#endif