From 6856d28ca3b95e9cc16ef7f1c8ab01bc789085e1 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 17 Feb 2016 00:59:12 -0600 Subject: [PATCH] more work on async copies [ROCm/clr commit: 0cdbe1ff0591b3cc10bdb34500936496bc6dca74] --- projects/clr/hipamd/bin/hipcc | 2 +- projects/clr/hipamd/src/hip_hcc.cpp | 70 +++++--- projects/clr/hipamd/tests/src/CMakeLists.txt | 2 + .../clr/hipamd/tests/src/hipMemcpyAsync.cpp | 149 ++++++++++++++++++ 4 files changed, 198 insertions(+), 25 deletions(-) create mode 100644 projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index 7537750ff6..1ab4cf2759 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/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" ; } diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index ff1f39d780..ae28947ef3 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/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); } diff --git a/projects/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt index 7e4736a99a..ec0b15ad62 100644 --- a/projects/clr/hipamd/tests/src/CMakeLists.txt +++ b/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 " " ) diff --git a/projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp b/projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp new file mode 100644 index 0000000000..e2968af2f2 --- /dev/null +++ b/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(); + +}