API

API#

struct AgentData#

#include <GetClosestNumaNode.hpp>

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

API

Contents

API#