[TransferBench] Updating for 2.8.3
This commit is contained in:
Gilbert Lee
@@ -0,0 +1,137 @@
|
||||
#ifndef ENVVARS_HPP
|
||||
#define ENVVARS_HPP
|
||||
|
||||
// This class manages environment variable that affect TransferBench
|
||||
class EnvVars
|
||||
{
|
||||
public:
|
||||
// Default configuration values
|
||||
int const DEFAULT_NUM_WARMUPS = 3;
|
||||
int const DEFAULT_NUM_ITERATIONS = 10;
|
||||
int const DEFAULT_SAMPLING_FACTOR = 1;
|
||||
int const DEFAULT_NUM_CPU_PER_LINK = 4;
|
||||
|
||||
// Environment variables
|
||||
int useHipCall; // Use hipMemcpy/hipMemset instead of custom shader kernels
|
||||
int useMemset; // Perform a memset instead of a copy (ignores source memory)
|
||||
int useSingleSync; // Perform synchronization only once after all iterations instead of per iteration
|
||||
int useInteractive; // Pause for user-input before starting transfer loop
|
||||
int useSleep; // Adds a 100ms sleep after each synchronization
|
||||
int combineTiming; // Combines the timing with kernel launch
|
||||
int showAddr; // Print out memory addresses for each Link
|
||||
int outputToCsv; // Output in CSV format
|
||||
int byteOffset; // Byte-offset for memory allocations
|
||||
int numWarmups; // Number of un-timed warmup iterations to perform
|
||||
int numIterations; // Number of timed iterations to perform
|
||||
int samplingFactor; // Affects how many different values of N are generated (when N set to 0)
|
||||
int numCpuPerLink; // Number of CPU child threads to use per CPU link
|
||||
|
||||
// Constructor that collects values
|
||||
EnvVars()
|
||||
{
|
||||
useHipCall = GetEnvVar("USE_HIP_CALL" , 0);
|
||||
useMemset = GetEnvVar("USE_MEMSET" , 0);
|
||||
useSingleSync = GetEnvVar("USE_SINGLE_SYNC" , 0);
|
||||
useInteractive = GetEnvVar("USE_INTERACTIVE" , 0);
|
||||
combineTiming = GetEnvVar("COMBINE_TIMING" , 0);
|
||||
showAddr = GetEnvVar("SHOW_ADDR" , 0);
|
||||
outputToCsv = GetEnvVar("OUTPUT_TO_CSV" , 0);
|
||||
byteOffset = GetEnvVar("BYTE_OFFSET" , 0);
|
||||
numWarmups = GetEnvVar("NUM_WARMUPS" , DEFAULT_NUM_WARMUPS);
|
||||
numIterations = GetEnvVar("NUM_ITERATIONS" , DEFAULT_NUM_ITERATIONS);
|
||||
samplingFactor = GetEnvVar("SAMPLING_FACTOR" , DEFAULT_SAMPLING_FACTOR);
|
||||
numCpuPerLink = GetEnvVar("NUM_CPU_PER_LINK" , DEFAULT_NUM_CPU_PER_LINK);
|
||||
|
||||
// Perform some basic validation
|
||||
if (byteOffset % sizeof(float))
|
||||
{
|
||||
printf("[ERROR] BYTE_OFFSET must be set to multiple of %lu\n", sizeof(float));
|
||||
exit(1);
|
||||
}
|
||||
if (numWarmups < 0)
|
||||
{
|
||||
printf("[ERROR] NUM_WARMUPS must be set to a non-negative number\n");
|
||||
exit(1);
|
||||
}
|
||||
if (numIterations <= 0)
|
||||
{
|
||||
printf("[ERROR] NUM_ITERATIONS must be set to a positive number\n");
|
||||
exit(1);
|
||||
}
|
||||
if (samplingFactor < 1)
|
||||
{
|
||||
printf("[ERROR] SAMPLING_FACTOR must be greater or equal to 1\n");
|
||||
exit(1);
|
||||
}
|
||||
if (numCpuPerLink < 1)
|
||||
{
|
||||
printf("[ERROR] NUM_CPU_PER_LINK must be greater or equal to 1\n");
|
||||
exit(1);
|
||||
}
|
||||
}
|
||||
|
||||
// Display info on the env vars that can be used
|
||||
static void DisplayUsage()
|
||||
{
|
||||
printf("Environment variables:\n");
|
||||
printf("======================\n");
|
||||
printf(" USE_HIP_CALL - Use hipMemcpy/hipMemset instead of custom shader kernels for GPU-executed copies\n");
|
||||
printf(" USE_MEMSET - Perform a memset instead of a copy (ignores source memory)\n");
|
||||
printf(" USE_SINGLE_SYNC - Perform synchronization only once after all iterations instead of per iteration\n");
|
||||
printf(" USE_INTERACTIVE - Pause for user-input before starting transfer loop\n");
|
||||
printf(" COMBINE_TIMING - Combines timing with launch (potentially lower timing overhead)\n");
|
||||
printf(" SHOW_ADDR - Print out memory addresses for each Link\n");
|
||||
printf(" OUTPUT_TO_CSV - Outputs to CSV format if set\n");
|
||||
printf(" BYTE_OFFSET - Initial byte-offset for memory allocations. Must be multiple of 4. Defaults to 0\n");
|
||||
printf(" NUM_WARMUPS=W - Perform W untimed warmup iteration(s) per test\n");
|
||||
printf(" NUM_ITERATIONS=I - Perform I timed iteration(s) per test\n");
|
||||
printf(" SAMPLING_FACTOR=F - Add F samples (when possible) between powers of 2 when auto-generating data sizes\n");
|
||||
printf(" NUM_CPU_PER_LINK=C - Use C threads per Link for CPU-executed copies\n");
|
||||
}
|
||||
|
||||
// Display env var settings
|
||||
void DisplayEnvVars() const
|
||||
{
|
||||
if (!outputToCsv)
|
||||
{
|
||||
printf("Run configuration\n");
|
||||
printf("=====================================================\n");
|
||||
printf("%-20s = %12d : Using %s for GPU-executed copies\n", "USE_HIP_CALL", useHipCall,
|
||||
useHipCall ? "HIP functions" : "custom kernels");
|
||||
printf("%-20s = %12d : Performing %s\n", "USE_MEMSET", useMemset,
|
||||
useMemset ? "memset" : "memcopy");
|
||||
if (useHipCall && !useMemset)
|
||||
{
|
||||
char* env = getenv("HSA_ENABLE_SDMA");
|
||||
printf("%-20s = %12s : %s\n", "HSA_ENABLE_SDMA", env,
|
||||
(env && !strcmp(env, "0")) ? "Using blit kernels for hipMemcpy" : "Using DMA copy engines");
|
||||
}
|
||||
printf("%-20s = %12d : %s\n", "USE_SINGLE_SYNC", useSingleSync,
|
||||
useSingleSync ? "Synchronizing only once, after all iterations" : "Synchronizing per iteration");
|
||||
printf("%-20s = %12d : Running in %s mode\n", "USE_INTERACTIVE", useInteractive,
|
||||
useInteractive ? "interactive" : "non-interactive");
|
||||
printf("%-20s = %12d : %s\n", "COMBINE_TIMING", combineTiming,
|
||||
combineTiming ? "Using combined timing+launch" : "Using separate timing / launch");
|
||||
printf("%-20s = %12d : %s\n", "SHOW_ADDR", showAddr,
|
||||
showAddr ? "Displaying src/dst mem addresses" : "Not displaying src/dst mem addresses");
|
||||
printf("%-20s = %12d : Output to %s\n", "OUTPUT_TO_CSV", outputToCsv,
|
||||
outputToCsv ? "CSV" : "console");
|
||||
printf("%-20s = %12d : Using byte offset of %d\n", "BYTE_OFFSET", byteOffset, byteOffset);
|
||||
printf("%-20s = %12d : Running %d warmup iteration(s) per topology\n", "NUM_WARMUPS", numWarmups, numWarmups);
|
||||
printf("%-20s = %12d : Running %d timed iteration(s) per topology\n", "NUM_ITERATIONS", numIterations, numIterations);
|
||||
printf("%-20s = %12d : Using %d CPU thread(s) per CPU-based-copy Link\n", "NUM_CPU_PER_LINK", numCpuPerLink, numCpuPerLink);
|
||||
printf("\n");
|
||||
}
|
||||
};
|
||||
|
||||
private:
|
||||
// Helper function that gets parses environment variable or sets to default value
|
||||
int GetEnvVar(std::string const varname, int defaultValue)
|
||||
{
|
||||
if (getenv(varname.c_str()))
|
||||
return atoi(getenv(varname.c_str()));
|
||||
return defaultValue;
|
||||
}
|
||||
};
|
||||
|
||||
#endif
|
||||
@@ -6,7 +6,7 @@ endif
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
|
||||
EXE=TransferBench
|
||||
CXXFLAGS = -O3 -I../../src/include -I.
|
||||
CXXFLAGS = -O3 -I../../src/include -I. -lnuma
|
||||
|
||||
all: $(EXE)
|
||||
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -25,6 +25,7 @@ THE SOFTWARE.
|
||||
#include <chrono>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <cstdint>
|
||||
#include <set>
|
||||
#include <unistd.h>
|
||||
#include <map>
|
||||
@@ -33,7 +34,22 @@ THE SOFTWARE.
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_ext.h>
|
||||
#include <hsa/hsa_ext_amd.h>
|
||||
#include "copy_kernel.h"
|
||||
#include <hip/hcc_detail/hip_fp16.h>
|
||||
|
||||
// Include common_kernel.h from RCCL for copy kernel
|
||||
// However define some variables to avoid extra includes / missing defines
|
||||
#define NCCL_DEVICE_H_ // Avoid loading devcomm.h
|
||||
#define WARP_SIZE 64
|
||||
|
||||
typedef uint64_t PackType;
|
||||
typedef ulong2 Pack128;
|
||||
typedef struct
|
||||
{
|
||||
uint16_t data;
|
||||
} rccl_bfloat16;
|
||||
|
||||
#include "../../src/collectives/device/common_kernel.h"
|
||||
#include "EnvVars.hpp"
|
||||
|
||||
// Helper macro for catching HIP errors
|
||||
#define HIP_CALL(cmd) \
|
||||
@@ -50,11 +66,12 @@ THE SOFTWARE.
|
||||
// Different src/dst memory types supported
|
||||
typedef enum
|
||||
{
|
||||
MEM_CPU = 0, // Pinned CPU memory
|
||||
MEM_GPU = 1 // Global GPU memory
|
||||
MEM_CPU = 0, // Pinned CPU memory
|
||||
MEM_GPU = 1, // Coarse-grained global GPU memory
|
||||
MEM_GPU_FINE = 2 // Fine-grained global GPU memory
|
||||
} MemType;
|
||||
|
||||
char const MemTypeStr[3] = "CG";
|
||||
char const MemTypeStr[4] = "CGF";
|
||||
|
||||
typedef enum
|
||||
{
|
||||
@@ -62,17 +79,6 @@ typedef enum
|
||||
MODE_CHECK = 1 // Check data against pattern
|
||||
} ModeType;
|
||||
|
||||
// Each Link is a uni-direction operation from a src memory to dst memory executed by a specific GPU
|
||||
struct Link
|
||||
{
|
||||
int exeIndex; // GPU to execute on
|
||||
MemType srcMemType; // Source memory type
|
||||
int srcIndex; // Source device index
|
||||
MemType dstMemType; // Destination memory type
|
||||
int dstIndex; // Destination device index
|
||||
int numBlocksToUse; // Number of threadblocks to use for this Link
|
||||
};
|
||||
|
||||
// Each threadblock copies N floats from src to dst
|
||||
struct BlockParam
|
||||
{
|
||||
@@ -81,46 +87,98 @@ struct BlockParam
|
||||
float* dst;
|
||||
};
|
||||
|
||||
// Each Link is a uni-direction operation from a src memory to dst memory executed by a specific GPU
|
||||
struct Link
|
||||
{
|
||||
// Link config
|
||||
MemType exeMemType; // Link executor type (CPU or GPU)
|
||||
int exeIndex; // Executor index (NUMA node for CPU / device ID for GPU)
|
||||
MemType srcMemType; // Source memory type
|
||||
int srcIndex; // Source device index
|
||||
MemType dstMemType; // Destination memory type
|
||||
int dstIndex; // Destination device index
|
||||
int numBlocksToUse; // Number of threadblocks to use for this Link
|
||||
|
||||
// Link implementation
|
||||
float* srcMem; // Source memory
|
||||
float* dstMem; // Destination memory
|
||||
|
||||
hipEvent_t startEvent;
|
||||
hipEvent_t stopEvent;
|
||||
hipStream_t stream;
|
||||
BlockParam* blockParam;
|
||||
|
||||
double totalTime;
|
||||
};
|
||||
|
||||
void DisplayUsage(char const* cmdName); // Display usage instructions
|
||||
void GenerateConfigFile(char const* cfgFile, int numBlocks); // Generate a sample config file
|
||||
void DisplayTopology(); // Display GPU topology
|
||||
void ParseLinks(char* line, std::vector<Link>& links); // Parse Link information
|
||||
void AllocateMemory(MemType memType, int devIndex, size_t numBytes, bool useFineGrainMem, float** memPtr);
|
||||
void PopulateTestSizes(size_t const numBytesPerLink, int const samplingFactor, std::vector<size_t>& valuesofN);
|
||||
void ParseMemType(std::string const& token, int const numCpus, int const numGpus, MemType* memType, int* memIndex);
|
||||
void ParseLinks(char* line, int numCpus, int numGpus, std::vector<Link>& links); // Parse Link information
|
||||
void EnablePeerAccess(int const deviceId, int const peerDeviceId);
|
||||
void AllocateMemory(MemType memType, int devIndex, size_t numBytes, float** memPtr);
|
||||
void DeallocateMemory(MemType memType, int devIndex, float* memPtr);
|
||||
void CheckPages(char* byteArray, size_t numBytes, int targetId);
|
||||
void CheckOrFill(ModeType mode, int N, bool isMemset, bool isHipCall, float* ptr);
|
||||
void RunLink(EnvVars const& ev, size_t const N, int const iteration, Link& link);
|
||||
|
||||
|
||||
std::string GetLinkTypeDesc(uint32_t linkType, uint32_t hopCount);
|
||||
std::string GetDesc(MemType srcMemType, int srcIndex,
|
||||
MemType dstMemType, int dstIndex);
|
||||
std::string GetLinkDesc(Link const& link);
|
||||
|
||||
|
||||
#define MAX_NAME_LEN 64
|
||||
#define BLOCKSIZE 256
|
||||
#define COPY_UNROLL 4
|
||||
#define MEMSET_UNROLL 4
|
||||
|
||||
// Dummy reduction function (not used because it's just a copy)
|
||||
struct FuncNull {
|
||||
__device__ float operator()(const float x, const float y) const {
|
||||
return 0;
|
||||
}
|
||||
};
|
||||
|
||||
// GPU copy kernel
|
||||
__global__ void __launch_bounds__(BLOCKSIZE)
|
||||
CopyKernel(BlockParam* blockParams)
|
||||
GpuCopyKernel(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;
|
||||
// Collect the arguments for this block
|
||||
int N = blockParams[blockIdx.x].N;
|
||||
const float* src[1] = {(float* )blockParams[blockIdx.x].src};
|
||||
float* dst[1] = {(float* )blockParams[blockIdx.x].dst};
|
||||
|
||||
Copy<COPY_UNROLL, BLOCKSIZE>(dst, src, N);
|
||||
ReduceOrCopyMulti<COPY_UNROLL, FuncNull, float, 1, 1, 1, 1>(
|
||||
threadIdx.x, BLOCKSIZE, 1, src, 1, dst, N);
|
||||
}
|
||||
|
||||
// GPU set kernel
|
||||
__global__ void __launch_bounds__(BLOCKSIZE)
|
||||
MemsetKernel(BlockParam* blockParams)
|
||||
GpuMemsetKernel(BlockParam* blockParams)
|
||||
{
|
||||
// Collect the arguments for this block
|
||||
int N = blockParams[blockIdx.x].N;
|
||||
float* __restrict__ dst = (float*)blockParams[blockIdx.x].dst;
|
||||
// Collect the arguments for this block
|
||||
int N = blockParams[blockIdx.x].N;
|
||||
float* __restrict__ dst = (float*)blockParams[blockIdx.x].dst;
|
||||
|
||||
// Use non-zero value
|
||||
#pragma unroll MEMSET_UNROLL
|
||||
for (int tid = threadIdx.x; tid < N; tid += BLOCKSIZE)
|
||||
{
|
||||
dst[tid] = 1234.0;
|
||||
}
|
||||
// Use non-zero value
|
||||
#pragma unroll MEMSET_UNROLL
|
||||
for (int tid = threadIdx.x; tid < N; tid += BLOCKSIZE)
|
||||
{
|
||||
dst[tid] = 1234.0;
|
||||
}
|
||||
}
|
||||
|
||||
// CPU copy kernel
|
||||
void CpuCopyKernel(BlockParam const& blockParams)
|
||||
{
|
||||
memcpy(blockParams.dst, blockParams.src, blockParams.N * sizeof(float));
|
||||
}
|
||||
|
||||
// CPU memset kernel
|
||||
void CpuMemsetKernel(BlockParam const& blockParams)
|
||||
{
|
||||
for (int i = 0; i < blockParams.N; i++)
|
||||
blockParams.dst[i] = 1234.0;
|
||||
}
|
||||
|
||||
@@ -1,310 +0,0 @@
|
||||
/*************************************************************************
|
||||
* 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_
|
||||
@@ -1,39 +1,43 @@
|
||||
#Configfile Format:
|
||||
#==================
|
||||
#A Link is defined as a uni-directional transfer from src memory location to dst memory location
|
||||
#Each single line in the configuration file defines a set of Links to run in parallel
|
||||
# Configfile Format:
|
||||
# ==================
|
||||
# A Link is defined as a uni-directional transfer from src memory location to dst memory location executed by either CPU or GPU
|
||||
# Each single line in the configuration file defines a set of Links to run in parallel
|
||||
|
||||
#There are two ways to specify the configuration file:
|
||||
# There are two ways to specify the configuration file:
|
||||
|
||||
#1) Basic
|
||||
# The basic specification assumes the same number of threadblocks/CUs used per link
|
||||
# A positive number of Links is specified followed by that number of triplets describing each Link
|
||||
# 1) Basic
|
||||
# The basic specification assumes the same number of threadblocks/CUs used per GPU-executed Link
|
||||
# A positive number of Links is specified followed by that number of triplets describing each Link
|
||||
|
||||
#Links #CUs (GPUIndex1 srcMem1 dstMem1) ... (GPUIndexL srcMemL dstMemL)
|
||||
# #Links #CUs (srcMem1->Executor1->dstMem1) ... (srcMemL->ExecutorL->dstMemL)
|
||||
|
||||
#2) Advanced
|
||||
# The advanced specification allows different number of threadblocks/CUs used per Link
|
||||
# A negative number of links is specified, followed by quadruples describing each Link
|
||||
# -#Links (GPUIndex1 #CUs1 srcMem1 dstMem1) ... (GPUIndexL #CUsL srcMemL dstMemL)
|
||||
# 2) Advanced
|
||||
# The advanced specification allows different number of threadblocks/CUs used per GPU-executed Link
|
||||
# A negative number of links is specified, followed by quadruples describing each Link
|
||||
# -#Links (srcMem1->Executor1->dstMem1 #CUs1) ... (srcMemL->ExecutorL->dstMemL #CUsL)
|
||||
|
||||
#Argument Details:
|
||||
# #Links : Number of Links to be run in parallel
|
||||
# #CUs : Number of threadblocks/CUs to use for a Link
|
||||
# GpuIndex: 0-indexed GPU id executing the Link
|
||||
# srcMemL : Source memory location (Where the data is to be read from). Ignored in memset mode
|
||||
# dstMemL : Destination memory location (Where the data is to be written to)
|
||||
# Memory locations are specified by a character indicating memory type, followed by GPU device index (0-indexed)
|
||||
# Supported memory locations are:
|
||||
# - P: Pinned host memory (on CPU, on NUMA node closest to provided GPU index)
|
||||
# - G: Global device memory (on GPU)
|
||||
#Round brackets may be included for human clarity, but will be ignored
|
||||
# Argument Details:
|
||||
# #Links : Number of Links to be run in parallel
|
||||
# #CUs : Number of threadblocks/CUs to use for a GPU-executed Link
|
||||
# srcMemL : Source memory location (Where the data is to be read from). Ignored in memset mode
|
||||
# Executor: Executor are specified by a character indicating executor type, followed by device index (0-indexed)
|
||||
# - C: CPU-executed (Indexed from 0 to 1)
|
||||
# - G: GPU-executed (Indexed from 0 to 3)
|
||||
# dstMemL : Destination memory location (Where the data is to be written to)
|
||||
|
||||
#Examples:
|
||||
#1 4 (0 G0 G1) Single Link that uses 4 CUs on GPU 0 that reads memory from GPU 0 and copies it to memory on GPU 1
|
||||
#1 4 (0 G1 G0) Single Link that uses 4 CUs on GPU 0 that reads memory from GPU 1 and copies it to memory on GPU 0
|
||||
#1 4 (2 P0 G2) Single Link that uses 4 CUs on GPU 2 that reads memory from CPU 0 and copies it to memory on GPU 2
|
||||
#2 4 (0 G0 G1) (1 G1 G0) Runs 2 Links in parallel. GPU 0 - > GPU1, and GP1 -> GPU 0, each with 4 CUs
|
||||
#-2 (0 G0 G1 4) (1 G1 G0 2) Runs 2 Links in parallel. GPU 0 - > GPU 1 using four CUs, and GPU1 -> GPU 0 using two CUs
|
||||
# Memory locations are specified by a character indicating memory type, followed by device index (0-indexed)
|
||||
# Supported memory locations are:
|
||||
# - C: Pinned host memory (on NUMA node, indexed from 0 to 1)
|
||||
# - G: Global device memory (on GPU device indexed from 0 to 3)
|
||||
# - F: Fine-grain device memory (on GPU device indexed from 0 to 3)
|
||||
|
||||
# Single link between GPUs 0 and 1
|
||||
1 1 (0 G0 G1)
|
||||
# Examples:
|
||||
# 1 4 (G0->G0->G1) Single Link that uses 4 CUs on GPU 0 that reads memory from GPU 0 and copies it to memory on GPU 1
|
||||
# 1 4 (G1->C0->G0) Single Link that uses 4 CUs on GPU 0 that reads memory from CPU 1 and copies it to memory on GPU 0
|
||||
# 1 4 (C0->G2->G2) Single Link that uses 4 CUs on GPU 2 that reads memory from CPU 0 and copies it to memory on GPU 2
|
||||
# 2 4 G0->G0->G1 G1->G1->G0 Runs 2 Links in parallel. GPU 0 - > GPU1, and GP1 -> GPU 0, each with 4 CUs
|
||||
# -2 (G0 G0 G1 4) (G1 G1 G0 2) Runs 2 Links in parallel. GPU 0 - > GPU 1 using four CUs, and GPU1 -> GPU 0 using two CUs
|
||||
# Round brackets and arrows' ->' may be included for human clarity, but will be ignored and are unnecessary
|
||||
|
||||
# Single GPU-executed link between GPUs 0 and 1 using 4 CUs
|
||||
1 4 (G0->G0->G1)
|
||||
|
||||
Reference in New Issue
Block a user