API#

struct AgentData#

Public Members

bool isInitialized#
std::vector<hsa_agent_t> cpuAgents#
std::vector<hsa_agent_t> gpuAgents#
std::vector<int> closestNumaNode#
class EnvVars#
#include <EnvVars.hpp>

Public Functions

inline EnvVars()#
inline void DisplayEnvVars() const#
inline void DisplayP2PBenchmarkEnvVars() const#
inline void DisplaySweepEnvVars() const#
inline void DisplayA2AEnvVars() const#
inline void DisplaySchmooEnvVars() const#
inline void DisplayRemoteWriteEnvVars() const#
inline std::string GetCuMaskDesc() const#

Public Members

int const DEFAULT_NUM_WARMUPS = 3#
int const DEFAULT_NUM_ITERATIONS = 10#
int const DEFAULT_SAMPLING_FACTOR = 1#
int const DEFAULT_P2P_NUM_CPU_SE = 4#
std::string const DEFAULT_SWEEP_SRC = "CG"#
std::string const DEFAULT_SWEEP_EXE = "CDG"#
std::string const DEFAULT_SWEEP_DST = "CG"#
int const DEFAULT_SWEEP_MIN = 1#
int const DEFAULT_SWEEP_MAX = 24#
int const DEFAULT_SWEEP_TEST_LIMIT = 0#
int const DEFAULT_SWEEP_TIME_LIMIT = 0#
int alwaysValidate#
int blockBytes#
int blockOrder#
int byteOffset#
int continueOnError#
int gfxBlockSize#
int gfxSingleTeam#
int gfxUnroll#
int gfxWaveOrder#
int hideEnv#
int numCpuDevices#
int numGpuDevices#
int numIterations#
int numWarmups#
int outputToCsv#
int samplingFactor#
int sharedMemBytes#
int showIterations#
int useInteractive#
int usePcieIndexing#
int usePrepSrcKernel#
int useSingleStream#
int useXccFilter#
int validateDirect#
std::vector<float> fillPattern#
std::vector<uint32_t> cuMask#
std::vector<std::vector<int>> prefXccTable#
int numCpuSubExecs#
int numGpuSubExecs#
int p2pMode#
int useDmaCopy#
int useRemoteRead#
int useFineGrain#
int sweepMin#
int sweepMax#
int sweepTestLimit#
int sweepTimeLimit#
int sweepXgmiMin#
int sweepXgmiMax#
int sweepSeed#
int sweepRandBytes#
std::string sweepSrc#
std::string sweepExe#
std::string sweepDst#
int a2aDirect#
int a2aMode#
int enableDebug#
ConfigModeEnum configMode#
std::default_random_engine *generator#
std::vector<int> numCpusPerNuma#
std::vector<int> wallClockPerDeviceMhz#
std::vector<std::set<int>> xccIdsPerDevice#

Public Static Functions

static inline void DisplayUsage()#
static inline int GetEnvVar(std::string const &varname, int defaultValue)#
static inline std::string GetEnvVar(std::string const &varname, std::string const &defaultValue)#
struct ExecutorInfo#
#include <TransferBench.hpp>

Public Members

std::vector<Transfer*> transfers#
size_t totalBytes#
int totalSubExecs#
SubExecParam *subExecParamGpu#
std::vector<hipStream_t> streams#
std::vector<hipEvent_t> startEvents#
std::vector<hipEvent_t> stopEvents#
double totalTime#
struct SubExecParam#
#include <Kernels.hpp>

Public Members

size_t N#
int numSrcs#
int numDsts#
float *src[MAX_SRCS]#
float *dst[MAX_DSTS]#
int32_t preferredXccId#
int teamSize#
int teamIdx#
long long startCycle#
long long stopCycle#
uint32_t hwId#
uint32_t xccId#
struct Transfer#
#include <TransferBench.hpp>

Public Functions

