more work on async copies

[ROCm/clr commit: 0cdbe1ff05]

projects/clr/hipamd/bin/hipcc

@@ -164,7 +164,7 @@ if ($needHipHcc) {
     if ((not -e $object) or ((stat($source))[9] > (stat($object))[9]))  {
         my $CMD = "$HCC $HCCFLAGS  -I$HSA_PATH/include -I$HIP_PATH/include -Wall -c $source -o $object";
         if ($verbose & 0x10) {
-            $CMD .= " -g" ;
+            $CMD .= " -g -O2" ;
         } else {
             $CMD .= " -O3" ;
         }

projects/clr/hipamd/src/hip_hcc.cpp

@@ -62,12 +62,13 @@ THE SOFTWARE.
 //static const int debug = 0;
 static const int release = 1;
 
+int HIP_LAUNCH_BLOCKING = 0;
+
 int HIP_PRINT_ENV = 0;
 int HIP_TRACE_API= 0;
-int HIP_LAUNCH_BLOCKING = 0;
 int HIP_STAGING_SIZE = 64;   /* size of staging buffers, in KB */
 int HIP_STAGING_BUFFERS = 2;  
-int HIP_STREAM_SIGNALS = 2;  /* number of signals to use when stream is created */
+int HIP_STREAM_SIGNALS = 2;  /* number of signals to allocate at stream creation */
 
 #define TRACE_API   0x1 /* trace API calls and return values */
 #define TRACE_SYNC  0x2 /* trace synchronization pieces */
@@ -128,7 +129,7 @@ public:
 
     inline ihipDevice_t *       getDevice() const;
 
-    hsa_signal_t                getSignal() ;
+    ihipSignal_t *               getSignal() ;
     void                        releaseSignal(ihipSignal_t *signal) ;
 
 private:
@@ -241,6 +242,13 @@ ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, unsig
 {
     _signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1);
 
+    auto s = this;
+
+    std::for_each(_signalPool.begin(), _signalPool.end(), 
+            [s](ihipSignal_t &iter) { 
+            printf ("  stream:%p allocated hsa_signal=%p\n", s, (iter._hsa_signal));
+            });
+
 };
 
 //---
@@ -259,18 +267,18 @@ inline ihipDevice_t * ihipStream_t::getDevice() const
 
 // Allocate a new signal from the signal pool.
 // Returned signals are initialized to a value of "1".
-hsa_signal_t ihipStream_t::getSignal() 
+ihipSignal_t *ihipStream_t::getSignal() 
 {
     int numToScan = _signalPool.size();
     do {
         auto thisCursor = _signalCursor;
-        if (++_signalCursor > _signalPool.size()) {
+        if (++_signalCursor == _signalPool.size()) {
             _signalCursor = 0;
         }
 
         if (_signalPool[thisCursor]._refCnt == 0) {
             _signalPool[thisCursor]._refCnt ++; // allocate it
-            return _signalPool[thisCursor]._hsa_signal;
+            return &_signalPool[thisCursor];
         }
 
         numToScan--;
@@ -336,6 +344,7 @@ void ihipDevice_t::init(unsigned device_index, hc::accelerator acc)
     this->reset();
 };
 
+
 ihipDevice_t::~ihipDevice_t()
 {
     if (_null_stream) {
@@ -628,12 +637,14 @@ void ihipInit()
     /*
      * Environment variables
      */
-    READ_ENV_I(release, HIP_PRINT_ENV, 0,  "Print HIP environment variables.");
-    READ_ENV_I(release, HIP_TRACE_API, 0,  "Trace each HIP API call.  Print function name and return code to stderr as program executes.");
+    READ_ENV_I(release, HIP_PRINT_ENV, 0,  "Print HIP environment variables.");  
+    //-- READ HIP_PRINT_ENV env first, since it has impact on later env var reading
+
     READ_ENV_I(release, HIP_LAUNCH_BLOCKING, CUDA_LAUNCH_BLOCKING, "Make HIP APIs 'host-synchronous', so they block until any kernel launches or data copy commands complete. Alias: CUDA_LAUNCH_BLOCKING." );
+    READ_ENV_I(release, HIP_TRACE_API, 0,  "Trace each HIP API call.  Print function name and return code to stderr as program executes.");
     READ_ENV_I(release, HIP_STAGING_SIZE, 0, "Size of each staging buffer (in KB)" );
     READ_ENV_I(release, HIP_STAGING_BUFFERS, 0, "Number of staging buffers to use in each direction");
-    READ_ENV_I(release, HIP_STREAM_SIGNALS, 0, "Number of signals to use when creating a new stream (pool can later grow)");
+    READ_ENV_I(release, HIP_STREAM_SIGNALS, 0, "Number of signals to allocate when new stream is created (signal pool will grow on demand)");
 
     /*
      * Build a table of valid compute devices.
@@ -791,7 +802,10 @@ inline bool ihipCheckCommandSwitchSync(hipStream_t stream, ihipCommand_t new_com
         addedSync = true;
         *marker = stream->_av.create_marker();
 
-        tprintf (TRACE_SYNC, "stream %p switch to %s (barrier pkt inserted)\n", (void*)stream, new_command == ihipCommandKernel ? "Kernel" : "Data");
+        tprintf (TRACE_SYNC, "stream %p switch %s to %s (barrier pkt inserted)\n", 
+                (void*)stream, 
+                stream->_last_command == ihipCommandKernel ? "Kernel" : "Data",
+                new_command == ihipCommandKernel ? "Kernel" : "Data");
         stream->_last_command = new_command;
     }
 
@@ -1908,10 +1922,12 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
 }
 
 
-//---
-/*
+#if USE_ASYNC_COPY==0
+/**
  * @warning on HCC hipMemcpyAsync uses a synchronous copy.
  */
+#endif
+//---
 hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream)
 {
     std::call_once(hip_initialized, ihipInit);
@@ -1927,9 +1943,6 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
 
     // Async - need to set up dependency on the last command queued to the device?
 
-    // TODO-hsart This routine needs to ensure that dst and src are mapped on the GPU.
-    // This is a synchronous copy - remove and replace with code below when we have appropriate LOCK APIs.
-    hc::am_copy(dst, src, sizeBytes);
 
 #if USE_ASYNC_COPY
 
@@ -1943,25 +1956,33 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
         } else {
             // Let HSA runtime handle it:
             // TODO - need buffer pool for the signals rather than lock:
-            device->_copy_lock[1].lock();
+            ihipSignal_t *ihip_signal = stream->getSignal();
 
-            hsa_signal_store_relaxed(device->_copy_signal, 1);
-            hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, src, sizeBytes, device->_hsa_agent, 0, NULL, device->_copy_signal);
+            //stream->saveLastSignal(ihipSignal);
+
+            hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, src, sizeBytes, device->_hsa_agent, 0, NULL, ihip_signal->_hsa_signal);
 
             if (hsa_status == HSA_STATUS_SUCCESS) {
-                hsa_signal_wait_relaxed(device->_copy_signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); 
+
+                if (HIP_LAUNCH_BLOCKING) {
+                    hsa_signal_wait_relaxed(ihip_signal->_hsa_signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); 
+                    stream->releaseSignal(ihip_signal);
+                } 
+            } else {
+                // This path can be hit if src or dst point to unpinned host memory.
+                // TODO - does async-copy fall back to sync if input pointers are not pinned?
+                e = hipErrorInvalidValue;
             }
-
-            device->_copy_lock[1].unlock();
-
         }
     } else {
         e = hipErrorInvalidValue;
     }
-
+#else
+    // TODO-hsart This routine needs to ensure that dst and src are mapped on the GPU.
+    // This is a synchronous copy - remove and replace with code below when we have appropriate LOCK APIs.
+    hc::am_copy(dst, src, sizeBytes);
 #endif
 
-    // TODO - if am_copy becomes async, and we have HIP_LAUNCH_BLOCKING set, then we would wait for copy operation to complete here.
 
     return ihipLogStatus(e);
 }
