From b8cf48fc16b4bf1d053e1208c443ca72b5ad02c3 Mon Sep 17 00:00:00 2001 From: gilbertlee-amd <44450918+gilbertlee-amd@users.noreply.github.com> Date: Wed, 7 Aug 2019 17:21:41 -0600 Subject: [PATCH] Adding TransferBench tool (#113) * Adding standalone TransferBench tool --- tools/TransferBench/Makefile | 16 ++ tools/TransferBench/TransferBench.cpp | 313 ++++++++++++++++++++++++++ tools/TransferBench/TransferBench.hpp | 111 +++++++++ tools/TransferBench/copy_kernel.h | 310 +++++++++++++++++++++++++ tools/TransferBench/example.cfg | 4 + 5 files changed, 754 insertions(+) create mode 100644 tools/TransferBench/Makefile create mode 100644 tools/TransferBench/TransferBench.cpp create mode 100644 tools/TransferBench/TransferBench.hpp create mode 100644 tools/TransferBench/copy_kernel.h create mode 100644 tools/TransferBench/example.cfg diff --git a/tools/TransferBench/Makefile b/tools/TransferBench/Makefile new file mode 100644 index 0000000000..d22ac47063 --- /dev/null +++ b/tools/TransferBench/Makefile @@ -0,0 +1,16 @@ +HIP_PATH?= $(wildcard /opt/rocm/hip) +ifeq (,$(HIP_PATH)) +HIP_PATH=../../.. +endif +HIPCC=$(HIP_PATH)/bin/hipcc + +EXE=TransferBench +CXXFLAGS = -O3 -fopenmp -I../../src/include -I. + +all: $(EXE) + +$(EXE): $(EXE).cpp $(shell find -regex ".*\.\hpp") + $(HIPCC) $(CXXFLAGS) $< -o $@ + +clean: + rm -f *.o $(EXE) diff --git a/tools/TransferBench/TransferBench.cpp b/tools/TransferBench/TransferBench.cpp new file mode 100644 index 0000000000..dca4f1376c --- /dev/null +++ b/tools/TransferBench/TransferBench.cpp @@ -0,0 +1,313 @@ +/* +Copyright (c) 2019 - present 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. +*/ + +// This program measures simultaneous copy performance across multiple GPUs +// on the same node + +#include +#include +#include +#include + +#include +#include "copy_kernel.h" +#include "TransferBench.hpp" + +int main(int argc, char **argv) +{ + // Display usage + if (argc <= 1) + { + printf("Usage: %s configFile \n", argv[0]); + printf("- configFile: file describing topologies to test\n"); + printf(" Each line should contain a single topology\n"); + printf(" L - number of links followed by L white-space separated triples (src, dst, # blocks)\n"); + printf(" For example:\n"); + printf(" 2 0 1 3 1 0 3\n"); + printf(" would define 2 links each using 3 threadblocks from GPU0 -> GPU1, and GPU1->GPU0\n"); + printf("- N: (Optional) Number of bytes to transfer per link.\n"); + printf(" If not specified, defaults to 2^28=256MB. Must be a multiple of 128 bytes\n"); + printf("Set env var USE_MEMCPY_ASYNC to use hipMemcpyAsync instead of copy kernel\n"); + exit(0); + } + + // Parse number of bytes to use (or use default if not specified) + size_t const numBytesPerLink = argc > 2 ? atoll(argv[2]) : (1<<28); + size_t N = numBytesPerLink / sizeof(float); + if (numBytesPerLink % 128) + { + printf("[ERROR] numBytesPerLink (%lu) must be a multiple of 128\n", numBytesPerLink); + exit(1); + } + + // Currently an environment variable is required in order to enable fine-grained VRAM allocations + if (!getenv("HSA_FORCE_FINE_GRAIN_PCIE")) + { + printf("[ERROR] Currently you must set HSA_FORCE_FINE_GRAIN_PCIE=1 prior to execution\n"); + exit(1); + } + + bool useMemcpy = getenv("USE_MEMCPY_ASYNC"); + printf("Using %s\n", useMemcpy ? "hipMemcpyAsync (USE_MEMCPY_ASYNC found) [# of blocks to use will be ignored]" : "copy kernel (USE_MEMCPY_ASYNC not found)"); + + // Collect the number of available GPUs on this machine + int numDevices; + HIP_CALL(hipGetDeviceCount(&numDevices)); + if (numDevices < 1) + { + printf("[ERROR] No GPU devices found\n"); + exit(1); + } + + // Print header + printf("%-*s(GB/s)", MAX_NAME_LEN - 6, "Configuration"); + for (int i = 0; i < numDevices; i++) + printf(" GPU %02d", i); + printf(" Total\n"); + for (int i = 0; i < MAX_NAME_LEN + 8 * (numDevices + 1); i++) printf("="); + printf("\n"); + + // Read configuration file + FILE* fp = fopen(argv[1], "r"); + if (!fp) + { + printf("[ERROR] Unable to open link configuration file: [%s]\n", argv[1]); + exit(1); + } + + // Track links that get used + std::map, int> linkMap; + + char line[2048]; + while(fgets(line, 2048, fp)) + { + // Parse links from configuration file + std::vector links; + ParseLinks(line, links); + + int const numLinks = links.size(); + if (numLinks == 0) continue; + + // Clear counters + int linkCount[numDevices]; + for (int i = 0; i < numDevices; i++) + linkCount[i] = 0; + + float* linkSrcMem[numLinks]; + float* linkDstMem[numLinks]; + hipStream_t streams[numLinks]; + hipEvent_t startEvents[numLinks]; + hipEvent_t stopEvents[numLinks]; + std::vector cpuBlockParams[numLinks]; + BlockParam* gpuBlockParams[numLinks]; + + char name[MAX_NAME_LEN+1] = {}; + + for (int i = 0; i < numLinks; i++) + { + int const src = links[i].srcGpu; + int const dst = links[i].dstGpu; + if (src < 0 || src >= numDevices || + dst < 0 || dst >= numDevices) + { + printf("[ERROR] Invalid link (%d to %d). Total devices: %d\n", src, dst, numDevices); + exit(1); + } + snprintf(name + strlen(name), MAX_NAME_LEN, "%d->%d:%d ", src, dst, links[i].numBlocksToUse); + + // Enable peer-to-peer access if this is the first time seeing this pair + auto linkPair = std::make_pair(src, dst); + linkMap[linkPair]++; + if (linkMap[linkPair] == 1) + { + int canAccess; + HIP_CALL(hipDeviceCanAccessPeer(&canAccess, src, dst)); + if (!canAccess) + { + printf("[ERROR] Unable to enable peer access between device %d and %d\n", src, dst); + exit(1); + } + HIP_CALL(hipSetDevice(src)); + HIP_CALL(hipDeviceEnablePeerAccess(dst, 0)); + } + + // Count # of links / total blocks each GPU will be working on + linkCount[src]++; + + // Allocate GPU memory on source GPU / streams / events + HIP_CALL(hipSetDevice(links[i].srcGpu)); + HIP_CALL(hipStreamCreate(&streams[i])); + HIP_CALL(hipEventCreate(&startEvents[i])); + HIP_CALL(hipEventCreate(&stopEvents[i])); + HIP_CALL(hipMalloc((void **)&linkSrcMem[i], numBytesPerLink)); + HIP_CALL(hipMalloc((void**)&gpuBlockParams[i], sizeof(BlockParam) * numLinks)); + CheckOrFill(N, linkSrcMem[i], false); + + // Allocate fine-grained GPU memory on destination GPU + HIP_CALL(hipSetDevice(links[i].dstGpu)); + HIP_CALL(hipExtMallocWithFlags((void**)&linkDstMem[i], numBytesPerLink, hipDeviceMallocFinegrained)); + + // 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 + // NOTE: Have each sub-array to work on multiple of 32-floats (128-bytes), + // but divide as evenly as possible + // NOTE: N is always a multiple of 32 + int blocksWithExtra = (N / 32) % links[i].numBlocksToUse; + int perBlockBaseN = (N / 32) / links[i].numBlocksToUse * 32; + for (int j = 0; j < links[i].numBlocksToUse; j++) + { + BlockParam param; + param.N = perBlockBaseN + ((j < blocksWithExtra) ? 32 : 0); + param.src = linkSrcMem[i] + ((j * perBlockBaseN) + ((j < blocksWithExtra) ? + j : blocksWithExtra) * 32); + param.dst = linkDstMem[i] + ((j * perBlockBaseN) + ((j < blocksWithExtra) ? + j : blocksWithExtra) * 32); + cpuBlockParams[i].push_back(param); + } + + HIP_CALL(hipMemcpy(gpuBlockParams[i], cpuBlockParams[i].data(), + sizeof(BlockParam) * links[i].numBlocksToUse, hipMemcpyHostToDevice)); + } + + // Launch kernels (warmup iterations are not counted) + int numWarmups = 3; + int numIterations = 10; + double totalCpuTime = 0; + double totalGpuTime[numDevices]; + for (int i = 0; i < numDevices; i++) totalGpuTime[i] = 0.0; + + for (int iteration = -numWarmups; iteration < numIterations; iteration++) + { + auto cpuStart = std::chrono::high_resolution_clock::now(); + #pragma omp parallel for num_threads(numLinks) + for (int i = 0; i < numLinks; i++) + { + HIP_CALL(hipSetDevice(links[i].srcGpu)); + HIP_CALL(hipEventRecord(startEvents[i], streams[i])); + if (useMemcpy) + { + HIP_CALL(hipMemcpyAsync(linkDstMem[i], linkSrcMem[i], + numBytesPerLink, hipMemcpyDeviceToDevice, + streams[i])); + } + else + { + hipLaunchKernelGGL(CopyKernel, + dim3(links[i].numBlocksToUse, 1, 1), + dim3(BLOCKSIZE, 1, 1), + 0, + streams[i], + gpuBlockParams[i]); + } + HIP_CALL(hipEventRecord(stopEvents[i], streams[i])); + } + + for (int i = 0; i < numLinks; i++) + hipStreamSynchronize(streams[i]); + + auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; + double deltaSec = std::chrono::duration_cast>(cpuDelta).count(); + + if (iteration >= 0) + { + totalCpuTime += deltaSec; + + for (int i = 0; i < numDevices; i++) + { + // Multiple links running on the same device may be running simultaneously + // so try to figure out the first/last event across all links + float maxTime = 0.0f; + for (int j = 0; j < numLinks; j++) + { + if (links[j].srcGpu != i) continue; + for (int k = 0; k < numLinks; k++) + { + if (links[k].srcGpu != i) continue; + + float gpuDeltaMsec; + HIP_CALL(hipEventElapsedTime(&gpuDeltaMsec, startEvents[j], stopEvents[k])); + maxTime = std::max(maxTime, gpuDeltaMsec); + } + } + totalGpuTime[i] += maxTime / 1000.0; + } + } + } + + // Validate that each link has transferred correctly + for (int i = 0; i < numLinks; i++) + CheckOrFill(N, linkDstMem[i], true); + + // Report timings + printf("%-*s", MAX_NAME_LEN, name); + for (int i = 0; i < numDevices; i++) + { + if (linkCount[i] == 0) + { + printf("%8.3f", 0.0); + } + else + { + totalGpuTime[i] /= (1.0 * numIterations); + printf("%8.3f", (linkCount[i] * numBytesPerLink / 1.0E9) / totalGpuTime[i]); + } + } + + // Print off bandwidth (based on CPU wall-time timer) + totalCpuTime /= numIterations; + printf("%8.3f\n", (numLinks * numBytesPerLink / 1.0E9) / totalCpuTime); + + // Release GPU memory + for (int i = 0; i < numLinks; i++) + { + HIP_CALL(hipFree(linkSrcMem[i])); + HIP_CALL(hipFree(linkDstMem[i])); + HIP_CALL(hipFree(gpuBlockParams[i])); + HIP_CALL(hipStreamDestroy(streams[i])); + HIP_CALL(hipEventDestroy(startEvents[i])); + HIP_CALL(hipEventDestroy(stopEvents[i])); + + } + } + fclose(fp); + + // Print link information + for (int i = 0; i < MAX_NAME_LEN + 8 * (numDevices + 1); i++) printf("="); + printf("\n"); + printf("Link topology:\n"); + uint32_t linkType; + uint32_t hopCount; + for (auto mapPair : linkMap) + { + int src = mapPair.first.first; + int dst = mapPair.first.second; + HIP_CALL(hipExtGetLinkTypeAndHopCount(src, dst, &linkType, &hopCount)); + printf("%d -> %d: %s [%d hop(s)]\n", src, dst, + linkType == HSA_AMD_LINK_INFO_TYPE_HYPERTRANSPORT ? "HYPERTRANSPORT" : + linkType == HSA_AMD_LINK_INFO_TYPE_QPI ? "QPI" : + linkType == HSA_AMD_LINK_INFO_TYPE_PCIE ? "PCIE" : + linkType == HSA_AMD_LINK_INFO_TYPE_INFINBAND ? "INFINIBAND" : + linkType == HSA_AMD_LINK_INFO_TYPE_XGMI ? "XGMI" : "UNKNOWN", + hopCount); + } + return 0; +} diff --git a/tools/TransferBench/TransferBench.hpp b/tools/TransferBench/TransferBench.hpp new file mode 100644 index 0000000000..25f61e9253 --- /dev/null +++ b/tools/TransferBench/TransferBench.hpp @@ -0,0 +1,111 @@ +/* +Copyright (c) 2019 - present 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. +*/ + +// Helper macro for catching HIP errors +#define HIP_CALL(cmd) \ + do { \ + hipError_t error = (cmd); \ + if (error != hipSuccess) \ + { \ + std::cerr << "Encountered HIP error (" << hipGetErrorString(error) << ") at line " \ + << __LINE__ << " in file " << __FILE__ << "\n"; \ + exit(-1); \ + } \ + } while (0) + +#define MAX_NAME_LEN 64 +#define BLOCKSIZE 256 +#define COPY_UNROLL 4 + +// Each link is defined between a source GPU and destination GPU +struct Link +{ + int srcGpu; // Source GPU (global memory source) + int dstGpu; // Destination GPU (fine-grained memory destination) + int numBlocksToUse; // Number of threadblocks to use for this link +}; + +// Each threadblock copies N floats from src to dst +struct BlockParam +{ + int N; + float* src; + float* dst; +}; + +// GPU copy kernel +__global__ void __launch_bounds__(BLOCKSIZE) +CopyKernel(BlockParam* blockParams) +{ + // Collect the arguments for this block + int N = blockParams[blockIdx.x].N; + const float* __restrict__ src = (float* )blockParams[blockIdx.x].src; + float* __restrict__ dst = (float* )blockParams[blockIdx.x].dst; + + Copy(dst, src, N); +} + +// Helper function to parse a link of link definitions +void ParseLinks(char const* line, std::vector& links) +{ + links.clear(); + int numLinks = 0; + + std::istringstream iss; + iss.clear(); + iss.str(line); + iss >> numLinks; + links.resize(numLinks); + if (iss.fail()) return; + + + for (int i = 0; i < numLinks; i++) + iss >> links[i].srcGpu >> links[i].dstGpu >> links[i].numBlocksToUse; +} + +// Helper function to either fill a device pointer with pseudo-random data, or to check to see if it matches +void CheckOrFill(int N, float* devPtr, bool doCheck) +{ + float* refBuffer = (float*)malloc(N * sizeof(float)); + + for (int i = 0; i < N; i++) + refBuffer[i] = i % 383 + 31; + + if (doCheck) + { + float* hostBuffer = (float*) malloc(N * sizeof(float)); + HIP_CALL(hipMemcpy(hostBuffer, devPtr, N * sizeof(float), hipMemcpyDeviceToHost)); + for (int i = 0; i < N; i++) + { + if (refBuffer[i] != hostBuffer[i]) + { + printf("[ERROR] Mismatch at element %d Ref: %f Actual: %f\n", i, refBuffer[i], hostBuffer[i]); + exit(1); + } + } + } + else + { + HIP_CALL(hipMemcpy(devPtr, refBuffer, N * sizeof(float), hipMemcpyHostToDevice)); + } + free(refBuffer); +} diff --git a/tools/TransferBench/copy_kernel.h b/tools/TransferBench/copy_kernel.h new file mode 100644 index 0000000000..e2377fdbb8 --- /dev/null +++ b/tools/TransferBench/copy_kernel.h @@ -0,0 +1,310 @@ +/************************************************************************* + * Copyright (c) 2015, NVIDIA CORPORATION. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + + +#ifndef COPY_KERNEL_H_ +#define COPY_KERNEL_H_ +#include +#include + +// Define min for ssize_t +static __device__ int min(int a, ssize_t b) { return (a < b) ? a : b; } + +typedef uint64_t PackType; + +template +struct MULTI { + __device__ PackType operator()(const PackType x, const PackType y) const + { + return FUNC()(x, y); + } +}; + +#define ALIGNUP(x, a) ((((x)-1) & ~((a)-1)) + (a)) + +template +__device__ inline volatile T* AlignUp(volatile T * ptr, size_t align) { + size_t ptrval = reinterpret_cast(ptr); + return reinterpret_cast(ALIGNUP(ptrval, align)); +} + +template inline __device__ +T vFetch(const volatile T* ptr) { + return *ptr; +} + +template inline __device__ +void vStore(volatile T* ptr, const T val) { + *ptr = val; +} + +template +__attribute__((noinline)) +__device__ inline void ReduceCopy( + const int tid, const int nthreads, + const volatile T * __restrict__ const src0, + const volatile T * __restrict__ const src1, + volatile T * __restrict__ const dest0, + volatile T * __restrict__ const dest1, const int N) { + for (int idx = tid; idx < N; idx += nthreads) { + T val = vFetch(src0+idx); + if (TWO_INPUTS) { + val = FUNC()(val, vFetch(src1+idx)); + } + vStore(dest0+idx, val); + if (TWO_OUTPUTS) { + vStore(dest1+idx, val); + } + } +} + +template +struct FuncPassA { + __device__ T operator()(const T x, const T y) const { + return x; + } +}; + +template +struct FuncSum { + __device__ T operator()(const T x, const T y) const { + return x + y; + } +}; + +template +struct MULTI { + static_assert(sizeof(PackType) == 2 * sizeof(float), + "PackType must be twice the size of float."); + union converter { + PackType storage; + struct { + float a, b; + }; + }; + + __device__ PackType operator()(const PackType x, const PackType y) const { + converter cx, cy, cr; + cx.storage = x; + cy.storage = y; + + cr.a = FUNC()(cx.a, cy.a); + cr.b = FUNC()(cx.b, cy.b); + + return cr.storage; + } +}; + + +typedef ulong2 Pack128; + +template +struct MULTI128 { + __device__ void operator()(Pack128& x, Pack128& y) { + x.x = MULTI()(x.x, y.x); + x.y = MULTI()(x.y, y.y); + } +}; + +inline __device__ void Fetch128(Pack128& v, const Pack128* p) { + v.x = p->x; + v.y = p->y; +} +inline __device__ void Store128(Pack128* p, Pack128& v) { + p->x = v.x; + p->y = v.y; +} + +template +__device__ void ReduceCopyMulti(const int tid, const int nthreads, + int nsrcs, const T* srcs[MAXSRCS], int ndsts, T* dsts[MAXDSTS], + const int offset, const int N) { + for (int idx = offset+tid; idx < offset+N; idx += nthreads) { + T val = vFetch(srcs[0]+idx); + #pragma unroll + for (int i=1; i +__device__ void ReduceCopy128bMulti( const int w, const int nw, const int t, + int nsrcs, const T* s[MAXSRCS], int ndsts, T* d[MAXDSTS], + const int elemOffset, const int Npack) { + const int inc = nw * UNROLL * WARP_SIZE; + int offset = w * UNROLL * WARP_SIZE + t; + + const Pack128* srcs[MAXSRCS]; + for (int i=0; i()(vals[u], vals2[u]); + } + #pragma unroll 1 + for (int i=MINSRCS; i()(vals[u], vals2[u]); + } + + // Store + for (int i = 0; i < MINDSTS; i++) { + for (int u = 0; u < UNROLL; ++u) Store128(dsts[i]+u*WARP_SIZE, vals[u]); + } + #pragma unroll 1 + for (int i=MINDSTS; i +__device__ int ptrAlign128(T* ptr) { return (uint64_t)ptr % alignof(Pack128); } + +// Try to limit consecutive load/stores to 8. +// Use UNROLL 8 when we have a single source and a single destination, 4 otherwise +#define AUTOUNROLL (UNROLL*(4/(MINDSTS+MINSRCS))) + +template +__device__ void ReduceOrCopyMulti(const int tid, const int nthreads, + int nsrcs, const T* srcs[MAXSRCS], int ndsts, T* dsts[MAXDSTS], + int N) { + int Nrem = N; + if (Nrem <= 0) return; + + int alignDiff = 0; + int align = ptrAlign128(srcs[0]); + #pragma unroll + for (int i=1; i(tid, nthreads, nsrcs, srcs, ndsts, dsts, 0, Npreamble); + Nrem -= Npreamble; + if (Nrem == 0) return; + } + int offset = Npreamble; + + // stage 2: fast path: use 128b loads/stores to do the bulk of the work, + // assuming the pointers we have are all 128-bit alignable. + int w = tid / WARP_SIZE; // Warp number + int nw = nthreads / WARP_SIZE; // Number of warps + int t = tid % WARP_SIZE; // Thread (inside the warp) + + const int packFactor = sizeof(Pack128) / sizeof(T); + + // stage 2a: main loop + int Npack2a = (Nrem / (packFactor * AUTOUNROLL * WARP_SIZE)) + * (AUTOUNROLL * WARP_SIZE); // round down + int Nelem2a = Npack2a * packFactor; + + ReduceCopy128bMulti(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Npack2a); + + Nrem -= Nelem2a; + if (Nrem == 0) return; + offset += Nelem2a; + + // stage 2b: slightly less optimized for section when we don't have full + // unrolling + + int Npack2b = Nrem / packFactor; + int Nelem2b = Npack2b * packFactor; + + ReduceCopy128bMulti(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Npack2b); + + Nrem -= Nelem2b; + if (Nrem == 0) return; + offset += Nelem2b; + + // stage 2c: tail + ReduceCopyMulti(tid, nthreads, nsrcs, srcs, ndsts, dsts, offset, Nrem); +} + +// Assumptions: +// - there is exactly 1 block +// - THREADS is the number of producer threads +// - this function is called by all producer threads +template +__device__ void Copy(volatile T * __restrict__ const dest, + const volatile T * __restrict__ const src, const int N) { + const T* srcs[2]; + T* dsts[2]; + srcs[0] = (const T*)src; + dsts[0] = (T*)dest; + ReduceOrCopyMulti, T, 1, 2, 1, 2>(threadIdx.x, THREADS, + 1, srcs, 1, dsts, N); +} + +template +__device__ void DoubleCopy(volatile T * __restrict__ const dest0, + volatile T * __restrict__ const dest1, + const volatile T * __restrict__ const src, const int N) { + const T* srcs[2]; + T* dsts[2]; + srcs[0] = (const T*)src; + dsts[0] = (T*)dest0; + dsts[1] = (T*)dest1; + ReduceOrCopyMulti, T, 1, 2, 1, 2>(threadIdx.x, THREADS, + 1, srcs, 2, dsts, N); +} + +template +__device__ void Reduce(volatile T * __restrict__ const dest, + const volatile T * __restrict__ const src0, + const volatile T * __restrict__ const src1, const int N) { + const T* srcs[2]; + T* dsts[2]; + srcs[0] = (const T*)src0; + srcs[1] = (const T*)src1; + dsts[0] = (T*)dest; + ReduceOrCopyMulti, T, 1, 2, 1, 2>(threadIdx.x, THREADS, + 2, srcs, 1, dsts, N); +} + +template +__device__ void ReduceCopy(volatile T * __restrict__ const dest0, + volatile T * __restrict__ const dest1, + const volatile T * __restrict__ const src0, + const volatile T * __restrict__ const src1, const int N) { + const T* srcs[2]; + T* dsts[2]; + srcs[0] = (const T*)src0; + srcs[1] = (const T*)src1; + dsts[0] = (T*)dest0; + dsts[1] = (T*)dest1; + ReduceOrCopyMulti, T, 1, 2, 1, 2>(threadIdx.x, THREADS, + 2, srcs, 2, dsts, N); +} +#endif // COPY_KERNEL_H_ diff --git a/tools/TransferBench/example.cfg b/tools/TransferBench/example.cfg new file mode 100644 index 0000000000..aad8d4c600 --- /dev/null +++ b/tools/TransferBench/example.cfg @@ -0,0 +1,4 @@ +# Each line consists of L (# of links) followed by L white-space-separated triples of (srcGpu, dstGpu, #blocks) + +# Single link between GPUs 0 and 1 +1 0 1 1