void PrepareSubExecParams(EnvVars const &ev)#
bool PrepareSrc(EnvVars const &ev)#
void ValidateDst(EnvVars const &ev)#
void PrepareReference(EnvVars const &ev, std::vector<float> &buffer, int bufferIdx)#
std::string SrcToStr() const#
std::string DstToStr() const#

Public Members

ExeType exeType#
int exeIndex#
int exeSubIndex#
int numSubExecs#
size_t numBytes#
int numSrcs#
std::vector<MemType> srcType#
std::vector<int> srcIndex#
int numDsts#
std::vector<MemType> dstType#
std::vector<int> dstIndex#
size_t numBytesActual#
double transferTime#
double transferBandwidth#
double executorBandwidth#
std::vector<double> perIterationTime#
std::vector<std::set<std::pair<int, int>>> perIterationCUs#
int transferIndex#
std::vector<float*> srcMem#
std::vector<float*> dstMem#
std::vector<SubExecParam> subExecParam#
SubExecParam *subExecParamGpuPtr#
std::vector<int> subExecIdx#
namespace std#

STL namespace.

file Compatibility.hpp
#include <hip/hip_ext.h>
#include <hip/hip_runtime.h>
#include <hsa/hsa_ext_amd.h>
file EnvVars.hpp
#include <algorithm>
#include <random>
#include <time.h>
#include “Compatibility.hpp
#include “Kernels.hpp

Defines

TB_VERSION#
PRINT_EV(NAME, VALUE, DESCRIPTION)#
PRINT_ES(NAME, VALUE, DESCRIPTION)#

Enums

enum ConfigModeEnum#

Values:

enumerator CFG_FILE#
enumerator CFG_P2P#
enumerator CFG_SWEEP#
enumerator CFG_SCALE#
enumerator CFG_A2A#
enumerator CFG_SCHMOO#
enumerator CFG_RWRITE#
enum BlockOrderEnum#

Values:

enumerator ORDER_SEQUENTIAL#
enumerator ORDER_INTERLEAVED#
enumerator ORDER_RANDOM#

Variables

char const MemTypeStr[]#
char const ExeTypeStr[]#
file GetClosestNumaNode.hpp

Defines

HSA_CHECK(cmd)#

Functions

hsa_status_t MemPoolInfoCallback(hsa_amd_memory_pool_t pool, void *data)#
hsa_status_t AgentInfoCallback(hsa_agent_t agent, void *data)#
AgentData &GetAgentData()#
int GetClosestNumaNode(int gpuIdx)#
file Kernels.hpp

Defines

PackedFloat_t#
MAX_BLOCKSIZE#
FLOATS_PER_PACK#
MEMSET_CHAR#
MEMSET_VAL#
MAX_WAVEGROUPS#
MAX_UNROLL#
NUM_WAVEORDERS#
MAX_SRCS#
MAX_DSTS#
__trace_hwreg()#
GetXccId(val)#
GPU_KERNEL_UNROLL_DECL(BLOCKSIZE)#

Typedefs

typedef void (*GpuKernelFuncPtr)(SubExecParam*, int)#

Functions