@@ -2015,6 +2036,7 @@ hipError_t hipMemsetAsync(void* dst, int  value, size_t sizeBytes, hipStream_t s
 
 hipError_t hipMemset(void* dst, int  value, size_t sizeBytes )
 {
+    // TODO - call an ihip memset so HIP_TRACE is correct.
     return hipMemsetAsync(dst, value, sizeBytes, hipStreamNull);
 }
 

projects/clr/hipamd/tests/src/CMakeLists.txt

@@ -104,6 +104,7 @@ make_hip_executable (hip_brev hip_brev.cpp)
 make_hip_executable (hip_ffs hip_ffs.cpp) 
 make_hip_executable (hipGetDeviceAttribute hipGetDeviceAttribute.cpp) 
 make_hip_executable (hipMemcpy hipMemcpy.cpp) 
+make_hip_executable (hipMemcpyAsync hipMemcpyAsync.cpp) 
 make_hip_executable (hipMemset hipMemset.cpp) 
 make_hip_executable (hipEventRecord hipEventRecord.cpp) 
 make_hip_executable (hipLanguageExtensions hipLanguageExtensions.cpp) 
@@ -131,6 +132,7 @@ make_test(hipGridLaunch " " )
 make_test(hipPointerAttrib " " )
 
 make_test(hipMemcpy  " " )
+make_test(hipMemcpyAsync  " " )
 
 make_test(hipHcc  " " )
 

projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp

@@ -0,0 +1,149 @@
+// Test under-development.  Calls async mem-copy API, experiment with functionality.
+
+#include "hip_runtime.h"
+#include "test_common.h"
+
+unsigned p_streams = 2;
+
+
+void simpleNegTest()
+{
+    printf ("testing: %s\n",__func__);
+    hipError_t e;
+    float *A_malloc, *A_pinned, *A_d;
+
+    size_t Nbytes = N*sizeof(float);
+    A_malloc = (float*)malloc(Nbytes);
+    HIPCHECK(hipMallocHost(&A_pinned, Nbytes));
+    HIPCHECK(hipMalloc(&A_d, Nbytes));
+
+
+    // Can't use default with async copy
+    e = hipMemcpyAsync(A_pinned, A_d, Nbytes, hipMemcpyDefault, NULL);
+    HIPASSERT (e==hipErrorInvalidMemcpyDirection);
+
+
+    // Not sure what happens here, the memory must be pinned.
+    e = hipMemcpyAsync(A_malloc, A_d, Nbytes, hipMemcpyHostToDevice, NULL);
+    HIPASSERT (e==hipErrorInvalidValue);
+
+
+}
+
+//---
+//Classic example showing how to overlap data transfer with compute.
+//We divide the work into "chunks" and create a stream for each chunk.
+//Each chunk then runs a H2D copy, followed by kernel execution, followed by D2H copyback.
+//Work in separate streams is independent which enables concurrency.
+
+// IN: nStreams : number of streams to use for the test
+// IN :useNullStream - use NULL stream.  Synchronizes everything.
+// IN: useSyncMemcpyH2D - use sync memcpy (no overlap) for H2D
+// IN: useSyncMemcpyD2H - use sync memcpy (no overlap) for D2H
+void chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemcpyH2D, bool useSyncMemcpyD2H)
+{
+
+    size_t Nbytes = N*sizeof(int);
+    printf ("testing: %s(useNullStream=%d, useSyncMemcpyH2D=%d, useSyncMemcpyD2H=%d)  ",__func__, useNullStream, useSyncMemcpyH2D, useSyncMemcpyD2H);
+    printf ("Nbytes=%zu (%6.1f MB)\n", Nbytes, (double)(Nbytes)/1024.0/1024.0);
+
+    int *A_d, *B_d, *C_d;
+    int *A_h, *B_h, *C_h;
+
+    HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, true);
+
+
+    unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+
+
+    hipStream_t *stream = (hipStream_t*)malloc(sizeof(hipStream_t) * nStreams);
+    if (useNullStream) { 
+        nStreams = 1;
+        stream[0] = NULL;
+    } else  {
+        for (int i = 0; i < nStreams; ++i) {
+            HIPCHECK (hipStreamCreate(&stream[i]));
+        }
+    }
+
+
+    size_t workLeft = N; 
+    size_t workPerStream = N / nStreams;
+    for (int i = 0; i < nStreams; ++i) {
+        size_t work = (workLeft < workPerStream) ? workLeft : workPerStream;
+        size_t workBytes = work * sizeof(int);
+
+        size_t offset = i*workPerStream;
+
+        if (useSyncMemcpyH2D) {
+            HIPCHECK ( hipMemcpy(&A_d[offset], &A_h[offset], workBytes, hipMemcpyHostToDevice));
+            HIPCHECK ( hipMemcpy(&B_d[offset], &B_h[offset], workBytes, hipMemcpyHostToDevice));
+        } else {
+            HIPCHECK ( hipMemcpyAsync(&A_d[offset], &A_h[offset], workBytes, hipMemcpyHostToDevice, stream[i]));
+            HIPCHECK ( hipMemcpyAsync(&B_d[offset], &B_h[offset], workBytes, hipMemcpyHostToDevice, stream[i]));
+        };
+
+        hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i], &A_d[offset], &B_d[offset], &C_d[offset], work);
+
+        if (useSyncMemcpyD2H) {
+            HIPCHECK ( hipMemcpy(&C_h[offset], &C_d[offset], workBytes, hipMemcpyDeviceToHost));
+        } else {
+            HIPCHECK ( hipMemcpyAsync(&C_h[offset], &C_d[offset], workBytes, hipMemcpyDeviceToHost, stream[i]));
+        }
+    } 
+
+
+    HIPCHECK (hipDeviceSynchronize());
+
+
+    HipTest::checkVectorADD(A_h, B_h, C_h, N);
+
+    HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, true);
+};
+
+
+//---
+//Parse arguments specific to this test.
+void parseMyArguments(int argc, char *argv[])
+{
+    int more_argc = HipTest::parseStandardArguments(argc, argv, false);
+
+    // parse args for this test:
+    for (int i = 1; i < more_argc; i++) {
+        const char *arg = argv[i];
+
+        if (!strcmp(arg, "--streams")) {
+            if (++i >= argc || !HipTest::parseUInt(argv[i], &p_streams)) {
+               failed("Bad streams argument");
+            }
+        } else {
+            failed("Bad argument '%s'", arg);
+        }
+    };
+};
+
+
+
+
+int main(int argc, char *argv[])
+{
+    HipTest::parseStandardArguments(argc, argv, true);
+    parseMyArguments(argc, argv);
+
+
+    printf ("info: set device to %d\n", p_gpuDevice);
+    HIPCHECK(hipSetDevice(p_gpuDevice));
+
+    simpleNegTest();
+
+
+    chunkedAsyncExample(p_streams, true, true, true); // Easy sync version
+    chunkedAsyncExample(p_streams, false, true, true); // Easy sync version
+    chunkedAsyncExample(p_streams, false, false, true); // Some async
+    chunkedAsyncExample(p_streams, false, false, false); // All async
+
+
+
+    passed();
+
+}