Adding TransferBench tool (#113)
* Adding standalone TransferBench tool
This commit is contained in:
gilbertlee-amd
GitHub
@@ -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)
|
||||
@@ -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 <chrono>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <set>
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "copy_kernel.h"
|
||||
#include "TransferBench.hpp"
|
||||
|
||||
int main(int argc, char **argv)
|
||||
{
|
||||
// Display usage
|
||||
if (argc <= 1)
|
||||
{
|
||||
printf("Usage: %s configFile <N>\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<std::pair<int, int>, int> linkMap;
|
||||
|
||||
char line[2048];
|
||||
while(fgets(line, 2048, fp))
|
||||
{
|
||||
// Parse links from configuration file
|
||||
std::vector<Link> 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<BlockParam> 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<std::chrono::duration<double>>(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;
|
||||
}
|
||||
@@ -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<COPY_UNROLL, BLOCKSIZE>(dst, src, N);
|
||||
}
|
||||
|
||||
// Helper function to parse a link of link definitions
|
||||
void ParseLinks(char const* line, std::vector<Link>& 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);
|
||||
}
|
||||
@@ -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 <cstdio>
|
||||
#include <cstdint>
|
||||
|
||||
// Define min for ssize_t
|
||||
static __device__ int min(int a, ssize_t b) { return (a < b) ? a : b; }
|
||||
|
||||
typedef uint64_t PackType;
|
||||
|
||||
template<class FUNC, typename T>
|
||||
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<typename T>
|
||||
__device__ inline volatile T* AlignUp(volatile T * ptr, size_t align) {
|
||||
size_t ptrval = reinterpret_cast<size_t>(ptr);
|
||||
return reinterpret_cast<volatile T*>(ALIGNUP(ptrval, align));
|
||||
}
|
||||
|
||||
template<typename T> inline __device__
|
||||
T vFetch(const volatile T* ptr) {
|
||||
return *ptr;
|
||||
}
|
||||
|
||||
template<typename T> inline __device__
|
||||
void vStore(volatile T* ptr, const T val) {
|
||||
*ptr = val;
|
||||
}
|
||||
|
||||
template<class FUNC, typename T, bool TWO_INPUTS, bool TWO_OUTPUTS>
|
||||
__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<typename T>
|
||||
struct FuncPassA {
|
||||
__device__ T operator()(const T x, const T y) const {
|
||||
return x;
|
||||
}
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
struct FuncSum {
|
||||
__device__ T operator()(const T x, const T y) const {
|
||||
return x + y;
|
||||
}
|
||||
};
|
||||
|
||||
template<class FUNC>
|
||||
struct MULTI<FUNC, float> {
|
||||
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<class FUNC, typename T>
|
||||
struct MULTI128 {
|
||||
__device__ void operator()(Pack128& x, Pack128& y) {
|
||||
x.x = MULTI<FUNC, T>()(x.x, y.x);
|
||||
x.y = MULTI<FUNC, T>()(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<class FUNC, typename T, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
|
||||
__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<MINSRCS; i++) val = FUNC()(val, vFetch(srcs[i]+idx));
|
||||
#pragma unroll 1
|
||||
for (int i=MINSRCS; i<MAXSRCS && i<nsrcs; i++) val = FUNC()(val, vFetch(srcs[i]+idx));
|
||||
|
||||
#pragma unroll
|
||||
for (int i=0; i<MINDSTS; i++) vStore(dsts[i]+idx, val);
|
||||
#pragma unroll 1
|
||||
for (int i=MINDSTS; i<MAXDSTS && i<ndsts; i++) vStore(dsts[i]+idx, val);
|
||||
}
|
||||
}
|
||||
|
||||
#define WARP_SIZE 64
|
||||
|
||||
template<class FUNC, typename T, int UNROLL, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
|
||||
__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<MAXSRCS; i++) srcs[i] = ((const Pack128*)(s[i]+elemOffset))+offset;
|
||||
Pack128* dsts[MAXDSTS];
|
||||
for (int i=0; i<MAXDSTS; i++) dsts[i] = ((Pack128*)(d[i]+elemOffset))+offset;
|
||||
|
||||
while (offset < Npack) {
|
||||
Pack128 vals[UNROLL];
|
||||
// Load and reduce
|
||||
for (int u = 0; u < UNROLL; ++u) Fetch128(vals[u], srcs[0]+u*WARP_SIZE);
|
||||
|
||||
for (int i=1; i<MINSRCS; i++) {
|
||||
Pack128 vals2[UNROLL];
|
||||
for (int u = 0; u < UNROLL; ++u) Fetch128(vals2[u], srcs[i]+u*WARP_SIZE);
|
||||
for (int u = 0; u < UNROLL; ++u) MULTI128<FUNC, T>()(vals[u], vals2[u]);
|
||||
}
|
||||
#pragma unroll 1
|
||||
for (int i=MINSRCS; i<MAXSRCS && i<nsrcs; i++) {
|
||||
Pack128 vals2[UNROLL];
|
||||
for (int u = 0; u < UNROLL; ++u) Fetch128(vals2[u], srcs[i]+u*WARP_SIZE);
|
||||
for (int u = 0; u < UNROLL; ++u) MULTI128<FUNC, T>()(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<MAXDSTS && i<ndsts; i++) {
|
||||
for (int u = 0; u < UNROLL; ++u) Store128(dsts[i]+u*WARP_SIZE, vals[u]);
|
||||
}
|
||||
for (int i=0; i<MAXSRCS; i++) srcs[i] += inc;
|
||||
for (int i=0; i<MAXDSTS; i++) dsts[i] += inc;
|
||||
offset += inc;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__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<int UNROLL, class FUNC, typename T, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
|
||||
__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<MINSRCS; i++) alignDiff |= (align ^ ptrAlign128(srcs[i]));
|
||||
for (int i=MINSRCS; i<MAXSRCS && i<nsrcs; i++) alignDiff |= (align ^ ptrAlign128(srcs[i]));
|
||||
#pragma unroll
|
||||
for (int i=0; i<MINDSTS; i++) alignDiff |= (align ^ ptrAlign128(dsts[i]));
|
||||
for (int i=MINDSTS; i<MAXDSTS && i<ndsts; i++) alignDiff |= (align ^ ptrAlign128(dsts[i]));
|
||||
|
||||
int Npreamble = alignDiff ? Nrem :
|
||||
N < alignof(Pack128) ? N :
|
||||
(alignof(Pack128) - align) % alignof(Pack128);
|
||||
|
||||
// stage 1: preamble: handle any elements up to the point of everything coming
|
||||
// into alignment
|
||||
if (Npreamble) {
|
||||
ReduceCopyMulti<FUNC, T, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(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<FUNC, T, AUTOUNROLL, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(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<FUNC, T, 1, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Npack2b);
|
||||
|
||||
Nrem -= Nelem2b;
|
||||
if (Nrem == 0) return;
|
||||
offset += Nelem2b;
|
||||
|
||||
// stage 2c: tail
|
||||
ReduceCopyMulti<FUNC, T, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(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<int UNROLL, int THREADS, typename T>
|
||||
__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<UNROLL, FuncPassA<T>, T, 1, 2, 1, 2>(threadIdx.x, THREADS,
|
||||
1, srcs, 1, dsts, N);
|
||||
}
|
||||
|
||||
template<int UNROLL, int THREADS, typename T>
|
||||
__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<UNROLL, FuncPassA<T>, T, 1, 2, 1, 2>(threadIdx.x, THREADS,
|
||||
1, srcs, 2, dsts, N);
|
||||
}
|
||||
|
||||
template<int UNROLL, int THREADS, typename T>
|
||||
__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<UNROLL, FuncPassA<T>, T, 1, 2, 1, 2>(threadIdx.x, THREADS,
|
||||
2, srcs, 1, dsts, N);
|
||||
}
|
||||
|
||||
template<int UNROLL, int THREADS, typename T>
|
||||
__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<UNROLL, FuncPassA<T>, T, 1, 2, 1, 2>(threadIdx.x, THREADS,
|
||||
2, srcs, 2, dsts, N);
|
||||
}
|
||||
#endif // COPY_KERNEL_H_
|
||||
@@ -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
|
||||
Viittaa uudesa ongelmassa
Block a user