void CpuReduceKernel(SubExecParam const &p)#
std::string PrepSrcValueString()#
__host__ __device__ float PrepSrcValue (int srcBufferIdx, size_t idx)
__global__ void CollectXccIdsKernel (int *xccIds)
__global__ void PrepSrcDataKernel (float *ptr, size_t N, int srcBufferIdx)
__device__ int64_t GetTimestamp ()
template<typename T> __device__ __forceinline__ T MemsetVal ()
template<> __device__ __forceinline__ float MemsetVal ()
template<int BLOCKSIZE, int UNROLL> __global__ void __launch_bounds__ (BLOCKSIZE) GpuReduceKernel(SubExecParam *params
if (threadIdx.x==0) startCycle
GetXccId(xccId)#
if (p.preferredXccId !=-1 &&xccId !=p.preferredXccId) return
for (int i=0;i< numSrcs;i++) srcFloat4[i]
switch (waveOrder)
if (numSrcs==0)

Variables

__global__ void int waveOrder  {int64_t startCycle
SubExecParam &p = params[blockIdx.y]#
int32_t xccId#
int32_t const numSrcs = p.numSrcs#
int32_t const numDsts = p.numDsts#
float4 const  *__restrict__ srcFloat4 [MAX_SRCS]
float4 *__restrict__ dstFloat4 [MAX_DSTS]
int32_t const nTeams = p.teamSize#
int32_t const teamIdx = p.teamIdx#
int32_t const nWaves = BLOCKSIZE / warpSize#
int32_t const waveIdx = threadIdx.x / warpSize#
int32_t const tIdx = threadIdx.x % warpSize#
size_t const numFloat4 = p.N / 4#
int32_t teamStride#
int32_t waveStride#
int32_t unrlStride#
int32_t teamStride2#
int32_t waveStride2#
size_t const loop1Stride = nTeams * nWaves * UNROLL * warpSize#
size_t const loop1Limit = numFloat4 / loop1Stride * loop1Stride#
GpuKernelFuncPtr GpuKernelTable[MAX_WAVEGROUPS][MAX_UNROLL] = {GPU_KERNEL_UNROLL_DECL(64), GPU_KERNEL_UNROLL_DECL(128), GPU_KERNEL_UNROLL_DECL(192), GPU_KERNEL_UNROLL_DECL(256), GPU_KERNEL_UNROLL_DECL(320), GPU_KERNEL_UNROLL_DECL(384), GPU_KERNEL_UNROLL_DECL(448),}#
file TransferBench.hpp
#include <vector>
#include <sstream>
#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <cstdint>
#include <set>
#include <unistd.h>
#include <map>
#include <iostream>
#include “Compatibility.hpp
#include “EnvVars.hpp

Defines

HIP_CALL(cmd)#
MAX_LINE_LEN#

Typedefs

typedef std::pair<ExeType, int> Executor#
typedef std::map<Executor, ExecutorInfo> TransferMap#

Enums

enum MemType#

Values:

enumerator MEM_CPU#
enumerator MEM_GPU#
enumerator MEM_CPU_FINE#
enumerator MEM_GPU_FINE#
enumerator MEM_CPU_UNPINNED#
enumerator MEM_NULL#
enum ExeType#

Values:

enumerator EXE_CPU#
enumerator EXE_GPU_GFX#
enumerator EXE_GPU_DMA#

Functions

bool IsGpuType(MemType m)#
bool IsCpuType(MemType m)#
bool IsGpuType(ExeType e)#
bool IsCpuType(ExeType e)#
inline MemType CharToMemType(char const c)#
inline ExeType CharToExeType(char const c)#
void DisplayUsage(char const *cmdName)#
void DisplayTopology(bool const outputToCsv)#
void PopulateTestSizes(size_t const numBytesPerTransfer, int const samplingFactor, std::vector<size_t> &valuesofN)#
void ParseMemType(EnvVars const &ev, std::string const &token, std::vector<MemType> &memType, std::vector<int> &memIndex)#
void ParseExeType(EnvVars const &ev, std::string const &token, ExeType &exeType, int &exeIndex, int &exeSubIndex)#
void ParseTransfers(EnvVars const &ev, char *line, std::vector<Transfer> &transfers)#
void ExecuteTransfers(EnvVars const &ev, int const testNum, size_t const N, std::vector<Transfer> &transfers, bool verbose = true, double *totalBandwidthCpu = nullptr)#
void EnablePeerAccess(int const deviceId, int const peerDeviceId)#
void AllocateMemory(MemType memType, int devIndex, size_t numBytes, void **memPtr)#
void DeallocateMemory(MemType memType, void *memPtr, size_t const size = 0)#
void CheckPages(char *byteArray, size_t numBytes, int targetId)#
void RunTransfer(EnvVars const &ev, int const iteration, ExecutorInfo &exeInfo, int const transferIdx)#
void RunPeerToPeerBenchmarks(EnvVars const &ev, size_t N)#
void RunScalingBenchmark(EnvVars const &ev, size_t N, int const exeIndex, int const maxSubExecs)#
void RunSweepPreset(EnvVars const &ev, size_t const numBytesPerTransfer, int const numGpuSubExec, int const numCpuSubExec, bool const isRandom)#
void RunAllToAllBenchmark(EnvVars const &ev, size_t const numBytesPerTransfer, int const numSubExecs)#
void RunSchmooBenchmark(EnvVars const &ev, size_t const numBytesPerTransfer, int const localIdx, int const remoteIdx, int const maxSubExecs)#
void RunRemoteWriteBenchmark(EnvVars const &ev, size_t const numBytesPerTransfer, int numSubExecs, int const srcIdx, int minGpus, int maxGpus)#
std::string GetLinkTypeDesc(uint32_t linkType, uint32_t hopCount)#
int RemappedIndex(int const origIdx, bool const isCpuType)#
void LogTransfers(FILE *fp, int const testNum, std::vector<Transfer> const &transfers)#
std::string PtrVectorToStr(std::vector<float*> const &strVector, int const initOffset)#

Variables

size_t const DEFAULT_BYTES_PER_TRANSFER = (1 << 26)#
char const MemTypeStr[7] = "CGBFUN"
char const ExeTypeStr[4] = "CGD"
char const ExeTypeName[3][4] = {"CPU", "GPU", "DMA"}#
file TransferBench.cpp
#include <numa.h>
#include <cmath>
#include <numaif.h>
#include <random>
#include <stack>
#include <thread>
#include “TransferBench.hpp
#include “GetClosestNumaNode.hpp

Functions

int main(int argc, char **argv)#
void ExecuteTransfers(EnvVars const &ev, int const testNum, size_t const N, std::vector<Transfer> &transfers, bool verbose, double *totalBandwidthCpu)
void DisplayUsage(char const *cmdName)
int RemappedIndex(int const origIdx, bool const isCpuType)
void DisplayTopology(bool const outputToCsv)
void ParseMemType(EnvVars const &ev, std::string const &token, std::vector<MemType> &memTypes, std::vector<int> &memIndices)
void ParseExeType(EnvVars const &ev, std::string const &token, ExeType &exeType, int &exeIndex, int &exeSubIndex)
void ParseTransfers(EnvVars const &ev, char *line, std::vector<Transfer> &transfers)
void EnablePeerAccess(int const deviceId, int const peerDeviceId)
void AllocateMemory(MemType memType, int devIndex, size_t numBytes, void **memPtr)
void DeallocateMemory(MemType memType, void *memPtr, size_t const bytes)
void CheckPages(char *array, size_t numBytes, int targetId)
uint32_t GetId(uint32_t hwId)#
void RunTransfer(EnvVars const &ev, int const iteration, ExecutorInfo &exeInfo, int const transferIdx)
void RunPeerToPeerBenchmarks(EnvVars const &ev, size_t N)
void RunScalingBenchmark(EnvVars const &ev, size_t N, int const exeIndex, int const maxSubExecs)
void RunAllToAllBenchmark(EnvVars const &ev, size_t const numBytesPerTransfer, int const numSubExecs)
void RunSchmooBenchmark(EnvVars const &ev, size_t const numBytesPerTransfer, int const localIdx, int const remoteIdx, int const maxSubExecs)
void RunRemoteWriteBenchmark(EnvVars const &ev, size_t const numBytesPerTransfer, int numSubExecs, int const srcIdx, int minGpus, int maxGpus)
void RunSweepPreset(EnvVars const &ev, size_t const numBytesPerTransfer, int const numGpuSubExecs, int const numCpuSubExecs, bool const isRandom)
void LogTransfers(FILE *fp, int const testNum, std::vector<Transfer> const &transfers)
std::string PtrVectorToStr(std::vector<float*> const &strVector, int const initOffset)
dir /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-transferbench/checkouts/latest/src/include
dir /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-transferbench/checkouts/latest/src