API

Contents

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 void DisplayParallelCopyEnvVars() 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
int gpuMaxHwQueues
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
hsa_agent_t dstAgent
hsa_agent_t srcAgent
hsa_signal_t signal
hsa_amd_sdma_engine_id_t sdmaEngineId
namespace std

STL namespace.

file Compatibility.hpp
#include <hip/hip_ext.h>
#include <hip/hip_runtime.h>
#include <hsa/hsa_ext_amd.h>

Defines

HIP_CALL(cmd)
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
GetHwId(hwId)
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

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
enumerator MEM_MANAGED
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)
void RunParallelCopyBenchmark(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[8] = "CGBFUNM"
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 RunParallelCopyBenchmark(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/docs-6.1.2/src/include
dir /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-transferbench/checkouts/docs-6.1.2/src