Skip to content

GitLab

Commit ddb6508f authored by Gilbert Lee's avatar Gilbert Lee
Browse files

TransferBench v1.02

parent 07ac2dce
Hide whitespace changes
Inline Side-by-side
Showing with 398 additions and 390 deletions
CHANGELOG.md
View file @ ddb6508f
# Changelog for TransferBench # Changelog for TransferBench
## v1.02
### Added
- Setting NUM_ITERATIONS to negative number indicates to run for -NUM_ITERATIONS seconds per Test
### Changed
- Copies are now refered to as Transfers instead of Links
- Re-ordering how env vars are displayed (alphabetically now)
### Removed
- Combined timing is now always on for kernel-based GPU copies. COMBINED_TIMING env var has been removed
- Use single sync is no longer supported to facility variable iterations. USE_SINGLE_SYNC env var has been removed
## v1.01 ## v1.01
### Added ### Added
- Adding USE_SINGLE_STREAM feature - Adding USE_SINGLE_STREAM feature
......
EnvVars.hpp
View file @ ddb6508f
...@@ -25,32 +25,32 @@ THE SOFTWARE. ...@@ -25,32 +25,32 @@ THE SOFTWARE.
#include <algorithm> #include <algorithm>
#define TB_VERSION "1.02"
// This class manages environment variable that affect TransferBench // This class manages environment variable that affect TransferBench
class EnvVars class EnvVars
{ {
public: public:
// Default configuration values // Default configuration values
int const DEFAULT_NUM_WARMUPS = 3; int const DEFAULT_NUM_WARMUPS = 3;
int const DEFAULT_NUM_ITERATIONS = 10; int const DEFAULT_NUM_ITERATIONS = 10;
int const DEFAULT_SAMPLING_FACTOR = 1; int const DEFAULT_SAMPLING_FACTOR = 1;
int const DEFAULT_NUM_CPU_PER_LINK = 4; int const DEFAULT_NUM_CPU_PER_TRANSFER = 4;
// Environment variables // Environment variables
int useHipCall; // Use hipMemcpy/hipMemset instead of custom shader kernels int blockBytes; // Each CU, except the last, gets a multiple of this many bytes to copy
int useMemset; // Perform a memset instead of a copy (ignores source memory) int byteOffset; // Byte-offset for memory allocations
int useSingleSync; // Perform synchronization only once after all iterations instead of per iteration int numCpuPerTransfer; // Number of CPU child threads to use per CPU Transfer
int useInteractive; // Pause for user-input before starting transfer loop int numIterations; // Number of timed iterations to perform. If negative, run for -numIterations seconds instead
int combineTiming; // Combines the timing with kernel launch int numWarmups; // Number of un-timed warmup iterations to perform
int outputToCsv; // Output in CSV format int outputToCsv; // Output in CSV format
int byteOffset; // Byte-offset for memory allocations int samplingFactor; // Affects how many different values of N are generated (when N set to 0)
int numWarmups; // Number of un-timed warmup iterations to perform int sharedMemBytes; // Amount of shared memory to use per threadblock
int numIterations; // Number of timed iterations to perform int useHipCall; // Use hipMemcpy/hipMemset instead of custom shader kernels
int samplingFactor; // Affects how many different values of N are generated (when N set to 0) int useInteractive; // Pause for user-input before starting transfer loop
int numCpuPerLink; // Number of CPU child threads to use per CPU link int useMemset; // Perform a memset instead of a copy (ignores source memory)
int sharedMemBytes; // Amount of shared memory to use per threadblock int usePcieIndexing; // Base GPU indexing on PCIe address instead of HIP device
int blockBytes; // Each CU, except the last, gets a multiple of this many bytes to copy int useSingleStream; // Use a single stream per device instead of per Tink. Can not be used with USE_HIP_CALL
int usePcieIndexing; // Base GPU indexing on PCIe address instead of HIP device
int useSingleStream; // Use a single stream per device instead of per Link. Can not be used with USE_HIP_CALL
std::vector<float> fillPattern; // Pattern of floats used to fill source data std::vector<float> fillPattern; // Pattern of floats used to fill source data
...@@ -61,21 +61,19 @@ public: ...@@ -61,21 +61,19 @@ public:
hipDeviceGetAttribute(&maxSharedMemBytes, hipDeviceGetAttribute(&maxSharedMemBytes,
hipDeviceAttributeMaxSharedMemoryPerMultiprocessor, 0); hipDeviceAttributeMaxSharedMemoryPerMultiprocessor, 0);
useHipCall = GetEnvVar("USE_HIP_CALL" , 0); blockBytes = GetEnvVar("BLOCK_BYTES" , 256);
useMemset = GetEnvVar("USE_MEMSET" , 0); byteOffset = GetEnvVar("BYTE_OFFSET" , 0);
useSingleSync = GetEnvVar("USE_SINGLE_SYNC" , 1); numCpuPerTransfer = GetEnvVar("NUM_CPU_PER_TRANSFER", DEFAULT_NUM_CPU_PER_TRANSFER);
useInteractive = GetEnvVar("USE_INTERACTIVE" , 0); numIterations = GetEnvVar("NUM_ITERATIONS" , DEFAULT_NUM_ITERATIONS);
combineTiming = GetEnvVar("COMBINE_TIMING" , 0); numWarmups = GetEnvVar("NUM_WARMUPS" , DEFAULT_NUM_WARMUPS);
outputToCsv = GetEnvVar("OUTPUT_TO_CSV" , 0); outputToCsv = GetEnvVar("OUTPUT_TO_CSV" , 0);
byteOffset = GetEnvVar("BYTE_OFFSET" , 0); samplingFactor = GetEnvVar("SAMPLING_FACTOR" , DEFAULT_SAMPLING_FACTOR);
numWarmups = GetEnvVar("NUM_WARMUPS" , DEFAULT_NUM_WARMUPS); sharedMemBytes = GetEnvVar("SHARED_MEM_BYTES" , maxSharedMemBytes / 2 + 1);
numIterations = GetEnvVar("NUM_ITERATIONS" , DEFAULT_NUM_ITERATIONS); useHipCall = GetEnvVar("USE_HIP_CALL" , 0);
samplingFactor = GetEnvVar("SAMPLING_FACTOR" , DEFAULT_SAMPLING_FACTOR); useInteractive = GetEnvVar("USE_INTERACTIVE" , 0);
numCpuPerLink = GetEnvVar("NUM_CPU_PER_LINK" , DEFAULT_NUM_CPU_PER_LINK); useMemset = GetEnvVar("USE_MEMSET" , 0);
sharedMemBytes = GetEnvVar("SHARED_MEM_BYTES" , maxSharedMemBytes / 2 + 1); usePcieIndexing = GetEnvVar("USE_PCIE_INDEX" , 0);
blockBytes = GetEnvVar("BLOCK_BYTES" , 256); useSingleStream = GetEnvVar("USE_SINGLE_STREAM" , 0);
usePcieIndexing = GetEnvVar("USE_PCIE_INDEX" , 0);
useSingleStream = GetEnvVar("USE_SINGLE_STREAM", 0);
// Check for fill pattern // Check for fill pattern
char* pattern = getenv("FILL_PATTERN"); char* pattern = getenv("FILL_PATTERN");
...@@ -146,19 +144,14 @@ public: ...@@ -146,19 +144,14 @@ public:
printf("[ERROR] NUM_WARMUPS must be set to a non-negative number\n"); printf("[ERROR] NUM_WARMUPS must be set to a non-negative number\n");
exit(1); exit(1);
} }
if (numIterations <= 0)
{
printf("[ERROR] NUM_ITERATIONS must be set to a positive number\n");
exit(1);
}
if (samplingFactor < 1) if (samplingFactor < 1)
{ {
printf("[ERROR] SAMPLING_FACTOR must be greater or equal to 1\n"); printf("[ERROR] SAMPLING_FACTOR must be greater or equal to 1\n");
exit(1); exit(1);
} }
if (numCpuPerLink < 1) if (numCpuPerTransfer < 1)
{ {
printf("[ERROR] NUM_CPU_PER_LINK must be greater or equal to 1\n"); printf("[ERROR] NUM_CPU_PER_TRANSFER must be greater or equal to 1\n");
exit(1); exit(1);
} }
if (sharedMemBytes < 0 || sharedMemBytes > maxSharedMemBytes) if (sharedMemBytes < 0 || sharedMemBytes > maxSharedMemBytes)
...@@ -183,22 +176,20 @@ public: ...@@ -183,22 +176,20 @@ public:
{ {
printf("Environment variables:\n"); printf("Environment variables:\n");
printf("======================\n"); printf("======================\n");
printf(" USE_HIP_CALL - Use hipMemcpy/hipMemset instead of custom shader kernels for GPU-executed copies\n"); printf(" BLOCK_BYTES=B - Each CU (except the last) receives a multiple of BLOCK_BYTES to copy\n");
printf(" USE_MEMSET - Perform a memset instead of a copy (ignores source memory)\n");
printf(" USE_SINGLE_SYNC - Perform synchronization only once after all iterations instead of per iteration\n");
printf(" USE_INTERACTIVE - Pause for user-input before starting transfer loop\n");
printf(" COMBINE_TIMING - Combines timing with launch (potentially lower timing overhead)\n");
printf(" OUTPUT_TO_CSV - Outputs to CSV format if set\n");
printf(" BYTE_OFFSET - Initial byte-offset for memory allocations. Must be multiple of 4. Defaults to 0\n"); printf(" BYTE_OFFSET - Initial byte-offset for memory allocations. Must be multiple of 4. Defaults to 0\n");
printf(" NUM_WARMUPS=W - Perform W untimed warmup iteration(s) per test\n"); printf(" FILL_PATTERN=STR - Fill input buffer with pattern specified in hex digits (0-9,a-f,A-F). Must be even number of digits, (byte-level big-endian)\n");
printf(" NUM_CPU_PER_TRANSFER=C - Use C threads per Transfer for CPU-executed copies\n");
printf(" NUM_ITERATIONS=I - Perform I timed iteration(s) per test\n"); printf(" NUM_ITERATIONS=I - Perform I timed iteration(s) per test\n");
printf(" NUM_WARMUPS=W - Perform W untimed warmup iteration(s) per test\n");
printf(" OUTPUT_TO_CSV - Outputs to CSV format if set\n");
printf(" SAMPLING_FACTOR=F - Add F samples (when possible) between powers of 2 when auto-generating data sizes\n"); printf(" SAMPLING_FACTOR=F - Add F samples (when possible) between powers of 2 when auto-generating data sizes\n");
printf(" NUM_CPU_PER_LINK=C - Use C threads per Link for CPU-executed copies\n");
printf(" FILL_PATTERN=STR - Fill input buffer with pattern specified in hex digits (0-9,a-f,A-F). Must be even number of digits, (byte-level big-endian)\n");
printf(" SHARED_MEM_BYTES=X - Use X shared mem bytes per threadblock, potentially to avoid multiple threadblocks per CU\n"); printf(" SHARED_MEM_BYTES=X - Use X shared mem bytes per threadblock, potentially to avoid multiple threadblocks per CU\n");
printf(" BLOCK_BYTES=B - Each CU (except the last) receives a multiple of BLOCK_BYTES to copy\n"); printf(" USE_HIP_CALL - Use hipMemcpy/hipMemset instead of custom shader kernels for GPU-executed copies\n");
printf(" USE_INTERACTIVE - Pause for user-input before starting transfer loop\n");
printf(" USE_MEMSET - Perform a memset instead of a copy (ignores source memory)\n");
printf(" USE_PCIE_INDEX - Index GPUs by PCIe address-ordering instead of HIP-provided indexing\n"); printf(" USE_PCIE_INDEX - Index GPUs by PCIe address-ordering instead of HIP-provided indexing\n");
printf(" USE_SINGLE_STREAM - Use single stream per device instead of per link. Cannot be used with USE_HIP_CALL\n"); printf(" USE_SINGLE_STREAM - Use single stream per device instead of per Transfer. Cannot be used with USE_HIP_CALL\n");
} }
// Display env var settings // Display env var settings
...@@ -206,45 +197,41 @@ public: ...@@ -206,45 +197,41 @@ public:
{ {
if (!outputToCsv) if (!outputToCsv)
{ {
printf("Run configuration\n"); printf("Run configuration (TransferBench v%s)\n", TB_VERSION);
printf("=====================================================\n"); printf("=====================================================\n");
printf("%-20s = %12d : Each CU gets a multiple of %d bytes to copy\n", "BLOCK_BYTES", blockBytes, blockBytes);
printf("%-20s = %12d : Using byte offset of %d\n", "BYTE_OFFSET", byteOffset, byteOffset);
printf("%-20s = %12s : ", "FILL_PATTERN", getenv("FILL_PATTERN") ? "(specified)" : "(unset)");
if (fillPattern.size())
printf("Pattern: %s", getenv("FILL_PATTERN"));
else
printf("Pseudo-random: (Element i = i modulo 383 + 31)");
printf("\n");
printf("%-20s = %12d : Using %d CPU thread(s) per CPU-based-copy Transfer\n", "NUM_CPU_PER_TRANSFER", numCpuPerTransfer, numCpuPerTransfer);
printf("%-20s = %12d : Running %d %s per topology\n", "NUM_ITERATIONS", numIterations,
numIterations > 0 ? numIterations : -numIterations,
numIterations > 0 ? "timed iteration(s)" : "second(s)");
printf("%-20s = %12d : Running %d warmup iteration(s) per topology\n", "NUM_WARMUPS", numWarmups, numWarmups);
printf("%-20s = %12d : Output to %s\n", "OUTPUT_TO_CSV", outputToCsv,
outputToCsv ? "CSV" : "console");
printf("%-20s = %12s : Using %d shared mem per threadblock\n", "SHARED_MEM_BYTES",
getenv("SHARED_MEM_BYTES") ? "(specified)" : "(unset)", sharedMemBytes);
printf("%-20s = %12d : Using %s for GPU-executed copies\n", "USE_HIP_CALL", useHipCall, printf("%-20s = %12d : Using %s for GPU-executed copies\n", "USE_HIP_CALL", useHipCall,
useHipCall ? "HIP functions" : "custom kernels"); useHipCall ? "HIP functions" : "custom kernels");
printf("%-20s = %12d : Performing %s\n", "USE_MEMSET", useMemset,
useMemset ? "memset" : "memcopy");
if (useHipCall && !useMemset) if (useHipCall && !useMemset)
{ {
char* env = getenv("HSA_ENABLE_SDMA"); char* env = getenv("HSA_ENABLE_SDMA");
printf("%-20s = %12s : %s\n", "HSA_ENABLE_SDMA", env, printf("%-20s = %12s : %s\n", "HSA_ENABLE_SDMA", env,
(env && !strcmp(env, "0")) ? "Using blit kernels for hipMemcpy" : "Using DMA copy engines"); (env && !strcmp(env, "0")) ? "Using blit kernels for hipMemcpy" : "Using DMA copy engines");
} }
printf("%-20s = %12d : %s\n", "USE_SINGLE_SYNC", useSingleSync,
useSingleSync ? "Synchronizing only once, after all iterations" : "Synchronizing per iteration");
printf("%-20s = %12d : Running in %s mode\n", "USE_INTERACTIVE", useInteractive, printf("%-20s = %12d : Running in %s mode\n", "USE_INTERACTIVE", useInteractive,
useInteractive ? "interactive" : "non-interactive"); useInteractive ? "interactive" : "non-interactive");
printf("%-20s = %12d : %s\n", "COMBINE_TIMING", combineTiming, printf("%-20s = %12d : Performing %s\n", "USE_MEMSET", useMemset,
combineTiming ? "Using combined timing+launch" : "Using separate timing / launch"); useMemset ? "memset" : "memcopy");
printf("%-20s = %12d : Output to %s\n", "OUTPUT_TO_CSV", outputToCsv, printf("%-20s = %12d : Using %s-based GPU indexing\n", "USE_PCIE_INDEX",
outputToCsv ? "CSV" : "console"); usePcieIndexing, (usePcieIndexing ? "PCIe" : "HIP"));
printf("%-20s = %12d : Using byte offset of %d\n", "BYTE_OFFSET", byteOffset, byteOffset); printf("%-20s = %12d : Using single stream per %s\n", "USE_SINGLE_STREAM",
printf("%-20s = %12d : Running %d warmup iteration(s) per topology\n", "NUM_WARMUPS", numWarmups, numWarmups); useSingleStream, (useSingleStream ? "device" : "Transfer"));
printf("%-20s = %12d : Running %d timed iteration(s) per topology\n", "NUM_ITERATIONS", numIterations, numIterations);
printf("%-20s = %12d : Using %d CPU thread(s) per CPU-based-copy Link\n", "NUM_CPU_PER_LINK", numCpuPerLink, numCpuPerLink);
printf("%-20s = %12s : ", "FILL_PATTERN", getenv("FILL_PATTERN") ? "(specified)" : "(unset)");
if (fillPattern.size())
{
printf("Pattern: %s", getenv("FILL_PATTERN"));
}
else
{
printf("Pseudo-random: (Element i = i modulo 383 + 31)");
}
printf("\n");
printf("%-20s = %12s : Using %d shared mem per threadblock\n", "SHARED_MEM_BYTES",
getenv("SHARED_MEM_BYTES") ? "(specified)" : "(unset)", sharedMemBytes);
printf("%-20s = %12d : Each CU gets a multiple of %d bytes to copy\n", "BLOCK_BYTES", blockBytes, blockBytes);
printf("%-20s = %12d : Using %s-based GPU indexing\n", "USE_PCIE_INDEX", usePcieIndexing, (usePcieIndexing ? "PCIe" : "HIP"));
printf("%-20s = %12d : Using single stream per %s\n", "USE_SINGLE_STREAM", useSingleStream, (useSingleStream ? "device" : "Link"));
printf("\n"); printf("\n");
} }
}; };
......
LICENSE.md 0 → 100644
View file @ ddb6508f
Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
TransferBench.cpp
View file @ ddb6508f
...@@ -45,25 +45,25 @@ int main(int argc, char **argv) ...@@ -45,25 +45,25 @@ int main(int argc, char **argv)
// Collect environment variables / display current run configuration // Collect environment variables / display current run configuration
EnvVars ev; EnvVars ev;
// Determine number of bytes to run per Link // Determine number of bytes to run per Transfer
// If a non-zero number of bytes is specified, use it // If a non-zero number of bytes is specified, use it
// Otherwise generate array of bytes values to execute over // Otherwise generate array of bytes values to execute over
std::vector<size_t> valuesOfN; std::vector<size_t> valuesOfN;
size_t numBytesPerLink = argc > 2 ? atoll(argv[2]) : DEFAULT_BYTES_PER_LINK; size_t numBytesPerTransfer = argc > 2 ? atoll(argv[2]) : DEFAULT_BYTES_PER_TRANSFER;
if (argc > 2) if (argc > 2)
{ {
// Adjust bytes if unit specified // Adjust bytes if unit specified
char units = argv[2][strlen(argv[2])-1]; char units = argv[2][strlen(argv[2])-1];
switch (units) switch (units)
{ {
case 'K': case 'k': numBytesPerLink *= 1024; break; case 'K': case 'k': numBytesPerTransfer *= 1024; break;
case 'M': case 'm': numBytesPerLink *= 1024*1024; break; case 'M': case 'm': numBytesPerTransfer *= 1024*1024; break;
case 'G': case 'g': numBytesPerLink *= 1024*1024*1024; break; case 'G': case 'g': numBytesPerTransfer *= 1024*1024*1024; break;
} }
} }
PopulateTestSizes(numBytesPerLink, ev.samplingFactor, valuesOfN); PopulateTestSizes(numBytesPerTransfer, ev.samplingFactor, valuesOfN);
// Find the largest N to be used - memory will only be allocated once per link config // Find the largest N to be used - memory will only be allocated once per set of simulatenous Transfers
size_t maxN = valuesOfN[0]; size_t maxN = valuesOfN[0];
for (auto N : valuesOfN) for (auto N : valuesOfN)
maxN = std::max(maxN, N); maxN = std::max(maxN, N);
...@@ -84,15 +84,15 @@ int main(int argc, char **argv) ...@@ -84,15 +84,15 @@ int main(int argc, char **argv)
int skipCpu = (!strcmp(argv[1], "g2g" ) || !strcmp(argv[1], "g2g_rr") ? 1 : 0); int skipCpu = (!strcmp(argv[1], "g2g" ) || !strcmp(argv[1], "g2g_rr") ? 1 : 0);
// Execute peer to peer benchmark mode // Execute peer to peer benchmark mode
RunPeerToPeerBenchmarks(ev, numBytesPerLink / sizeof(float), numBlocksToUse, readMode, skipCpu); RunPeerToPeerBenchmarks(ev, numBytesPerTransfer / sizeof(float), numBlocksToUse, readMode, skipCpu);
exit(0); exit(0);
} }
// Check that Link configuration file can be opened // Check that Transfer configuration file can be opened
FILE* fp = fopen(argv[1], "r"); FILE* fp = fopen(argv[1], "r");
if (!fp) if (!fp)
{ {
printf("[ERROR] Unable to open link configuration file: [%s]\n", argv[1]); printf("[ERROR] Unable to open transfer configuration file: [%s]\n", argv[1]);
exit(1); exit(1);
} }
...@@ -112,17 +112,17 @@ int main(int argc, char **argv) ...@@ -112,17 +112,17 @@ int main(int argc, char **argv)
HIP_CALL(hipGetDeviceCount(&numGpuDevices)); HIP_CALL(hipGetDeviceCount(&numGpuDevices));
int const numCpuDevices = numa_num_configured_nodes(); int const numCpuDevices = numa_num_configured_nodes();
// Track unique pair of links that get used // Track unique pair of transfers that get used
std::set<std::pair<int, int>> peerAccessTracker; std::set<std::pair<int, int>> peerAccessTracker;
// Print CSV header // Print CSV header
if (ev.outputToCsv) if (ev.outputToCsv)
{ {
printf("Test,NumBytes,SrcMem,Executor,DstMem,CUs,BW(GB/s),Time(ms)," printf("Test,NumBytes,SrcMem,Executor,DstMem,CUs,BW(GB/s),Time(ms),"
"LinkDesc,SrcAddr,DstAddr,ByteOffset,numWarmups,numIters\n"); "TransferDesc,SrcAddr,DstAddr,ByteOffset,numWarmups,numIters\n");
} }
// Loop over each line in the Link configuration file // Loop over each line in the Transfer configuration file
int testNum = 0; int testNum = 0;
char line[2048]; char line[2048];
while(fgets(line, 2048, fp)) while(fgets(line, 2048, fp))
...@@ -130,33 +130,33 @@ int main(int argc, char **argv) ...@@ -130,33 +130,33 @@ int main(int argc, char **argv)
// Check if line is a comment to be echoed to output (starts with ##) // Check if line is a comment to be echoed to output (starts with ##)
if (!ev.outputToCsv && line[0] == '#' && line[1] == '#') printf("%s", line); if (!ev.outputToCsv && line[0] == '#' && line[1] == '#') printf("%s", line);
// Parse links from configuration file // Parse transfers from configuration file
LinkMap linkMap; TransferMap transferMap;
ParseLinks(line, numCpuDevices, numGpuDevices, linkMap); ParseTransfers(line, numCpuDevices, numGpuDevices, transferMap);
if (linkMap.size() == 0) continue; if (transferMap.size() == 0) continue;
testNum++; testNum++;
// Prepare (maximum) memory for each link // Prepare (maximum) memory for each transfer
std::vector<Link*> linkList; std::vector<Transfer*> transferList;
for (auto& exeInfoPair : linkMap) for (auto& exeInfoPair : transferMap)
{ {
ExecutorInfo& exeInfo = exeInfoPair.second; ExecutorInfo& exeInfo = exeInfoPair.second;
exeInfo.totalTime = 0.0; exeInfo.totalTime = 0.0;
exeInfo.totalBlocks = 0; exeInfo.totalBlocks = 0;
for (Link& link : exeInfo.links) for (Transfer& transfer : exeInfo.transfers)
{ {
// Get some aliases to link variables // Get some aliases to transfer variables
MemType const& exeMemType = link.exeMemType; MemType const& exeMemType = transfer.exeMemType;
MemType const& srcMemType = link.srcMemType; MemType const& srcMemType = transfer.srcMemType;
MemType const& dstMemType = link.dstMemType; MemType const& dstMemType = transfer.dstMemType;
int const& blocksToUse = link.numBlocksToUse; int const& blocksToUse = transfer.numBlocksToUse;
// Get potentially remapped device indices // Get potentially remapped device indices
int const srcIndex = RemappedIndex(link.srcIndex, srcMemType); int const srcIndex = RemappedIndex(transfer.srcIndex, srcMemType);
int const exeIndex = RemappedIndex(link.exeIndex, exeMemType); int const exeIndex = RemappedIndex(transfer.exeIndex, exeMemType);
int const dstIndex = RemappedIndex(link.dstIndex, dstMemType); int const dstIndex = RemappedIndex(transfer.dstIndex, dstMemType);
// Enable peer-to-peer access if necessary (can only be called once per unique pair) // Enable peer-to-peer access if necessary (can only be called once per unique pair)
if (exeMemType == MEM_GPU) if (exeMemType == MEM_GPU)
...@@ -185,11 +185,11 @@ int main(int argc, char **argv) ...@@ -185,11 +185,11 @@ int main(int argc, char **argv)
} }
// Allocate (maximum) source / destination memory based on type / device index // Allocate (maximum) source / destination memory based on type / device index
AllocateMemory(srcMemType, srcIndex, maxN * sizeof(float) + ev.byteOffset, (void**)&link.srcMem); AllocateMemory(srcMemType, srcIndex, maxN * sizeof(float) + ev.byteOffset, (void**)&transfer.srcMem);
AllocateMemory(dstMemType, dstIndex, maxN * sizeof(float) + ev.byteOffset, (void**)&link.dstMem); AllocateMemory(dstMemType, dstIndex, maxN * sizeof(float) + ev.byteOffset, (void**)&transfer.dstMem);
link.blockParam.resize(exeMemType == MEM_CPU ? ev.numCpuPerLink : blocksToUse); transfer.blockParam.resize(exeMemType == MEM_CPU ? ev.numCpuPerTransfer : blocksToUse);
exeInfo.totalBlocks += link.blockParam.size(); exeInfo.totalBlocks += transfer.blockParam.size();
linkList.push_back(&link); transferList.push_back(&transfer);
} }
// Prepare GPU resources for GPU executors // Prepare GPU resources for GPU executors
...@@ -200,11 +200,11 @@ int main(int argc, char **argv) ...@@ -200,11 +200,11 @@ int main(int argc, char **argv)
AllocateMemory(exeMemType, exeIndex, exeInfo.totalBlocks * sizeof(BlockParam), AllocateMemory(exeMemType, exeIndex, exeInfo.totalBlocks * sizeof(BlockParam),
(void**)&exeInfo.blockParamGpu); (void**)&exeInfo.blockParamGpu);
int const numLinksToRun = ev.useSingleStream ? 1 : exeInfo.links.size(); int const numTransfersToRun = ev.useSingleStream ? 1 : exeInfo.transfers.size();
exeInfo.streams.resize(numLinksToRun); exeInfo.streams.resize(numTransfersToRun);
exeInfo.startEvents.resize(numLinksToRun); exeInfo.startEvents.resize(numTransfersToRun);
exeInfo.stopEvents.resize(numLinksToRun); exeInfo.stopEvents.resize(numTransfersToRun);
for (int i = 0; i < numLinksToRun; ++i) for (int i = 0; i < numTransfersToRun; ++i)
{ {
HIP_CALL(hipSetDevice(exeIndex)); HIP_CALL(hipSetDevice(exeIndex));
HIP_CALL(hipStreamCreate(&exeInfo.streams[i])); HIP_CALL(hipStreamCreate(&exeInfo.streams[i]));
...@@ -212,48 +212,52 @@ int main(int argc, char **argv) ...@@ -212,48 +212,52 @@ int main(int argc, char **argv)
HIP_CALL(hipEventCreate(&exeInfo.stopEvents[i])); HIP_CALL(hipEventCreate(&exeInfo.stopEvents[i]));
} }
int linkOffset = 0; int transferOffset = 0;
for (int i = 0; i < exeInfo.links.size(); i++) for (int i = 0; i < exeInfo.transfers.size(); i++)
{ {
exeInfo.links[i].blockParamGpuPtr = exeInfo.blockParamGpu + linkOffset; exeInfo.transfers[i].blockParamGpuPtr = exeInfo.blockParamGpu + transferOffset;
linkOffset += exeInfo.links[i].blockParam.size(); transferOffset += exeInfo.transfers[i].blockParam.size();
} }
} }
} }
// Loop over all the different number of bytes to use per Link // Loop over all the different number of bytes to use per Transfer
for (auto N : valuesOfN) for (auto N : valuesOfN)
{ {
if (!ev.outputToCsv) printf("Test %d: [%lu bytes]\n", testNum, N * sizeof(float)); if (!ev.outputToCsv) printf("Test %d: [%lu bytes]\n", testNum, N * sizeof(float));
// Prepare input memory and block parameters for current N // Prepare input memory and block parameters for current N
for (auto& exeInfoPair : linkMap) for (auto& exeInfoPair : transferMap)
{ {
ExecutorInfo& exeInfo = exeInfoPair.second; ExecutorInfo& exeInfo = exeInfoPair.second;
int linkOffset = 0; int transferOffset = 0;
for (int i = 0; i < exeInfo.links.size(); ++i) for (int i = 0; i < exeInfo.transfers.size(); ++i)
{ {
Link& link = exeInfo.links[i]; Transfer& transfer = exeInfo.transfers[i];
link.PrepareBlockParams(ev, N); transfer.PrepareBlockParams(ev, N);
// Copy block parameters to GPU for GPU executors // Copy block parameters to GPU for GPU executors
if (link.exeMemType == MEM_GPU) if (transfer.exeMemType == MEM_GPU)
{ {
HIP_CALL(hipMemcpy(&exeInfo.blockParamGpu[linkOffset], HIP_CALL(hipMemcpy(&exeInfo.blockParamGpu[transferOffset],
link.blockParam.data(), transfer.blockParam.data(),
link.blockParam.size() * sizeof(BlockParam), transfer.blockParam.size() * sizeof(BlockParam),
hipMemcpyHostToDevice)); hipMemcpyHostToDevice));
linkOffset += link.blockParam.size(); transferOffset += transfer.blockParam.size();
} }
} }
} }
// Launch kernels (warmup iterations are not counted) // Launch kernels (warmup iterations are not counted)
double totalCpuTime = 0; double totalCpuTime = 0;
for (int iteration = -ev.numWarmups; iteration < ev.numIterations; iteration++) size_t numTimedIterations = 0;
for (int iteration = -ev.numWarmups; ; iteration++)
{ {
if (ev.numIterations > 0 && iteration >= ev.numIterations) break;
if (ev.numIterations < 0 && totalCpuTime > -ev.numIterations) break;
// Pause before starting first timed iteration in interactive mode // Pause before starting first timed iteration in interactive mode
if (ev.useInteractive && iteration == 0) if (ev.useInteractive && iteration == 0)
{ {
...@@ -265,18 +269,18 @@ int main(int argc, char **argv) ...@@ -265,18 +269,18 @@ int main(int argc, char **argv)
// Start CPU timing for this iteration // Start CPU timing for this iteration
auto cpuStart = std::chrono::high_resolution_clock::now(); auto cpuStart = std::chrono::high_resolution_clock::now();
// Execute all links in parallel // Execute all Transfers in parallel
for (auto& exeInfoPair : linkMap) for (auto& exeInfoPair : transferMap)
{ {
ExecutorInfo& exeInfo = exeInfoPair.second; ExecutorInfo& exeInfo = exeInfoPair.second;
int const numLinksToRun = ev.useSingleStream ? 1 : exeInfo.links.size(); int const numTransfersToRun = ev.useSingleStream ? 1 : exeInfo.transfers.size();
for (int i = 0; i < numLinksToRun; ++i) for (int i = 0; i < numTransfersToRun; ++i)
threads.push(std::thread(RunLink, std::ref(ev), N, iteration, std::ref(exeInfo), i)); threads.push(std::thread(RunTransfer, std::ref(ev), N, iteration, std::ref(exeInfo), i));
} }
// Wait for all threads to finish // Wait for all threads to finish
int const numLinks = threads.size(); int const numTransfers = threads.size();
for (int i = 0; i < numLinks; i++) for (int i = 0; i < numTransfers; i++)
{ {
threads.top().join(); threads.top().join();
threads.pop(); threads.pop();
...@@ -286,9 +290,11 @@ int main(int argc, char **argv) ...@@ -286,9 +290,11 @@ int main(int argc, char **argv)
auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
double deltaSec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count(); double deltaSec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count();
if (iteration >= 0)
{
if (iteration >= 0) totalCpuTime += deltaSec; ++numTimedIterations;
totalCpuTime += deltaSec;
}
} }
// Pause for interactive mode // Pause for interactive mode
...@@ -299,89 +305,89 @@ int main(int argc, char **argv) ...@@ -299,89 +305,89 @@ int main(int argc, char **argv)
printf("\n"); printf("\n");
} }
// Validate that each link has transferred correctly // Validate that each transfer has transferred correctly
int const numLinks = linkList.size(); int const numTransfers = transferList.size();
for (auto link : linkList) for (auto transfer : transferList)
CheckOrFill(MODE_CHECK, N, ev.useMemset, ev.useHipCall, ev.fillPattern, link->dstMem + initOffset); CheckOrFill(MODE_CHECK, N, ev.useMemset, ev.useHipCall, ev.fillPattern, transfer->dstMem + initOffset);
// Report timings // Report timings
totalCpuTime = totalCpuTime / (1.0 * ev.numIterations) * 1000; totalCpuTime = totalCpuTime / (1.0 * numTimedIterations) * 1000;
double totalBandwidthGbs = (numLinks * N * sizeof(float) / 1.0E6) / totalCpuTime; double totalBandwidthGbs = (numTransfers * N * sizeof(float) / 1.0E6) / totalCpuTime;
double maxGpuTime = 0; double maxGpuTime = 0;
if (ev.useSingleStream) if (ev.useSingleStream)
{ {
for (auto& exeInfoPair : linkMap) for (auto& exeInfoPair : transferMap)
{ {
ExecutorInfo const& exeInfo = exeInfoPair.second; ExecutorInfo const& exeInfo = exeInfoPair.second;
MemType const exeMemType = exeInfoPair.first.first; MemType const exeMemType = exeInfoPair.first.first;
int const exeIndex = exeInfoPair.first.second; int const exeIndex = exeInfoPair.first.second;
double exeDurationMsec = exeInfo.totalTime / (1.0 * ev.numIterations); double exeDurationMsec = exeInfo.totalTime / (1.0 * numTimedIterations);
double exeBandwidthGbs = (exeInfo.links.size() * N * sizeof(float) / 1.0E9) / exeDurationMsec * 1000.0f; double exeBandwidthGbs = (exeInfo.transfers.size() * N * sizeof(float) / 1.0E9) / exeDurationMsec * 1000.0f;
maxGpuTime = std::max(maxGpuTime, exeDurationMsec); maxGpuTime = std::max(maxGpuTime, exeDurationMsec);
if (!ev.outputToCsv) if (!ev.outputToCsv)
{ {
printf(" Executor: %cPU %02d (# Links %02lu)| %9.3f GB/s | %8.3f ms |\n", printf(" Executor: %cPU %02d (# Transfers %02lu)| %9.3f GB/s | %8.3f ms |\n",
MemTypeStr[exeMemType], exeIndex, exeInfo.links.size(), exeBandwidthGbs, exeDurationMsec); MemTypeStr[exeMemType], exeIndex, exeInfo.transfers.size(), exeBandwidthGbs, exeDurationMsec);
for (auto link : exeInfo.links) for (auto transfer : exeInfo.transfers)
{ {
double linkDurationMsec = link.linkTime / (1.0 * ev.numIterations); double transferDurationMsec = transfer.transferTime / (1.0 * numTimedIterations);
double linkBandwidthGbs = (N * sizeof(float) / 1.0E9) / linkDurationMsec * 1000.0f; double transferBandwidthGbs = (N * sizeof(float) / 1.0E9) / transferDurationMsec * 1000.0f;
printf(" Link %02d | %9.3f GB/s | %8.3f ms | %c%02d -> %c%02d:(%02d) -> %c%02d\n", printf(" Transfer %02d | %9.3f GB/s | %8.3f ms | %c%02d -> %c%02d:(%03d) -> %c%02d\n",
link.linkIndex, transfer.transferIndex,
linkBandwidthGbs, transferBandwidthGbs,
linkDurationMsec, transferDurationMsec,
MemTypeStr[link.srcMemType], link.srcIndex, MemTypeStr[transfer.srcMemType], transfer.srcIndex,
MemTypeStr[link.exeMemType], link.exeIndex, MemTypeStr[transfer.exeMemType], transfer.exeIndex,
link.exeMemType == MEM_CPU ? ev.numCpuPerLink : link.numBlocksToUse, transfer.exeMemType == MEM_CPU ? ev.numCpuPerTransfer : transfer.numBlocksToUse,
MemTypeStr[link.dstMemType], link.dstIndex); MemTypeStr[transfer.dstMemType], transfer.dstIndex);
} }
} }
else else
{ {
printf("%d,%lu,ALL,%c%02d,ALL,ALL,%.3f,%.3f,ALL,ALL,ALL,%d,%d,%d\n", printf("%d,%lu,ALL,%c%02d,ALL,ALL,%.3f,%.3f,ALL,ALL,ALL,%d,%d,%lu\n",
testNum, N * sizeof(float), testNum, N * sizeof(float),
MemTypeStr[exeMemType], exeIndex, MemTypeStr[exeMemType], exeIndex,
exeBandwidthGbs, exeDurationMsec, exeBandwidthGbs, exeDurationMsec,
ev.byteOffset, ev.byteOffset,
ev.numWarmups, ev.numIterations); ev.numWarmups, numTimedIterations);
} }
} }
} }
else else
{ {
for (auto link : linkList) for (auto transfer : transferList)
{ {
double linkDurationMsec = link->linkTime / (1.0 * ev.numIterations); double transferDurationMsec = transfer->transferTime / (1.0 * numTimedIterations);
double linkBandwidthGbs = (N * sizeof(float) / 1.0E9) / linkDurationMsec * 1000.0f; double transferBandwidthGbs = (N * sizeof(float) / 1.0E9) / transferDurationMsec * 1000.0f;
maxGpuTime = std::max(maxGpuTime, linkDurationMsec); maxGpuTime = std::max(maxGpuTime, transferDurationMsec);
if (!ev.outputToCsv) if (!ev.outputToCsv)
{ {
printf(" Link %02d: %c%02d -> [%cPU %02d:%02d] -> %c%02d | %9.3f GB/s | %8.3f ms | %-16s\n", printf(" Transfer %02d: %c%02d -> [%cPU %02d:%03d] -> %c%02d | %9.3f GB/s | %8.3f ms | %-16s\n",
link->linkIndex, transfer->transferIndex,
MemTypeStr[link->srcMemType], link->srcIndex, MemTypeStr[transfer->srcMemType], transfer->srcIndex,
MemTypeStr[link->exeMemType], link->exeIndex, MemTypeStr[transfer->exeMemType], transfer->exeIndex,
link->exeMemType == MEM_CPU ? ev.numCpuPerLink : link->numBlocksToUse, transfer->exeMemType == MEM_CPU ? ev.numCpuPerTransfer : transfer->numBlocksToUse,
MemTypeStr[link->dstMemType], link->dstIndex, MemTypeStr[transfer->dstMemType], transfer->dstIndex,
linkBandwidthGbs, linkDurationMsec, transferBandwidthGbs, transferDurationMsec,
GetLinkDesc(*link).c_str()); GetTransferDesc(*transfer).c_str());
} }
else else
{ {
printf("%d,%lu,%c%02d,%c%02d,%c%02d,%d,%.3f,%.3f,%s,%p,%p,%d,%d,%d\n", printf("%d,%lu,%c%02d,%c%02d,%c%02d,%d,%.3f,%.3f,%s,%p,%p,%d,%d,%lu\n",
testNum, N * sizeof(float), testNum, N * sizeof(float),
MemTypeStr[link->srcMemType], link->srcIndex, MemTypeStr[transfer->srcMemType], transfer->srcIndex,
MemTypeStr[link->exeMemType], link->exeIndex, MemTypeStr[transfer->exeMemType], transfer->exeIndex,
MemTypeStr[link->dstMemType], link->dstIndex, MemTypeStr[transfer->dstMemType], transfer->dstIndex,
link->exeMemType == MEM_CPU ? ev.numCpuPerLink : link->numBlocksToUse, transfer->exeMemType == MEM_CPU ? ev.numCpuPerTransfer : transfer->numBlocksToUse,
linkBandwidthGbs, linkDurationMsec, transferBandwidthGbs, transferDurationMsec,
GetLinkDesc(*link).c_str(), GetTransferDesc(*transfer).c_str(),
link->srcMem + initOffset, link->dstMem + initOffset, transfer->srcMem + initOffset, transfer->dstMem + initOffset,
ev.byteOffset, ev.byteOffset,
ev.numWarmups, ev.numIterations); ev.numWarmups, numTimedIterations);
} }
} }
} }
...@@ -389,32 +395,32 @@ int main(int argc, char **argv) ...@@ -389,32 +395,32 @@ int main(int argc, char **argv)
// Display aggregate statistics // Display aggregate statistics
if (!ev.outputToCsv) if (!ev.outputToCsv)
{ {
printf(" Aggregate Bandwidth (CPU timed) | %9.3f GB/s | %8.3f ms | Overhead: %.3f ms\n", totalBandwidthGbs, totalCpuTime, printf(" Aggregate Bandwidth (CPU timed) | %9.3f GB/s | %8.3f ms | Overhead: %.3f ms\n", totalBandwidthGbs, totalCpuTime,
totalCpuTime - maxGpuTime); totalCpuTime - maxGpuTime);
} }
else else
{ {
printf("%d,%lu,ALL,ALL,ALL,ALL,%.3f,%.3f,ALL,ALL,ALL,%d,%d,%d\n", printf("%d,%lu,ALL,ALL,ALL,ALL,%.3f,%.3f,ALL,ALL,ALL,%d,%d,%lu\n",
testNum, N * sizeof(float), totalBandwidthGbs, totalCpuTime, ev.byteOffset, testNum, N * sizeof(float), totalBandwidthGbs, totalCpuTime, ev.byteOffset,
ev.numWarmups, ev.numIterations); ev.numWarmups, numTimedIterations);
} }
} }
// Release GPU memory // Release GPU memory
for (auto exeInfoPair : linkMap) for (auto exeInfoPair : transferMap)
{ {
ExecutorInfo& exeInfo = exeInfoPair.second; ExecutorInfo& exeInfo = exeInfoPair.second;
for (auto& link : exeInfo.links) for (auto& transfer : exeInfo.transfers)
{ {
// Get some aliases to link variables // Get some aliases to Transfer variables
MemType const& exeMemType = link.exeMemType; MemType const& exeMemType = transfer.exeMemType;
MemType const& srcMemType = link.srcMemType; MemType const& srcMemType = transfer.srcMemType;
MemType const& dstMemType = link.dstMemType; MemType const& dstMemType = transfer.dstMemType;
// Allocate (maximum) source / destination memory based on type / device index // Allocate (maximum) source / destination memory based on type / device index
DeallocateMemory(srcMemType, link.srcMem); DeallocateMemory(srcMemType, transfer.srcMem);
DeallocateMemory(dstMemType, link.dstMem); DeallocateMemory(dstMemType, transfer.dstMem);
link.blockParam.clear(); transfer.blockParam.clear();
} }
MemType const exeMemType = exeInfoPair.first.first; MemType const exeMemType = exeInfoPair.first.first;
...@@ -422,8 +428,8 @@ int main(int argc, char **argv) ...@@ -422,8 +428,8 @@ int main(int argc, char **argv)
if (exeMemType == MEM_GPU) if (exeMemType == MEM_GPU)
{ {
DeallocateMemory(exeMemType, exeInfo.blockParamGpu); DeallocateMemory(exeMemType, exeInfo.blockParamGpu);
int const numLinksToRun = ev.useSingleStream ? 1 : exeInfo.links.size(); int const numTransfersToRun = ev.useSingleStream ? 1 : exeInfo.transfers.size();
for (int i = 0; i < numLinksToRun; ++i) for (int i = 0; i < numTransfersToRun; ++i)
{ {
HIP_CALL(hipEventDestroy(exeInfo.startEvents[i])); HIP_CALL(hipEventDestroy(exeInfo.startEvents[i]));
HIP_CALL(hipEventDestroy(exeInfo.stopEvents[i])); HIP_CALL(hipEventDestroy(exeInfo.stopEvents[i]));
...@@ -439,7 +445,7 @@ int main(int argc, char **argv) ...@@ -439,7 +445,7 @@ int main(int argc, char **argv)
void DisplayUsage(char const* cmdName) void DisplayUsage(char const* cmdName)
{ {
printf("TransferBench v. %s\n", TB_VERSION); printf("TransferBench v%s\n", TB_VERSION);
printf("========================================\n"); printf("========================================\n");
if (numa_available() == -1) if (numa_available() == -1)
...@@ -453,16 +459,16 @@ void DisplayUsage(char const* cmdName) ...@@ -453,16 +459,16 @@ void DisplayUsage(char const* cmdName)
printf("Usage: %s config <N>\n", cmdName); printf("Usage: %s config <N>\n", cmdName);
printf(" config: Either:\n"); printf(" config: Either:\n");
printf(" - Filename of configFile containing Links to execute (see example.cfg for format)\n"); printf(" - Filename of configFile containing Transfers to execute (see example.cfg for format)\n");
printf(" - Name of preset benchmark:\n"); printf(" - Name of preset benchmark:\n");
printf(" p2p - All CPU/GPU pairs benchmark\n"); printf(" p2p - All CPU/GPU pairs benchmark\n");
printf(" p2p_rr - All CPU/GPU pairs benchmark with remote reads\n"); printf(" p2p_rr - All CPU/GPU pairs benchmark with remote reads\n");
printf(" g2g - All GPU/GPU pairs benchmark\n"); printf(" g2g - All GPU/GPU pairs benchmark\n");
printf(" g2g_rr - All GPU/GPU pairs benchmark with remote reads\n"); printf(" g2g_rr - All GPU/GPU pairs benchmark with remote reads\n");
printf(" - 3rd optional argument will be used as # of CUs to use (uses all by default)\n"); printf(" - 3rd optional argument will be used as # of CUs to use (uses all by default)\n");
printf(" N : (Optional) Number of bytes to transfer per link.\n"); printf(" N : (Optional) Number of bytes to copy per Transfer.\n");
printf(" If not specified, defaults to %lu bytes. Must be a multiple of 4 bytes\n", printf(" If not specified, defaults to %lu bytes. Must be a multiple of 4 bytes\n",
DEFAULT_BYTES_PER_LINK); DEFAULT_BYTES_PER_TRANSFER);
printf(" If 0 is specified, a range of Ns will be benchmarked\n"); printf(" If 0 is specified, a range of Ns will be benchmarked\n");
printf(" May append a suffix ('K', 'M', 'G') for kilobytes / megabytes / gigabytes\n"); printf(" May append a suffix ('K', 'M', 'G') for kilobytes / megabytes / gigabytes\n");
printf("\n"); printf("\n");
...@@ -574,21 +580,21 @@ void DisplayTopology(bool const outputToCsv) ...@@ -574,21 +580,21 @@ void DisplayTopology(bool const outputToCsv)
} }
} }
void PopulateTestSizes(size_t const numBytesPerLink, void PopulateTestSizes(size_t const numBytesPerTransfer,
int const samplingFactor, int const samplingFactor,
std::vector<size_t>& valuesOfN) std::vector<size_t>& valuesOfN)
{ {
valuesOfN.clear(); valuesOfN.clear();
// If the number of bytes is specified, use it // If the number of bytes is specified, use it
if (numBytesPerLink != 0) if (numBytesPerTransfer != 0)
{ {
if (numBytesPerLink % 4) if (numBytesPerTransfer % 4)
{ {
printf("[ERROR] numBytesPerLink (%lu) must be a multiple of 4\n", numBytesPerLink); printf("[ERROR] numBytesPerTransfer (%lu) must be a multiple of 4\n", numBytesPerTransfer);
exit(1); exit(1);
} }
size_t N = numBytesPerLink / sizeof(float); size_t N = numBytesPerTransfer / sizeof(float);
valuesOfN.push_back(N); valuesOfN.push_back(N);
} }
else else
...@@ -642,24 +648,24 @@ void ParseMemType(std::string const& token, int const numCpus, int const numGpus ...@@ -642,24 +648,24 @@ void ParseMemType(std::string const& token, int const numCpus, int const numGpus
} }
} }
// Helper function to parse a list of link definitions // Helper function to parse a list of Transfer definitions
void ParseLinks(char* line, int numCpus, int numGpus, LinkMap& linkMap) void ParseTransfers(char* line, int numCpus, int numGpus, TransferMap& transferMap)
{ {
// Replace any round brackets or '->' with spaces, // Replace any round brackets or '->' with spaces,
for (int i = 1; line[i]; i++) for (int i = 1; line[i]; i++)
if (line[i] == '(' || line[i] == ')' || line[i] == '-' || line[i] == '>' ) line[i] = ' '; if (line[i] == '(' || line[i] == ')' || line[i] == '-' || line[i] == '>' ) line[i] = ' ';
linkMap.clear(); transferMap.clear();
int numLinks = 0; int numTransfers = 0;
std::istringstream iss(line); std::istringstream iss(line);
iss >> numLinks; iss >> numTransfers;
if (iss.fail()) return; if (iss.fail()) return;
std::string exeMem; std::string exeMem;
std::string srcMem; std::string srcMem;
std::string dstMem; std::string dstMem;
if (numLinks > 0) if (numTransfers > 0)
{ {
// Method 1: Take in triples (srcMem, exeMem, dstMem) // Method 1: Take in triples (srcMem, exeMem, dstMem)
int numBlocksToUse; int numBlocksToUse;
...@@ -669,64 +675,64 @@ void ParseLinks(char* line, int numCpus, int numGpus, LinkMap& linkMap) ...@@ -669,64 +675,64 @@ void ParseLinks(char* line, int numCpus, int numGpus, LinkMap& linkMap)
printf("Parsing error: Number of blocks to use (%d) must be greater than 0\n", numBlocksToUse); printf("Parsing error: Number of blocks to use (%d) must be greater than 0\n", numBlocksToUse);
exit(1); exit(1);
} }
for (int i = 0; i < numLinks; i++) for (int i = 0; i < numTransfers; i++)
{ {
Link link; Transfer transfer;
link.linkIndex = i; transfer.transferIndex = i;
iss >> srcMem >> exeMem >> dstMem; iss >> srcMem >> exeMem >> dstMem;
if (iss.fail()) if (iss.fail())
{ {
printf("Parsing error: Unable to read valid Link triplet (possibly missing a SRC or EXE or DST)\n"); printf("Parsing error: Unable to read valid Transfer triplet (possibly missing a SRC or EXE or DST)\n");
exit(1); exit(1);
} }
ParseMemType(srcMem, numCpus, numGpus, &link.srcMemType, &link.srcIndex); ParseMemType(srcMem, numCpus, numGpus, &transfer.srcMemType, &transfer.srcIndex);
ParseMemType(exeMem, numCpus, numGpus, &link.exeMemType, &link.exeIndex); ParseMemType(exeMem, numCpus, numGpus, &transfer.exeMemType, &transfer.exeIndex);
ParseMemType(dstMem, numCpus, numGpus, &link.dstMemType, &link.dstIndex); ParseMemType(dstMem, numCpus, numGpus, &transfer.dstMemType, &transfer.dstIndex);
link.numBlocksToUse = numBlocksToUse; transfer.numBlocksToUse = numBlocksToUse;
// Ensure executor is either CPU or GPU // Ensure executor is either CPU or GPU
if (link.exeMemType != MEM_CPU && link.exeMemType != MEM_GPU) if (transfer.exeMemType != MEM_CPU && transfer.exeMemType != MEM_GPU)
{ {
printf("[ERROR] Executor must either be CPU ('C') or GPU ('G'), (from (%s->%s->%s %d))\n", printf("[ERROR] Executor must either be CPU ('C') or GPU ('G'), (from (%s->%s->%s %d))\n",
srcMem.c_str(), exeMem.c_str(), dstMem.c_str(), link.numBlocksToUse); srcMem.c_str(), exeMem.c_str(), dstMem.c_str(), transfer.numBlocksToUse);
exit(1); exit(1);
} }
Executor executor(link.exeMemType, link.exeIndex); Executor executor(transfer.exeMemType, transfer.exeIndex);
ExecutorInfo& executorInfo = linkMap[executor]; ExecutorInfo& executorInfo = transferMap[executor];
executorInfo.totalBlocks += link.numBlocksToUse; executorInfo.totalBlocks += transfer.numBlocksToUse;
executorInfo.links.push_back(link); executorInfo.transfers.push_back(transfer);
} }
} }
else else
{ {
// Method 2: Read in quads (srcMem, exeMem, dstMem, Read common # blocks to use, then read (src, dst) doubles // Method 2: Read in quads (srcMem, exeMem, dstMem, Read common # blocks to use, then read (src, dst) doubles
numLinks *= -1; numTransfers *= -1;
for (int i = 0; i < numLinks; i++) for (int i = 0; i < numTransfers; i++)
{ {
Link link; Transfer transfer;
link.linkIndex = i; transfer.transferIndex = i;
iss >> srcMem >> exeMem >> dstMem >> link.numBlocksToUse; iss >> srcMem >> exeMem >> dstMem >> transfer.numBlocksToUse;
if (iss.fail()) if (iss.fail())
{ {
printf("Parsing error: Unable to read valid Link quadruple (possibly missing a SRC or EXE or DST or #CU)\n"); printf("Parsing error: Unable to read valid Transfer quadruple (possibly missing a SRC or EXE or DST or #CU)\n");
exit(1); exit(1);
} }
ParseMemType(srcMem, numCpus, numGpus, &link.srcMemType, &link.srcIndex); ParseMemType(srcMem, numCpus, numGpus, &transfer.srcMemType, &transfer.srcIndex);
ParseMemType(exeMem, numCpus, numGpus, &link.exeMemType, &link.exeIndex); ParseMemType(exeMem, numCpus, numGpus, &transfer.exeMemType, &transfer.exeIndex);
ParseMemType(dstMem, numCpus, numGpus, &link.dstMemType, &link.dstIndex); ParseMemType(dstMem, numCpus, numGpus, &transfer.dstMemType, &transfer.dstIndex);
if (link.exeMemType != MEM_CPU && link.exeMemType != MEM_GPU) if (transfer.exeMemType != MEM_CPU && transfer.exeMemType != MEM_GPU)
{ {
printf("[ERROR] Executor must either be CPU ('C') or GPU ('G'), (from (%s->%s->%s %d))\n" printf("[ERROR] Executor must either be CPU ('C') or GPU ('G'), (from (%s->%s->%s %d))\n"
, srcMem.c_str(), exeMem.c_str(), dstMem.c_str(), link.numBlocksToUse); , srcMem.c_str(), exeMem.c_str(), dstMem.c_str(), transfer.numBlocksToUse);
exit(1); exit(1);
} }
Executor executor(link.exeMemType, link.exeIndex); Executor executor(transfer.exeMemType, transfer.exeIndex);
ExecutorInfo& executorInfo = linkMap[executor]; ExecutorInfo& executorInfo = transferMap[executor];
executorInfo.totalBlocks += link.numBlocksToUse; executorInfo.totalBlocks += transfer.numBlocksToUse;
executorInfo.links.push_back(link); executorInfo.transfers.push_back(transfer);
} }
} }
} }
...@@ -772,6 +778,7 @@ void AllocateMemory(MemType memType, int devIndex, size_t numBytes, void** memPt ...@@ -772,6 +778,7 @@ void AllocateMemory(MemType memType, int devIndex, size_t numBytes, void** memPt
} }
// Allocate host-pinned memory (should respect NUMA mem policy) // Allocate host-pinned memory (should respect NUMA mem policy)
if (memType == MEM_CPU_FINE) if (memType == MEM_CPU_FINE)
{ {
HIP_CALL(hipHostMalloc((void **)memPtr, numBytes, hipHostMallocNumaUser)); HIP_CALL(hipHostMalloc((void **)memPtr, numBytes, hipHostMallocNumaUser));
...@@ -969,108 +976,95 @@ error: ...@@ -969,108 +976,95 @@ error:
exit(1); exit(1);
} }
std::string GetLinkDesc(Link const& link) std::string GetTransferDesc(Transfer const& transfer)
{ {
return GetDesc(link.srcMemType, link.srcIndex, link.exeMemType, link.exeIndex) + "-" return GetDesc(transfer.srcMemType, transfer.srcIndex, transfer.exeMemType, transfer.exeIndex) + "-"
+ GetDesc(link.exeMemType, link.exeIndex, link.dstMemType, link.dstIndex); + GetDesc(transfer.exeMemType, transfer.exeIndex, transfer.dstMemType, transfer.dstIndex);
} }
void RunLink(EnvVars const& ev, size_t const N, int const iteration, ExecutorInfo& exeInfo, int const linkIdx) void RunTransfer(EnvVars const& ev, size_t const N, int const iteration, ExecutorInfo& exeInfo, int const transferIdx)
{ {
Link& link = exeInfo.links[linkIdx]; Transfer& transfer = exeInfo.transfers[transferIdx];
// GPU execution agent // GPU execution agent
if (link.exeMemType == MEM_GPU) if (transfer.exeMemType == MEM_GPU)
{ {
// Switch to executing GPU // Switch to executing GPU
int const exeIndex = RemappedIndex(link.exeIndex, MEM_GPU); int const exeIndex = RemappedIndex(transfer.exeIndex, MEM_GPU);
HIP_CALL(hipSetDevice(exeIndex)); HIP_CALL(hipSetDevice(exeIndex));
hipStream_t& stream = exeInfo.streams[linkIdx]; hipStream_t& stream = exeInfo.streams[transferIdx];
hipEvent_t& startEvent = exeInfo.startEvents[linkIdx]; hipEvent_t& startEvent = exeInfo.startEvents[transferIdx];
hipEvent_t& stopEvent = exeInfo.stopEvents[linkIdx]; hipEvent_t& stopEvent = exeInfo.stopEvents[transferIdx];
bool recordStart = (!ev.useSingleSync || iteration == 0 || ev.useSingleStream);
bool recordStop = (!ev.useSingleSync || iteration == ev.numIterations - 1 || ev.useSingleStream);
int const initOffset = ev.byteOffset / sizeof(float); int const initOffset = ev.byteOffset / sizeof(float);
if (ev.useHipCall) if (ev.useHipCall)
{ {
// Record start event // Record start event
if (recordStart) HIP_CALL(hipEventRecord(startEvent, stream)); HIP_CALL(hipEventRecord(startEvent, stream));
// Execute hipMemset / hipMemcpy // Execute hipMemset / hipMemcpy
if (ev.useMemset) if (ev.useMemset)
HIP_CALL(hipMemsetAsync(link.dstMem + initOffset, 42, N * sizeof(float), stream)); HIP_CALL(hipMemsetAsync(transfer.dstMem + initOffset, 42, N * sizeof(float), stream));
else else
HIP_CALL(hipMemcpyAsync(link.dstMem + initOffset, HIP_CALL(hipMemcpyAsync(transfer.dstMem + initOffset,
link.srcMem + initOffset, transfer.srcMem + initOffset,
N * sizeof(float), hipMemcpyDefault, N * sizeof(float), hipMemcpyDefault,
stream)); stream));
// Record stop event // Record stop event
if (recordStop) HIP_CALL(hipEventRecord(stopEvent, stream)); HIP_CALL(hipEventRecord(stopEvent, stream));
} }
else else
{ {
if (!ev.combineTiming && recordStart) HIP_CALL(hipEventRecord(startEvent, stream)); int const numBlocksToRun = ev.useSingleStream ? exeInfo.totalBlocks : transfer.numBlocksToUse;
int const numBlocksToRun = ev.useSingleStream ? exeInfo.totalBlocks : link.numBlocksToUse;
hipExtLaunchKernelGGL(ev.useMemset ? GpuMemsetKernel : GpuCopyKernel, hipExtLaunchKernelGGL(ev.useMemset ? GpuMemsetKernel : GpuCopyKernel,
dim3(numBlocksToRun, 1, 1), dim3(numBlocksToRun, 1, 1),
dim3(BLOCKSIZE, 1, 1), dim3(BLOCKSIZE, 1, 1),
ev.sharedMemBytes, stream, ev.sharedMemBytes, stream,
(ev.combineTiming && recordStart) ? startEvent : NULL, startEvent, stopEvent,
(ev.combineTiming && recordStop) ? stopEvent : NULL, 0, transfer.blockParamGpuPtr);
0, link.blockParamGpuPtr);
if (!ev.combineTiming & recordStop) HIP_CALL(hipEventRecord(stopEvent, stream));
} }
// Synchronize per iteration, unless in single sync mode, in which case // Synchronize per iteration, unless in single sync mode, in which case
// synchronize during last warmup / last actual iteration // synchronize during last warmup / last actual iteration
if (!ev.useSingleSync || iteration == -1 || iteration == ev.numIterations - 1) HIP_CALL(hipStreamSynchronize(stream));
{
HIP_CALL(hipStreamSynchronize(stream));
}
if (iteration >= 0) if (iteration >= 0)
{ {
// Record GPU timing // Record GPU timing
if (!ev.useSingleSync || iteration == ev.numIterations - 1 || ev.useSingleStream) float gpuDeltaMsec;
{ HIP_CALL(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent));
HIP_CALL(hipEventSynchronize(stopEvent));
float gpuDeltaMsec;
HIP_CALL(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent));
if (ev.useSingleStream) if (ev.useSingleStream)
{
for (Transfer& currTransfer : exeInfo.transfers)
{ {
for (Link& currLink : exeInfo.links) long long minStartCycle = currTransfer.blockParamGpuPtr[0].startCycle;
long long maxStopCycle = currTransfer.blockParamGpuPtr[0].stopCycle;
for (int i = 1; i < currTransfer.numBlocksToUse; i++)
{ {
long long minStartCycle = currLink.blockParamGpuPtr[0].startCycle; minStartCycle = std::min(minStartCycle, currTransfer.blockParamGpuPtr[i].startCycle);
long long maxStopCycle = currLink.blockParamGpuPtr[0].stopCycle; maxStopCycle = std::max(maxStopCycle, currTransfer.blockParamGpuPtr[i].stopCycle);
for (int i = 1; i < currLink.numBlocksToUse; i++)
{
minStartCycle = std::min(minStartCycle, currLink.blockParamGpuPtr[i].startCycle);
maxStopCycle = std::max(maxStopCycle, currLink.blockParamGpuPtr[i].stopCycle);
}
int const wallClockRate = GetWallClockRate(exeIndex);
double iterationTimeMs = (maxStopCycle - minStartCycle) / (double)(wallClockRate);
currLink.linkTime += iterationTimeMs;
} }
exeInfo.totalTime += gpuDeltaMsec; int const wallClockRate = GetWallClockRate(exeIndex);
} double iterationTimeMs = (maxStopCycle - minStartCycle) / (double)(wallClockRate);
else currTransfer.transferTime += iterationTimeMs;
{
link.linkTime += gpuDeltaMsec;
} }
exeInfo.totalTime += gpuDeltaMsec;
}
else
{
transfer.transferTime += gpuDeltaMsec;
} }
} }
} }
else if (link.exeMemType == MEM_CPU) // CPU execution agent else if (transfer.exeMemType == MEM_CPU) // CPU execution agent
{ {
// Force this thread and all child threads onto correct NUMA node // Force this thread and all child threads onto correct NUMA node
if (numa_run_on_node(link.exeIndex)) if (numa_run_on_node(transfer.exeIndex))
{ {
printf("[ERROR] Unable to set CPU to NUMA node %d\n", link.exeIndex); printf("[ERROR] Unable to set CPU to NUMA node %d\n", transfer.exeIndex);
exit(1); exit(1);
} }
...@@ -1079,18 +1073,18 @@ void RunLink(EnvVars const& ev, size_t const N, int const iteration, ExecutorInf ...@@ -1079,18 +1073,18 @@ void RunLink(EnvVars const& ev, size_t const N, int const iteration, ExecutorInf
auto cpuStart = std::chrono::high_resolution_clock::now(); auto cpuStart = std::chrono::high_resolution_clock::now();
// Launch child-threads to perform memcopies // Launch child-threads to perform memcopies
for (int i = 0; i < ev.numCpuPerLink; i++) for (int i = 0; i < ev.numCpuPerTransfer; i++)
childThreads.push_back(std::thread(ev.useMemset ? CpuMemsetKernel : CpuCopyKernel, std::ref(link.blockParam[i]))); childThreads.push_back(std::thread(ev.useMemset ? CpuMemsetKernel : CpuCopyKernel, std::ref(transfer.blockParam[i])));
// Wait for child-threads to finish // Wait for child-threads to finish
for (int i = 0; i < ev.numCpuPerLink; i++) for (int i = 0; i < ev.numCpuPerTransfer; i++)
childThreads[i].join(); childThreads[i].join();
auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
// Record time if not a warmup iteration // Record time if not a warmup iteration
if (iteration >= 0) if (iteration >= 0)
link.linkTime += (std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() * 1000.0); transfer.transferTime += (std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() * 1000.0);
} }
} }
...@@ -1110,7 +1104,7 @@ void RunPeerToPeerBenchmarks(EnvVars const& ev, size_t N, int numBlocksToUse, in ...@@ -1110,7 +1104,7 @@ void RunPeerToPeerBenchmarks(EnvVars const& ev, size_t N, int numBlocksToUse, in
if (!ev.outputToCsv) if (!ev.outputToCsv)
{ {
printf("Performing copies in each direction of %lu bytes\n", N * sizeof(float)); printf("Performing copies in each direction of %lu bytes\n", N * sizeof(float));
printf("Using %d threads per NUMA node for CPU copies\n", ev.numCpuPerLink); printf("Using %d threads per NUMA node for CPU copies\n", ev.numCpuPerTransfer);
printf("Using %d CUs per transfer\n", numBlocksToUse); printf("Using %d CUs per transfer\n", numBlocksToUse);
} }
else else
...@@ -1195,53 +1189,53 @@ double GetPeakBandwidth(EnvVars const& ev, ...@@ -1195,53 +1189,53 @@ double GetPeakBandwidth(EnvVars const& ev,
int const initOffset = ev.byteOffset / sizeof(float); int const initOffset = ev.byteOffset / sizeof(float);
// Prepare Links // Prepare Transfers
std::vector<Link*> links; std::vector<Transfer*> transfers;
ExecutorInfo exeInfo[2]; ExecutorInfo exeInfo[2];
for (int i = 0; i < 2; i++) for (int i = 0; i < 2; i++)
{ {
exeInfo[i].links.resize(1); exeInfo[i].transfers.resize(1);
exeInfo[i].streams.resize(1); exeInfo[i].streams.resize(1);
exeInfo[i].startEvents.resize(1); exeInfo[i].startEvents.resize(1);
exeInfo[i].stopEvents.resize(1); exeInfo[i].stopEvents.resize(1);
links.push_back(&exeInfo[i].links[0]); transfers.push_back(&exeInfo[i].transfers[0]);
} }
links[0]->srcMemType = links[1]->dstMemType = srcMemType; transfers[0]->srcMemType = transfers[1]->dstMemType = srcMemType;
links[0]->dstMemType = links[1]->srcMemType = dstMemType; transfers[0]->dstMemType = transfers[1]->srcMemType = dstMemType;
links[0]->srcIndex = links[1]->dstIndex = RemappedIndex(srcIndex, srcMemType); transfers[0]->srcIndex = transfers[1]->dstIndex = RemappedIndex(srcIndex, srcMemType);
links[0]->dstIndex = links[1]->srcIndex = RemappedIndex(dstIndex, dstMemType); transfers[0]->dstIndex = transfers[1]->srcIndex = RemappedIndex(dstIndex, dstMemType);
// Either perform (local read + remote write), or (remote read + local write) // Either perform (local read + remote write), or (remote read + local write)
links[0]->exeMemType = (readMode == 0 ? srcMemType : dstMemType); transfers[0]->exeMemType = (readMode == 0 ? srcMemType : dstMemType);
links[1]->exeMemType = (readMode == 0 ? dstMemType : srcMemType); transfers[1]->exeMemType = (readMode == 0 ? dstMemType : srcMemType);
links[0]->exeIndex = RemappedIndex((readMode == 0 ? srcIndex : dstIndex), links[0]->exeMemType); transfers[0]->exeIndex = RemappedIndex((readMode == 0 ? srcIndex : dstIndex), transfers[0]->exeMemType);
links[1]->exeIndex = RemappedIndex((readMode == 0 ? dstIndex : srcIndex), links[1]->exeMemType); transfers[1]->exeIndex = RemappedIndex((readMode == 0 ? dstIndex : srcIndex), transfers[1]->exeMemType);
for (int i = 0; i <= isBidirectional; i++) for (int i = 0; i <= isBidirectional; i++)
{ {
AllocateMemory(links[i]->srcMemType, links[i]->srcIndex, AllocateMemory(transfers[i]->srcMemType, transfers[i]->srcIndex,
N * sizeof(float) + ev.byteOffset, (void**)&links[i]->srcMem); N * sizeof(float) + ev.byteOffset, (void**)&transfers[i]->srcMem);
AllocateMemory(links[i]->dstMemType, links[i]->dstIndex, AllocateMemory(transfers[i]->dstMemType, transfers[i]->dstIndex,
N * sizeof(float) + ev.byteOffset, (void**)&links[i]->dstMem); N * sizeof(float) + ev.byteOffset, (void**)&transfers[i]->dstMem);
// Prepare block parameters on CPU // Prepare block parameters on CPU
links[i]->numBlocksToUse = (links[i]->exeMemType == MEM_GPU) ? numBlocksToUse : ev.numCpuPerLink; transfers[i]->numBlocksToUse = (transfers[i]->exeMemType == MEM_GPU) ? numBlocksToUse : ev.numCpuPerTransfer;
links[i]->blockParam.resize(links[i]->numBlocksToUse); transfers[i]->blockParam.resize(transfers[i]->numBlocksToUse);
links[i]->PrepareBlockParams(ev, N); transfers[i]->PrepareBlockParams(ev, N);
if (links[i]->exeMemType == MEM_GPU) if (transfers[i]->exeMemType == MEM_GPU)
{ {
// Copy block parameters onto GPU // Copy block parameters onto GPU
AllocateMemory(MEM_GPU, links[i]->exeIndex, numBlocksToUse * sizeof(BlockParam), AllocateMemory(MEM_GPU, transfers[i]->exeIndex, numBlocksToUse * sizeof(BlockParam),
(void **)&links[i]->blockParamGpuPtr); (void **)&transfers[i]->blockParamGpuPtr);
HIP_CALL(hipMemcpy(links[i]->blockParamGpuPtr, HIP_CALL(hipMemcpy(transfers[i]->blockParamGpuPtr,
links[i]->blockParam.data(), transfers[i]->blockParam.data(),
numBlocksToUse * sizeof(BlockParam), numBlocksToUse * sizeof(BlockParam),
hipMemcpyHostToDevice)); hipMemcpyHostToDevice));
// Prepare GPU resources // Prepare GPU resources
HIP_CALL(hipSetDevice(links[i]->exeIndex)); HIP_CALL(hipSetDevice(transfers[i]->exeIndex));
HIP_CALL(hipStreamCreate(&exeInfo[i].streams[0])); HIP_CALL(hipStreamCreate(&exeInfo[i].streams[0]));
HIP_CALL(hipEventCreate(&exeInfo[i].startEvents[0])); HIP_CALL(hipEventCreate(&exeInfo[i].startEvents[0]));
HIP_CALL(hipEventCreate(&exeInfo[i].stopEvents[0])); HIP_CALL(hipEventCreate(&exeInfo[i].stopEvents[0]));
...@@ -1255,7 +1249,7 @@ double GetPeakBandwidth(EnvVars const& ev, ...@@ -1255,7 +1249,7 @@ double GetPeakBandwidth(EnvVars const& ev,
{ {
// Perform timed iterations // Perform timed iterations
for (int i = 0; i <= isBidirectional; i++) for (int i = 0; i <= isBidirectional; i++)
threads.push(std::thread(RunLink, std::ref(ev), N, iteration, std::ref(exeInfo[i]), 0)); threads.push(std::thread(RunTransfer, std::ref(ev), N, iteration, std::ref(exeInfo[i]), 0));
// Wait for all threads to finish // Wait for all threads to finish
for (int i = 0; i <= isBidirectional; i++) for (int i = 0; i <= isBidirectional; i++)
...@@ -1265,28 +1259,28 @@ double GetPeakBandwidth(EnvVars const& ev, ...@@ -1265,28 +1259,28 @@ double GetPeakBandwidth(EnvVars const& ev,
} }
} }
// Validate that each link has transferred correctly // Validate that each Transfer has transferred correctly
for (int i = 0; i <= isBidirectional; i++) for (int i = 0; i <= isBidirectional; i++)
CheckOrFill(MODE_CHECK, N, ev.useMemset, ev.useHipCall, ev.fillPattern, links[i]->dstMem + initOffset); CheckOrFill(MODE_CHECK, N, ev.useMemset, ev.useHipCall, ev.fillPattern, transfers[i]->dstMem + initOffset);
// Collect aggregate bandwidth // Collect aggregate bandwidth
double totalBandwidth = 0; double totalBandwidth = 0;
for (int i = 0; i <= isBidirectional; i++) for (int i = 0; i <= isBidirectional; i++)
{ {
double linkDurationMsec = links[i]->linkTime / (1.0 * ev.numIterations); double transferDurationMsec = transfers[i]->transferTime / (1.0 * ev.numIterations);
double linkBandwidthGbs = (N * sizeof(float) / 1.0E9) / linkDurationMsec * 1000.0f; double transferBandwidthGbs = (N * sizeof(float) / 1.0E9) / transferDurationMsec * 1000.0f;
totalBandwidth += linkBandwidthGbs; totalBandwidth += transferBandwidthGbs;
} }
// Release GPU memory // Release GPU memory
for (int i = 0; i <= isBidirectional; i++) for (int i = 0; i <= isBidirectional; i++)
{ {
DeallocateMemory(links[i]->srcMemType, links[i]->srcMem); DeallocateMemory(transfers[i]->srcMemType, transfers[i]->srcMem);
DeallocateMemory(links[i]->dstMemType, links[i]->dstMem); DeallocateMemory(transfers[i]->dstMemType, transfers[i]->dstMem);
if (links[i]->exeMemType == MEM_GPU) if (transfers[i]->exeMemType == MEM_GPU)
{ {
DeallocateMemory(MEM_GPU, links[i]->blockParamGpuPtr); DeallocateMemory(MEM_GPU, transfers[i]->blockParamGpuPtr);
HIP_CALL(hipStreamDestroy(exeInfo[i].streams[0])); HIP_CALL(hipStreamDestroy(exeInfo[i].streams[0]));
HIP_CALL(hipEventDestroy(exeInfo[i].startEvents[0])); HIP_CALL(hipEventDestroy(exeInfo[i].startEvents[0]));
HIP_CALL(hipEventDestroy(exeInfo[i].stopEvents[0])); HIP_CALL(hipEventDestroy(exeInfo[i].stopEvents[0]));
...@@ -1295,7 +1289,7 @@ double GetPeakBandwidth(EnvVars const& ev, ...@@ -1295,7 +1289,7 @@ double GetPeakBandwidth(EnvVars const& ev,
return totalBandwidth; return totalBandwidth;
} }
void Link::PrepareBlockParams(EnvVars const& ev, size_t const N) void Transfer::PrepareBlockParams(EnvVars const& ev, size_t const N)
{ {
int const initOffset = ev.byteOffset / sizeof(float); int const initOffset = ev.byteOffset / sizeof(float);
...@@ -1303,7 +1297,7 @@ void Link::PrepareBlockParams(EnvVars const& ev, size_t const N) ...@@ -1303,7 +1297,7 @@ void Link::PrepareBlockParams(EnvVars const& ev, size_t const N)
CheckOrFill(MODE_FILL, N, ev.useMemset, ev.useHipCall, ev.fillPattern, this->srcMem + initOffset); CheckOrFill(MODE_FILL, N, ev.useMemset, ev.useHipCall, ev.fillPattern, this->srcMem + initOffset);
// Each block needs to know src/dst pointers and how many elements to transfer // Each block needs to know src/dst pointers and how many elements to transfer
// Figure out the sub-array each block does for this Link // Figure out the sub-array each block does for this Transfer
// - Partition N as evenly as possible, but try to keep blocks as multiples of BLOCK_BYTES bytes, // - Partition N as evenly as possible, but try to keep blocks as multiples of BLOCK_BYTES bytes,
// except the very last one, for alignment reasons // except the very last one, for alignment reasons
int const targetMultiple = ev.blockBytes / sizeof(float); int const targetMultiple = ev.blockBytes / sizeof(float);
...@@ -1324,7 +1318,7 @@ void Link::PrepareBlockParams(EnvVars const& ev, size_t const N) ...@@ -1324,7 +1318,7 @@ void Link::PrepareBlockParams(EnvVars const& ev, size_t const N)
assigned += param.N; assigned += param.N;
} }
this->linkTime = 0.0; this->transferTime = 0.0;
} }
// NOTE: This is a stop-gap solution until HIP provides wallclock values // NOTE: This is a stop-gap solution until HIP provides wallclock values
......
TransferBench.hpp
View file @ ddb6508f
...@@ -37,8 +37,6 @@ THE SOFTWARE. ...@@ -37,8 +37,6 @@ THE SOFTWARE.
#include "EnvVars.hpp" #include "EnvVars.hpp"
#define TB_VERSION "1.01"
// Helper macro for catching HIP errors // Helper macro for catching HIP errors
#define HIP_CALL(cmd) \ #define HIP_CALL(cmd) \
do { \ do { \
...@@ -52,7 +50,7 @@ THE SOFTWARE. ...@@ -52,7 +50,7 @@ THE SOFTWARE.
} while (0) } while (0)
// Simple configuration parameters // Simple configuration parameters
size_t const DEFAULT_BYTES_PER_LINK = (1<<26); // Amount of data transferred per Link size_t const DEFAULT_BYTES_PER_TRANSFER = (1<<26); // Amount of data transferred per Transfer
// Different src/dst memory types supported // Different src/dst memory types supported
typedef enum typedef enum
...@@ -81,19 +79,19 @@ struct BlockParam ...@@ -81,19 +79,19 @@ struct BlockParam
long long stopCycle; long long stopCycle;
}; };
// Each Link is a uni-direction operation from a src memory to dst memory // Each Transfer is a uni-direction operation from a src memory to dst memory
struct Link struct Transfer
{ {
int linkIndex; // Link identifier int transferIndex; // Transfer identifier
// Link config // Transfer config
MemType exeMemType; // Link executor type (CPU or GPU) MemType exeMemType; // Transfer executor type (CPU or GPU)
int exeIndex; // Executor index (NUMA node for CPU / device ID for GPU) int exeIndex; // Executor index (NUMA node for CPU / device ID for GPU)
MemType srcMemType; // Source memory type MemType srcMemType; // Source memory type
int srcIndex; // Source device index int srcIndex; // Source device index
MemType dstMemType; // Destination memory type MemType dstMemType; // Destination memory type
int dstIndex; // Destination device index int dstIndex; // Destination device index
int numBlocksToUse; // Number of threadblocks to use for this Link int numBlocksToUse; // Number of threadblocks to use for this Transfer
// Memory // Memory
float* srcMem; // Source memory float* srcMem; // Source memory
...@@ -104,7 +102,7 @@ struct Link ...@@ -104,7 +102,7 @@ struct Link
BlockParam* blockParamGpuPtr; BlockParam* blockParamGpuPtr;
// Results // Results
double linkTime; double transferTime;
// Prepares src memory and how to divide N elements across threadblocks/threads // Prepares src memory and how to divide N elements across threadblocks/threads
void PrepareBlockParams(EnvVars const& ev, size_t const N); void PrepareBlockParams(EnvVars const& ev, size_t const N);
...@@ -114,7 +112,7 @@ typedef std::pair<MemType, int> Executor; ...@@ -114,7 +112,7 @@ typedef std::pair<MemType, int> Executor;
struct ExecutorInfo struct ExecutorInfo
{ {
std::vector<Link> links; // Links to execute std::vector<Transfer> transfers; // Transfers to execute
// For GPU-Executors // For GPU-Executors
int totalBlocks; // Total number of CUs/CPU threads to use int totalBlocks; // Total number of CUs/CPU threads to use
...@@ -127,7 +125,7 @@ struct ExecutorInfo ...@@ -127,7 +125,7 @@ struct ExecutorInfo
double totalTime; double totalTime;
}; };
typedef std::map<Executor, ExecutorInfo> LinkMap; typedef std::map<Executor, ExecutorInfo> TransferMap;
// Display usage instructions // Display usage instructions
void DisplayUsage(char const* cmdName); void DisplayUsage(char const* cmdName);
...@@ -136,21 +134,21 @@ void DisplayUsage(char const* cmdName); ...@@ -136,21 +134,21 @@ void DisplayUsage(char const* cmdName);
void DisplayTopology(bool const outputToCsv); void DisplayTopology(bool const outputToCsv);
// Build array of test sizes based on sampling factor // Build array of test sizes based on sampling factor
void PopulateTestSizes(size_t const numBytesPerLink, int const samplingFactor, void PopulateTestSizes(size_t const numBytesPerTransfer, int const samplingFactor,
std::vector<size_t>& valuesofN); std::vector<size_t>& valuesofN);
void ParseMemType(std::string const& token, int const numCpus, int const numGpus, void ParseMemType(std::string const& token, int const numCpus, int const numGpus,
MemType* memType, int* memIndex); MemType* memType, int* memIndex);
void ParseLinks(char* line, int numCpus, int numGpus, void ParseTransfers(char* line, int numCpus, int numGpus,
LinkMap& linkMap); TransferMap& transferMap);
void EnablePeerAccess(int const deviceId, int const peerDeviceId); void EnablePeerAccess(int const deviceId, int const peerDeviceId);
void AllocateMemory(MemType memType, int devIndex, size_t numBytes, void** memPtr); void AllocateMemory(MemType memType, int devIndex, size_t numBytes, void** memPtr);
void DeallocateMemory(MemType memType, void* memPtr); void DeallocateMemory(MemType memType, void* memPtr);
void CheckPages(char* byteArray, size_t numBytes, int targetId); void CheckPages(char* byteArray, size_t numBytes, int targetId);
void CheckOrFill(ModeType mode, int N, bool isMemset, bool isHipCall, std::vector<float> const& fillPattern, float* ptr); void CheckOrFill(ModeType mode, int N, bool isMemset, bool isHipCall, std::vector<float> const& fillPattern, float* ptr);
void RunLink(EnvVars const& ev, size_t const N, int const iteration, ExecutorInfo& exeInfo, int const linkIdx); void RunTransfer(EnvVars const& ev, size_t const N, int const iteration, ExecutorInfo& exeInfo, int const transferIdx);
void RunPeerToPeerBenchmarks(EnvVars const& ev, size_t N, int numBlocksToUse, int readMode, int skipCpu); void RunPeerToPeerBenchmarks(EnvVars const& ev, size_t N, int numBlocksToUse, int readMode, int skipCpu);
// Return the maximum bandwidth measured for given (src/dst) pair // Return the maximum bandwidth measured for given (src/dst) pair
...@@ -167,6 +165,6 @@ double GetPeakBandwidth(EnvVars const& ev, ...@@ -167,6 +165,6 @@ double GetPeakBandwidth(EnvVars const& ev,
std::string GetLinkTypeDesc(uint32_t linkType, uint32_t hopCount); std::string GetLinkTypeDesc(uint32_t linkType, uint32_t hopCount);
std::string GetDesc(MemType srcMemType, int srcIndex, std::string GetDesc(MemType srcMemType, int srcIndex,
MemType dstMemType, int dstIndex); MemType dstMemType, int dstIndex);
std::string GetLinkDesc(Link const& link); std::string GetTransferDesc(Transfer const& transfer);
int RemappedIndex(int const origIdx, MemType const memType); int RemappedIndex(int const origIdx, MemType const memType);
int GetWallClockRate(int deviceId); int GetWallClockRate(int deviceId);
example.cfg
View file @ ddb6508f
# ConfigFile Format: # ConfigFile Format:
# ================== # ==================
# A Link is defined as a uni-directional transfer from src memory location to dst memory location # A Transfer is defined as a uni-directional transfer from src memory location to dst memory location
# executed by either CPU or GPU # executed by either CPU or GPU
# Each single line in the configuration file defines a set of Links to run in parallel # Each single line in the configuration file defines a set of Transfers (a Test) to run in parallel
# There are two ways to specify the configuration file: # There are two ways to specify the configuration file:
# 1) Basic # 1) Basic
# The basic specification assumes the same number of threadblocks/CUs used per GPU-executed Link # The basic specification assumes the same number of threadblocks/CUs used per GPU-executed Transfer
# A positive number of Links is specified followed by that number of triplets describing each Link # A positive number of Transfers is specified followed by that number of triplets describing each Transfer
# #Links #CUs (srcMem1->Executor1->dstMem1) ... (srcMemL->ExecutorL->dstMemL) # #Transfers #CUs (srcMem1->Executor1->dstMem1) ... (srcMemL->ExecutorL->dstMemL)
# 2) Advanced # 2) Advanced
# The advanced specification allows different number of threadblocks/CUs used per GPU-executed Link # The advanced specification allows different number of threadblocks/CUs used per GPU-executed Transfer
# A negative number of links is specified, followed by quadruples describing each Link # A negative number of Transfers is specified, followed by quadruples describing each Transfer
# -#Links (srcMem1->Executor1->dstMem1 #CUs1) ... (srcMemL->ExecutorL->dstMemL #CUsL) # -#Transfers (srcMem1->Executor1->dstMem1 #CUs1) ... (srcMemL->ExecutorL->dstMemL #CUsL)
# Argument Details: # Argument Details:
# #Links : Number of Links to be run in parallel # #Transfers: Number of Transfers to be run in parallel
# #CUs : Number of threadblocks/CUs to use for a GPU-executed Link # #CUs : Number of threadblocks/CUs to use for a GPU-executed Transfer
# srcMemL : Source memory location (Where the data is to be read from). Ignored in memset mode # srcMemL : Source memory location (Where the data is to be read from). Ignored in memset mode
# Executor: Executor are specified by a character indicating executor type, followed by device index (0-indexed) # Executor : Executor is specified by a character indicating type, followed by device index (0-indexed)
# - C: CPU-executed (Indexed from 0 to 1) # - C: CPU-executed (Indexed from 0 to # NUMA nodes - 1)
# - G: GPU-executed (Indexed from 0 to 3) # - G: GPU-executed (Indexed from 0 to # GPUs - 1)
# dstMemL : Destination memory location (Where the data is to be written to) # dstMemL : Destination memory location (Where the data is to be written to)
# Memory locations are specified by a character indicating memory type, # Memory locations are specified by a character indicating memory type,
# followed by device index (0-indexed) # followed by device index (0-indexed)
# Supported memory locations are: # Supported memory locations are:
# - C: Pinned host memory (on NUMA node, indexed from 0 to [# NUMA nodes-1]) # - C: Pinned host memory (on NUMA node, indexed from 0 to [# NUMA nodes-1])
# - B: Fine-grain host memory (on NUMA node, indexed from 0 to [# NUMA nodes-1]) # - B: Fine-grain host memory (on NUMA node, indexed from 0 to [# NUMA nodes-1])
# - G: Global device memory (on GPU device indexed from 0 to [# GPUs - 1]) # - G: Global device memory (on GPU device indexed from 0 to [# GPUs - 1])
# - F: Fine-grain device memory (on GPU device indexed from 0 to [# GPUs - 1]) # - F: Fine-grain device memory (on GPU device indexed from 0 to [# GPUs - 1])
# Examples: # Examples:
# 1 4 (G0->G0->G1) Single link using 4 CUs on GPU0 to copy from GPU0 to GPU1 # 1 4 (G0->G0->G1) Single Transfer using 4 CUs on GPU0 to copy from GPU0 to GPU1
# 1 4 (C1->G2->G0) Single link using 4 CUs on GPU2 to copy from CPU1 to GPU0 # 1 4 (C1->G2->G0) Single Transfer using 4 CUs on GPU2 to copy from CPU1 to GPU0
# 2 4 G0->G0->G1 G1->G1->G0 Runs 2 Links in parallel. GPU0 to GPU1, and GPU1 to GPU0, each with 4 CUs # 2 4 G0->G0->G1 G1->G1->G0 Runs 2 Transfers in parallel. GPU0 to GPU1, and GPU1 to GPU0, each with 4 CUs
# -2 (G0 G0 G1 4) (G1 G1 G0 2) Runs 2 Links in parallel. GPU0 to GPU1 with 4 CUs, and GPU1 to GPU0 with 2 CUs # -2 (G0 G0 G1 4) (G1 G1 G0 2) Runs 2 Transfers in parallel. GPU0 to GPU1 with 4 CUs, and GPU1 to GPU0 with 2 CUs
# Round brackets and arrows' ->' may be included for human clarity, but will be ignored and are unnecessary # Round brackets and arrows' ->' may be included for human clarity, but will be ignored and are unnecessary
# Lines starting with # will be ignored. Lines starting with ## will be echoed to output # Lines starting with # will be ignored. Lines starting with ## will be echoed to output
# Single GPU-executed link between GPUs 0 and 1 using 4 CUs # Single GPU-executed Transfer between GPUs 0 and 1 using 4 CUs
1 4 (G0->G0->G1) 1 4 (G0->G0->G1)
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment