diff --git a/CHANGELOG.md b/CHANGELOG.md index e5c22af8..5d745b00 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,10 @@ Documentation for TransferBench is available at [https://rocm.docs.amd.com/projects/TransferBench](https://rocm.docs.amd.com/projects/TransferBench). +## v1.69.00 +### Fixed +- Fix for non-zero byte offsets used with DMA executor + ## v1.68.00 ### Fixed - Improper draining of writes that could artificially inflate transfer timing diff --git a/src/client/Client.cpp b/src/client/Client.cpp index 3074f3b1..95a0f9a5 100644 --- a/src/client/Client.cpp +++ b/src/client/Client.cpp @@ -259,7 +259,7 @@ void DisplayVersion() #ifndef TB_GIT_COMMIT #define TB_GIT_COMMIT "unknown" #endif - Print("TransferBench v%s.%s (%s:%s)%s%s\n", VERSION, CLIENT_VERSION, TB_GIT_BRANCH, TB_GIT_COMMIT, support.c_str(), multiNodeMode.c_str()); + Print("TransferBench v%s.%sC (%s:%s)%s%s\n", VERSION, CLIENT_VERSION, TB_GIT_BRANCH, TB_GIT_COMMIT, support.c_str(), multiNodeMode.c_str()); Print("=============================================================================================================\n"); } diff --git a/src/client/EnvVars.hpp b/src/client/EnvVars.hpp index 34940b69..6fd27d22 100644 --- a/src/client/EnvVars.hpp +++ b/src/client/EnvVars.hpp @@ -79,7 +79,7 @@ class EnvVars int useInteractive; // Pause for user-input before starting transfer loop // Data options - int alwaysValidate; // Validate after each iteration instead of once after all iterations + int alwaysValidate; // <0 = disable validation, 0 = validate once after all iterations, >0 = validate after each iteration int blockBytes; // Each subexecutor, except the last, gets a multiple of this many bytes to copy int byteOffset; // Byte-offset for memory allocations vector fillPattern; // Pattern of floats used to fill source data @@ -149,7 +149,7 @@ class EnvVars else if (archName == "gfx942") defaultGfxUnroll = 4; else if (archName == "gfx950") defaultGfxUnroll = 4; - alwaysValidate = GetEnvVar("ALWAYS_VALIDATE" , 0); + alwaysValidate = GetEnvVar("ALWAYS_VALIDATE", 0); blockBytes = GetEnvVar("BLOCK_BYTES" , 256); byteOffset = GetEnvVar("BYTE_OFFSET" , 0); fillCompress = GetEnvVarArray("FILL_COMPRESS" , {}); @@ -352,7 +352,7 @@ class EnvVars { printf("Environment variables (client):\n"); printf("======================\n"); - printf(" ALWAYS_VALIDATE - Validate after each iteration instead of once after all iterations\n"); + printf(" ALWAYS_VALIDATE - Data validation mode: <0=disabled, 0=validate once after all iterations (default), >0=validate after each iteration\n"); printf(" BLOCK_BYTES - Controls granularity of how work is divided across subExecutors\n"); printf(" BYTE_OFFSET - Initial byte-offset for memory allocations. Must be multiple of 4\n"); printf(" CU_MASK - CU mask for streams. Can specify ranges e.g '5,10-12,14'\n"); @@ -454,7 +454,8 @@ class EnvVars if (hideEnv) return; Print("ALWAYS_VALIDATE", alwaysValidate, - "Validating after %s", (alwaysValidate ? "each iteration" : "all iterations")); + "Validation %s", (alwaysValidate < 0 ? "disabled" : + alwaysValidate > 0 ? "after each iteration" : "after all iterations")); Print("BLOCK_BYTES", blockBytes, "Each CU gets a mulitple of %d bytes to copy", blockBytes); Print("BYTE_OFFSET", byteOffset, diff --git a/src/client/Utilities.hpp b/src/client/Utilities.hpp index 9730ace6..dafa7541 100644 --- a/src/client/Utilities.hpp +++ b/src/client/Utilities.hpp @@ -551,6 +551,8 @@ namespace TransferBench::Utils if (!RankDoesOutput()) return; if (!ev.outputToCsv) printf("Test %d:\n", testNum); + if (ev.alwaysValidate < 0 && !ev.outputToCsv && !ev.useInteractive) + printf("Validation disabled (ALWAYS_VALIDATE < 0)\n"); bool isMultiRank = TransferBench::GetNumRanks() > 1; diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index b2f3ffdc..b8438d43 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -30,10 +30,13 @@ THE SOFTWARE. #include #include #include +#include #include #include #include +#include #include +#include #include #include #include @@ -92,7 +95,7 @@ namespace TransferBench using std::set; using std::vector; - constexpr char VERSION[] = "1.68"; + constexpr char VERSION[] = "1.69"; /** * Enumeration of supported Executor types @@ -217,7 +220,7 @@ namespace TransferBench */ struct DataOptions { - int alwaysValidate = 0; ///< Validate after each iteration instead of once at end + int alwaysValidate = 0; ///< <0 = disable validation, 0 = validate once at end, >0 = validate after each iteration int blockBytes = 256; ///< Each subexecutor works on a multiple of this many bytes int byteOffset = 0; ///< Byte-offset for memory allocations vector fillPattern = {}; ///< Pattern of floats used to fill source data @@ -5401,11 +5404,106 @@ static bool IsConfiguredGid(union ibv_gid const& gid) return ERR_NONE; } + // Persistent worker-thread pool. + + class ThreadPool { + public: + explicit ThreadPool(size_t numThreads) : stop_(false) { + for (size_t i = 0; i < numThreads; ++i) { + workers_.emplace_back([this] { + for (;;) { + std::function task; + { + std::unique_lock lock(mutex_); + cv_.wait(lock, [this] { return stop_ || !tasks_.empty(); }); + if (stop_ && tasks_.empty()) return; + task = std::move(tasks_.front()); + tasks_.pop(); + } + task(); + } + }); + } + } + + ThreadPool(ThreadPool const&) = delete; + ThreadPool& operator=(ThreadPool const&) = delete; + + ~ThreadPool() { + { + std::lock_guard lock(mutex_); + stop_ = true; + } + cv_.notify_all(); + for (auto& w : workers_) + w.join(); + } + + size_t size() const { return workers_.size(); } + + std::future submit(std::function fn) { + auto task = std::make_shared>(std::move(fn)); + std::future fut = task->get_future(); + { + std::lock_guard lock(mutex_); + tasks_.emplace([task] { (*task)(); }); + } + cv_.notify_one(); + return fut; + } + + private: + std::vector workers_; + std::queue> tasks_; + std::mutex mutex_; + std::condition_variable cv_; + bool stop_; + }; + + // Pools are created once per RunTransfers call (sized from the run's executor/transfer counts) + // and destroyed when it returns. Two separate pools avoid a nested-blocking deadlock: executor + // workers submit their transfers to the transfer pool and block on the results, so the two + // levels must never share threads. Pools are passed explicitly through the call chain to avoid + // global state (reentrancy/thread-safety). + + // Runs "count" units of work. Falls to serial execution if the pool is unavailable/empty. + // Keeps the most severe error across all units (ERR_FATAL > ERR_WARN > ERR_NONE). + template + static inline ErrResult ParallelRun(ThreadPool* pool, size_t count, MakeTask makeTask) { + if (count == 0) return ERR_NONE; + + bool const usePool = (pool != nullptr && pool->size() > 0 && count > 1); + + std::vector> futures; + if (usePool) { + futures.reserve(count - 1); + for (size_t i = 1; i < count; ++i) + futures.emplace_back(pool->submit(makeTask(i))); + } + + // Unit 0 always runs inline on the calling thread. + ErrResult result = makeTask(0)(); + + if (usePool) { + for (auto& f : futures) { + ErrResult const r = f.get(); + if (r.errType > result.errType) result = r; + } + } else { + for (size_t i = 1; i < count; ++i) { + ErrResult const r = makeTask(i)(); + if (r.errType > result.errType) result = r; + } + } + return result; + } + // Execute a single GPU executor static ErrResult RunGpuExecutor(int const iteration, ConfigOptions const& cfg, int const exeIndex, - ExeInfo& exeInfo) + ExeInfo& exeInfo, + ThreadPool* transferPool) { auto cpuStart = std::chrono::high_resolution_clock::now(); ERR_CHECK(hipSetDevice(exeIndex)); @@ -5413,24 +5511,22 @@ static bool IsConfiguredGid(union ibv_gid const& gid) int xccDim = exeInfo.useSubIndices ? exeInfo.numSubIndices : 1; if (cfg.gfx.useMultiStream) { - // Launch one thread per Transfer in separate streams - vector> asyncTransfers; - for (int i = 0; i < exeInfo.streams.size(); i++) { - asyncTransfers.emplace_back(std::async(std::launch::async, - ExecuteGpuTransfer, - iteration, - exeInfo.totalSubExecs, - exeInfo.subExecParamGpu, - exeInfo.streams[i], - cfg.gfx.useHipEvents ? exeInfo.startEvents[i] : NULL, - cfg.gfx.useHipEvents ? exeInfo.stopEvents[i] : NULL, - xccDim, - std::cref(cfg), - exeInfo.gfxKernelToUse, - std::ref(exeInfo.resources[i]))); - } - for (auto& asyncTransfer : asyncTransfers) - ERR_CHECK(asyncTransfer.get()); + // Launch one Transfer per stream + ERR_CHECK(ParallelRun(transferPool, exeInfo.streams.size(), + [&](size_t i) -> std::function { + return [&, i] { + return ExecuteGpuTransfer(iteration, + exeInfo.totalSubExecs, + exeInfo.subExecParamGpu, + exeInfo.streams[i], + cfg.gfx.useHipEvents ? exeInfo.startEvents[i] : NULL, + cfg.gfx.useHipEvents ? exeInfo.stopEvents[i] : NULL, + xccDim, + cfg, + exeInfo.gfxKernelToUse, + exeInfo.resources[i]); + }; + })); } else { // Launch all Transfers in one kernel launch (avoid extra thread creation) ExecuteGpuTransfer(iteration, exeInfo.totalSubExecs, exeInfo.subExecParamGpu, exeInfo.streams[0], @@ -5507,6 +5603,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) int numDsts = (int)resources.dstMem.size(); ERR_CHECK(hipSetDevice(exeIndex)); int subIterations = 0; + size_t const initOffset = cfg.data.byteOffset / sizeof(float); + float* const src = resources.srcMem[0] + initOffset; if (!useSubIndices && !cfg.dma.useHsaCopy) { if (cfg.dma.useHipEvents) ERR_CHECK(hipEventRecord(startEvent, stream)); @@ -5520,12 +5618,13 @@ static bool IsConfiguredGid(union ibv_gid const& gid) do { // Queue for each output location for (int dstIdx = 0; dstIdx < numDsts; dstIdx++) { + float* const dst = resources.dstMem[dstIdx] + initOffset; #if defined(CUMEM_ENABLED) - ERR_CHECK(cuMemcpyAsync((CUdeviceptr)resources.dstMem[dstIdx], - (CUdeviceptr)resources.srcMem[0], + ERR_CHECK(cuMemcpyAsync((CUdeviceptr)dst, + (CUdeviceptr)src, resources.numBytes, stream)); #else - ERR_CHECK(hipMemcpyAsync(resources.dstMem[dstIdx], resources.srcMem[0], resources.numBytes, + ERR_CHECK(hipMemcpyAsync(dst, src, resources.numBytes, memcpyKind, stream)); #endif } @@ -5542,14 +5641,15 @@ static bool IsConfiguredGid(union ibv_gid const& gid) do { hsa_signal_store_screlease(resources.signal, numDsts); for (int dstIdx = 0; dstIdx < numDsts; dstIdx++) { + float* const dst = resources.dstMem[dstIdx] + initOffset; if (!useSubIndices) { - ERR_CHECK(hsa_amd_memory_async_copy(resources.dstMem[dstIdx], resources.dstAgent[dstIdx], - resources.srcMem[0], resources.srcAgent, + ERR_CHECK(hsa_amd_memory_async_copy(dst, resources.dstAgent[dstIdx], + src, resources.srcAgent, resources.numBytes, 0, NULL, resources.signal)); } else { - HSA_CALL(hsa_amd_memory_async_copy_on_engine(resources.dstMem[dstIdx], resources.dstAgent[dstIdx], - resources.srcMem[0], resources.srcAgent, + HSA_CALL(hsa_amd_memory_async_copy_on_engine(dst, resources.dstAgent[dstIdx], + src, resources.srcAgent, resources.numBytes, 0, NULL, resources.signal, resources.sdmaEngineId, true)); @@ -5583,27 +5683,25 @@ static bool IsConfiguredGid(union ibv_gid const& gid) static ErrResult RunDmaExecutor(int const iteration, ConfigOptions const& cfg, int const exeIndex, - ExeInfo& exeInfo) + ExeInfo& exeInfo, + ThreadPool* transferPool) { auto cpuStart = std::chrono::high_resolution_clock::now(); ERR_CHECK(hipSetDevice(exeIndex)); - vector> asyncTransfers; - for (int i = 0; i < exeInfo.resources.size(); i++) { - asyncTransfers.emplace_back(std::async(std::launch::async, - ExecuteDmaTransfer, - iteration, - exeInfo.useSubIndices, - exeIndex, - exeInfo.streams[i], - cfg.dma.useHipEvents ? exeInfo.startEvents[i] : NULL, - cfg.dma.useHipEvents ? exeInfo.stopEvents[i] : NULL, - std::cref(cfg), - std::ref(exeInfo.resources[i]))); - } - - for (auto& asyncTransfer : asyncTransfers) - ERR_CHECK(asyncTransfer.get()); + ERR_CHECK(ParallelRun(transferPool, exeInfo.resources.size(), + [&](size_t i) -> std::function { + return [&, i] { + return ExecuteDmaTransfer(iteration, + exeInfo.useSubIndices, + exeIndex, + exeInfo.streams[i], + cfg.dma.useHipEvents ? exeInfo.startEvents[i] : NULL, + cfg.dma.useHipEvents ? exeInfo.stopEvents[i] : NULL, + cfg, + exeInfo.resources[i]); + }; + })); auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; double deltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 / cfg.general.numSubIterations; @@ -5670,26 +5768,24 @@ static bool IsConfiguredGid(union ibv_gid const& gid) static ErrResult RunBmaExecutor(int const iteration, ConfigOptions const& cfg, int const exeIndex, - ExeInfo& exeInfo) + ExeInfo& exeInfo, + ThreadPool* transferPool) { auto cpuStart = std::chrono::high_resolution_clock::now(); ERR_CHECK(hipSetDevice(exeIndex)); - vector> asyncTransfers; - for (int i = 0; i < exeInfo.resources.size(); i++) { - asyncTransfers.emplace_back(std::async(std::launch::async, - ExecuteBatchDmaTransfer, - iteration, - exeIndex, - exeInfo.streams[i], - cfg.dma.useHipEvents ? exeInfo.startEvents[i] : NULL, - cfg.dma.useHipEvents ? exeInfo.stopEvents[i] : NULL, - std::cref(cfg), - std::ref(exeInfo.resources[i]))); - } - - for (auto& asyncTransfer : asyncTransfers) - ERR_CHECK(asyncTransfer.get()); + ERR_CHECK(ParallelRun(transferPool, exeInfo.resources.size(), + [&](size_t i) -> std::function { + return [&, i] { + return ExecuteBatchDmaTransfer(iteration, + exeIndex, + exeInfo.streams[i], + cfg.dma.useHipEvents ? exeInfo.startEvents[i] : NULL, + cfg.dma.useHipEvents ? exeInfo.stopEvents[i] : NULL, + cfg, + exeInfo.resources[i]); + }; + })); auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; double deltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 / cfg.general.numSubIterations; @@ -5704,17 +5800,18 @@ static bool IsConfiguredGid(union ibv_gid const& gid) static ErrResult RunExecutor(int const iteration, ConfigOptions const& cfg, ExeDevice const& exeDevice, - ExeInfo& exeInfo) + ExeInfo& exeInfo, + ThreadPool* transferPool) { switch (exeDevice.exeType) { case EXE_CPU: return RunCpuExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); - case EXE_GPU_GFX: return RunGpuExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); - case EXE_GPU_DMA: return RunDmaExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); + case EXE_GPU_GFX: return RunGpuExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo, transferPool); + case EXE_GPU_DMA: return RunDmaExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo, transferPool); #ifdef NIC_EXEC_ENABLED case EXE_NIC: return RunNicExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); #endif #ifdef BMA_EXEC_ENABLED - case EXE_GPU_BDMA: return RunBmaExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); + case EXE_GPU_BDMA: return RunBmaExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo, transferPool); #endif default: return {ERR_FATAL, "Unsupported executor (%d)", exeDevice.exeType}; } @@ -5874,24 +5971,30 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } - // Prepare reference src/dst arrays - only once for largest size + // Prepare reference src/dst arrays - only once for largest size. + // dstReference (expected results) is only needed when validation is enabled. + bool const validateEnabled = (cfg.data.alwaysValidate >= 0); size_t maxN = maxNumBytes / sizeof(float); - vector outputBuffer(maxN); - vector> dstReference(maxNumSrcs + 1, vector(maxN)); + vector outputBuffer(validateEnabled ? maxN : 0); + vector> dstReference; { size_t initOffset = cfg.data.byteOffset / sizeof(float); vector> srcReference(maxNumSrcs, vector(maxN)); - memset(dstReference[0].data(), MEMSET_CHAR, maxNumBytes); + if (validateEnabled) { + dstReference.assign(maxNumSrcs + 1, vector(maxN)); + memset(dstReference[0].data(), MEMSET_CHAR, maxNumBytes); + } for (int numSrcs = 0; numSrcs < maxNumSrcs; numSrcs++) { PrepareReference(cfg, srcReference[numSrcs], numSrcs); - for (int i = 0; i < maxN; i++) { - dstReference[numSrcs+1][i] = (numSrcs == 0 ? 0 : dstReference[numSrcs][i]) + srcReference[numSrcs][i]; - } + if (validateEnabled) + for (int i = 0; i < maxN; i++) + dstReference[numSrcs+1][i] = (numSrcs == 0 ? 0 : dstReference[numSrcs][i]) + srcReference[numSrcs][i]; } // Release un-used partial sums - for (int numSrcs = 0; numSrcs < minNumSrcs; numSrcs++) - dstReference[numSrcs].clear(); + if (validateEnabled) + for (int numSrcs = 0; numSrcs < minNumSrcs; numSrcs++) + dstReference[numSrcs].clear(); // Initialize all src memory buffers (if on local rank) bool const verbose = System::Get().IsVerbose(); @@ -5968,6 +6071,18 @@ static bool IsConfiguredGid(union ibv_gid const& gid) System::Get().Barrier(); } + // Create persistent worker pools once for this run (see ThreadPool). Sized so every executor + // and every transfer can run concurrently without the two levels sharing threads: + // - executor pool: one worker per local executor (unit 0 runs inline, so N-1 are farmed) + // - transfer pool: one worker per local transfer across all executors + size_t numLocalExecutors = localExecutors.size(); + size_t numLocalTransfers = 0; + for (auto const& exeDevice : localExecutors) + numLocalTransfers += executorMap[exeDevice].resources.size(); + + ThreadPool executorPool(numLocalExecutors); + ThreadPool transferPool(numLocalTransfers); + // Perform iterations size_t numTimedIterations = 0; double totalCpuTimeSec = 0.0; @@ -5986,20 +6101,18 @@ static bool IsConfiguredGid(union ibv_gid const& gid) // Start CPU timing for this iteration auto cpuStart = std::chrono::high_resolution_clock::now(); - // Execute all Transfers in parallel - std::vector> asyncExecutors; - for (auto const& exeDevice : localExecutors) { - asyncExecutors.emplace_back(std::async(std::launch::async, RunExecutor, - iteration, - std::cref(cfg), - std::cref(exeDevice), - std::ref(executorMap[exeDevice]))); - } - - // Wait for all threads to finish - for (auto& asyncExecutor : asyncExecutors) { - ERR_APPEND(asyncExecutor.get(), errResults); - } + // Execute all executors in parallel via the persistent executor pool. Each ExeInfo& is + // resolved on the calling thread (inside makeTask) to keep std::map access single-threaded, + // matching the previous behaviour. + ERR_APPEND(ParallelRun(&executorPool, localExecutors.size(), + [&](size_t i) -> std::function { + ExeDevice const& exeDevice = localExecutors[i]; + ExeInfo& exeInfo = executorMap[exeDevice]; + return [&cfg, &exeDevice, &exeInfo, &transferPool, iteration] { + return RunExecutor(iteration, cfg, exeDevice, exeInfo, &transferPool); + }; + }), + errResults); // Wait for all ranks to finish System::Get().Barrier(); @@ -6008,7 +6121,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; double deltaSec = std::chrono::duration_cast>(cpuDelta).count() / cfg.general.numSubIterations; - if (cfg.data.alwaysValidate) { + if (cfg.data.alwaysValidate > 0) { ERR_APPEND(ValidateAllTransfers(cfg, transfers, transferResources, dstReference, outputBuffer), errResults); } @@ -6019,11 +6132,32 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } - // Pause for interactive mode - run validation and show per-transfer result before prompt - if (cfg.general.useInteractive) { + // Interactive per-transfer PASS/FAIL display (skipped when validation disabled). This does its + // own comparison pass to show results; authoritative error reporting still comes from + // ValidateAllTransfers (mode 0) or the inline per-iteration validation (mode >0). + if (cfg.general.useInteractive && cfg.data.alwaysValidate >= 0) { + bool inputOk = true; if (localRank == 0) { - System::Get().Log("Transfers complete. Validation results:\n"); + if (cfg.data.alwaysValidate == 0) { + System::Get().Log("Transfers complete. Hit to run validation: "); + fflush(stdout); + if (getchar() == EOF) { + System::Get().Log("[ERROR] Unexpected EOF while waiting for input\n"); + inputOk = false; + } else { + System::Get().Log("\nValidation results:\n"); + } + } else { + System::Get().Log("Transfers complete.\nValidation results:\n"); + } + } + + // Coordinate any EOF abort so no rank is left waiting at the Barrier below + System::Get().Broadcast(0, sizeof(inputOk), &inputOk); + if (!inputOk) + ERR_APPEND((ErrResult{ERR_FATAL, "Unexpected EOF while waiting for interactive input"}), errResults); + if (localRank == 0) { size_t initOffset = cfg.data.byteOffset / sizeof(float); int numPass = 0, numFail = 0; for (auto rss : transferResources) { @@ -6070,20 +6204,17 @@ static bool IsConfiguredGid(union ibv_gid const& gid) if (anyLocalDst) { if (transferOk) ++numPass; else ++numFail; } } System::Get().Log(" Summary: %d PASS %d FAIL\n", numPass, numFail); - System::Get().Log("Hit to continue: "); - fflush(stdout); - if (scanf("%*c") != 0) { - System::Get().Log("[ERROR] Unexpected input\n"); - exit(1); - } - System::Get().Log("\n"); fflush(stdout); } System::Get().Barrier(); + } else if (cfg.general.useInteractive) { + if (localRank == 0) + System::Get().Log("Transfers complete. Validation disabled (ALWAYS_VALIDATE < 0)\n"); + System::Get().Barrier(); } // Validate results - if (!cfg.data.alwaysValidate) { + if (cfg.data.alwaysValidate == 0) { ERR_APPEND(ValidateAllTransfers(cfg, transfers, transferResources, dstReference, outputBuffer), errResults); }