TransferBench.cpp 78.5 KB
Newer Older
Gilbert Lee's avatar
Gilbert Lee committed
1
/*
gilbertlee-amd's avatar
gilbertlee-amd committed
2
Copyright (c) 2019-2023 Advanced Micro Devices, Inc. All rights reserved.
Gilbert Lee's avatar
Gilbert Lee committed
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24

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
25
26
#include <numa.h>     // If not found, try installing libnuma-dev (e.g apt-get install libnuma-dev)
#include <cmath>      // If not found, try installing g++-12      (e.g apt-get install g++-12)
Gilbert Lee's avatar
Gilbert Lee committed
27
#include <numaif.h>
Gilbert Lee's avatar
Gilbert Lee committed
28
#include <random>
Gilbert Lee's avatar
Gilbert Lee committed
29
30
31
32
33
34
35
36
#include <stack>
#include <thread>

#include "TransferBench.hpp"
#include "GetClosestNumaNode.hpp"

int main(int argc, char **argv)
{
Gilbert Lee's avatar
Gilbert Lee committed
37
38
39
40
41
42
43
  // Check for NUMA library support
  if (numa_available() == -1)
  {
    printf("[ERROR] NUMA library not supported. Check to see if libnuma has been installed on this system\n");
    exit(1);
  }

Gilbert Lee's avatar
Gilbert Lee committed
44
45
46
47
48
49
50
51
52
53
54
55
  // Display usage instructions and detected topology
  if (argc <= 1)
  {
    int const outputToCsv = EnvVars::GetEnvVar("OUTPUT_TO_CSV", 0);
    if (!outputToCsv) DisplayUsage(argv[0]);
    DisplayTopology(outputToCsv);
    exit(0);
  }

  // Collect environment variables / display current run configuration
  EnvVars ev;

Gilbert Lee's avatar
Gilbert Lee committed
56
57
  // Determine number of bytes to run per Transfer
  size_t numBytesPerTransfer = argc > 2 ? atoll(argv[2]) : DEFAULT_BYTES_PER_TRANSFER;
Gilbert Lee's avatar
Gilbert Lee committed
58
59
60
61
62
63
  if (argc > 2)
  {
    // Adjust bytes if unit specified
    char units = argv[2][strlen(argv[2])-1];
    switch (units)
    {
Gilbert Lee's avatar
Gilbert Lee committed
64
65
66
    case 'K': case 'k': numBytesPerTransfer *= 1024; break;
    case 'M': case 'm': numBytesPerTransfer *= 1024*1024; break;
    case 'G': case 'g': numBytesPerTransfer *= 1024*1024*1024; break;
Gilbert Lee's avatar
Gilbert Lee committed
67
68
    }
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
69
70
71
72
73
  if (numBytesPerTransfer % 4)
  {
    printf("[ERROR] numBytesPerTransfer (%lu) must be a multiple of 4\n", numBytesPerTransfer);
    exit(1);
  }
Gilbert Lee's avatar
Gilbert Lee committed
74

Gilbert Lee's avatar
Gilbert Lee committed
75
76
77
78
  // Check for preset tests
  // - Tests that sweep across possible sets of Transfers
  if (!strcmp(argv[1], "sweep") || !strcmp(argv[1], "rsweep"))
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
79
80
    int numGpuSubExecs = (argc > 3 ? atoi(argv[3]) : 4);
    int numCpuSubExecs = (argc > 4 ? atoi(argv[4]) : 4);
gilbertlee-amd's avatar
gilbertlee-amd committed
81

82
    ev.configMode = CFG_SWEEP;
gilbertlee-amd's avatar
gilbertlee-amd committed
83
    RunSweepPreset(ev, numBytesPerTransfer, numGpuSubExecs, numCpuSubExecs, !strcmp(argv[1], "rsweep"));
Gilbert Lee's avatar
Gilbert Lee committed
84
85
86
    exit(0);
  }
  // - Tests that benchmark peer-to-peer performance
gilbertlee-amd's avatar
gilbertlee-amd committed
87
  else if (!strcmp(argv[1], "p2p"))
Gilbert Lee's avatar
Gilbert Lee committed
88
  {
89
    ev.configMode = CFG_P2P;
gilbertlee-amd's avatar
gilbertlee-amd committed
90
    RunPeerToPeerBenchmarks(ev, numBytesPerTransfer / sizeof(float));
Gilbert Lee's avatar
Gilbert Lee committed
91
92
    exit(0);
  }
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
  // - Test SubExecutor scaling
  else if (!strcmp(argv[1], "scaling"))
  {
    int maxSubExecs = (argc > 3 ? atoi(argv[3]) : 32);
    int exeIndex    = (argc > 4 ? atoi(argv[4]) : 0);

    if (exeIndex >= ev.numGpuDevices)
    {
      printf("[ERROR] Cannot execute scaling test with GPU device %d\n", exeIndex);
      exit(1);
    }
    ev.configMode = CFG_SCALE;
    RunScalingBenchmark(ev, numBytesPerTransfer / sizeof(float), exeIndex, maxSubExecs);
    exit(0);
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
108
109
110
111
112
113
114
115
116
117
118
  // - Test all2all benchmark
  else if (!strcmp(argv[1], "a2a"))
  {
    int numSubExecs = (argc > 3 ? atoi(argv[3]) : 4);

    // Force single-stream mode for all-to-all benchmark
    ev.useSingleStream = 1;
    ev.configMode = CFG_A2A;
    RunAllToAllBenchmark(ev, numBytesPerTransfer, numSubExecs);
    exit(0);
  }
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
  else if (!strcmp(argv[1], "cmdline"))
  {
    // Print environment variables and CSV header
    ev.DisplayEnvVars();
    if (ev.outputToCsv)
    {
      printf("Test#,Transfer#,NumBytes,Src,Exe,Dst,CUs,BW(GB/s),Time(ms),SrcAddr,DstAddr\n");
    }

    // Read Transfer from command line
    std::string cmdlineTransfer;
    for (int i = 3; i < argc; i++)
      cmdlineTransfer += std::string(argv[i]) + " ";

    char line[2048];
    sprintf(line, "%s", cmdlineTransfer.c_str());
    std::vector<Transfer> transfers;
    ParseTransfers(line, ev.numCpuDevices, ev.numGpuDevices, transfers);
    if (transfers.empty()) exit(0);

    // If the number of bytes is specified, use it
    if (numBytesPerTransfer != 0)
    {
      size_t N = numBytesPerTransfer / sizeof(float);
      ExecuteTransfers(ev, 1, N, transfers);
    }
    else
    {
      // Otherwise generate a range of values
      for (int N = 256; N <= (1<<27); N *= 2)
      {
        int delta = std::max(1, N / ev.samplingFactor);
        int curr = N;
        while (curr < N * 2)
        {
          ExecuteTransfers(ev, 1, curr, transfers);
          curr += delta;
        }
      }
    }
    exit(0);
  }
Gilbert Lee's avatar
Gilbert Lee committed
161

Gilbert Lee's avatar
Gilbert Lee committed
162
  // Check that Transfer configuration file can be opened
163
  ev.configMode = CFG_FILE;
Gilbert Lee's avatar
Gilbert Lee committed
164
165
166
  FILE* fp = fopen(argv[1], "r");
  if (!fp)
  {
Gilbert Lee's avatar
Gilbert Lee committed
167
    printf("[ERROR] Unable to open transfer configuration file: [%s]\n", argv[1]);
Gilbert Lee's avatar
Gilbert Lee committed
168
169
170
    exit(1);
  }

Gilbert Lee's avatar
Gilbert Lee committed
171
  // Print environment variables and CSV header
Gilbert Lee's avatar
Gilbert Lee committed
172
173
174
  ev.DisplayEnvVars();
  if (ev.outputToCsv)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
175
    printf("Test#,Transfer#,NumBytes,Src,Exe,Dst,CUs,BW(GB/s),Time(ms),SrcAddr,DstAddr\n");
Gilbert Lee's avatar
Gilbert Lee committed
176
177
178
179
180
181
182
183
184
  }

  int testNum = 0;
  char line[2048];
  while(fgets(line, 2048, fp))
  {
    // Check if line is a comment to be echoed to output (starts with ##)
    if (!ev.outputToCsv && line[0] == '#' && line[1] == '#') printf("%s", line);

Gilbert Lee's avatar
Gilbert Lee committed
185
186
187
188
    // Parse set of parallel Transfers to execute
    std::vector<Transfer> transfers;
    ParseTransfers(line, ev.numCpuDevices, ev.numGpuDevices, transfers);
    if (transfers.empty()) continue;
Gilbert Lee's avatar
Gilbert Lee committed
189

gilbertlee-amd's avatar
gilbertlee-amd committed
190
191
192
193
194
195
196
197
198
199
200
    // If the number of bytes is specified, use it
    if (numBytesPerTransfer != 0)
    {
      size_t N = numBytesPerTransfer / sizeof(float);
      ExecuteTransfers(ev, ++testNum, N, transfers);
    }
    else
    {
      // Otherwise generate a range of values
      for (int N = 256; N <= (1<<27); N *= 2)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
201
        int delta = std::max(1, N / ev.samplingFactor);
gilbertlee-amd's avatar
gilbertlee-amd committed
202
203
204
        int curr = N;
        while (curr < N * 2)
        {
gilbertlee-amd's avatar
gilbertlee-amd committed
205
          ExecuteTransfers(ev, ++testNum, curr, transfers);
gilbertlee-amd's avatar
gilbertlee-amd committed
206
207
208
209
          curr += delta;
        }
      }
    }
Gilbert Lee's avatar
Gilbert Lee committed
210
211
  }
  fclose(fp);
Gilbert Lee's avatar
Gilbert Lee committed
212

Gilbert Lee's avatar
Gilbert Lee committed
213
214
  return 0;
}
Gilbert Lee's avatar
Gilbert Lee committed
215

Gilbert Lee's avatar
Gilbert Lee committed
216
void ExecuteTransfers(EnvVars const& ev,
gilbertlee-amd's avatar
gilbertlee-amd committed
217
218
219
                      int const testNum,
                      size_t const N,
                      std::vector<Transfer>& transfers,
gilbertlee-amd's avatar
gilbertlee-amd committed
220
221
                      bool verbose,
                      double* totalBandwidthCpu)
Gilbert Lee's avatar
Gilbert Lee committed
222
223
{
  int const initOffset = ev.byteOffset / sizeof(float);
Gilbert Lee's avatar
Gilbert Lee committed
224

Gilbert Lee's avatar
Gilbert Lee committed
225
226
  // Map transfers by executor
  TransferMap transferMap;
gilbertlee-amd's avatar
gilbertlee-amd committed
227
  for (int i = 0; i < transfers.size(); i++)
Gilbert Lee's avatar
Gilbert Lee committed
228
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
229
230
    Transfer& transfer = transfers[i];
    transfer.transferIndex = i;
gilbertlee-amd's avatar
gilbertlee-amd committed
231
    Executor executor(transfer.exeType, transfer.exeIndex);
Gilbert Lee's avatar
Gilbert Lee committed
232
    ExecutorInfo& executorInfo = transferMap[executor];
gilbertlee-amd's avatar
gilbertlee-amd committed
233
    executorInfo.transfers.push_back(&transfer);
Gilbert Lee's avatar
Gilbert Lee committed
234
  }
Gilbert Lee's avatar
Gilbert Lee committed
235

gilbertlee-amd's avatar
gilbertlee-amd committed
236
  // Loop over each executor and prepare sub-executors
gilbertlee-amd's avatar
gilbertlee-amd committed
237
  std::map<int, Transfer*> transferList;
Gilbert Lee's avatar
Gilbert Lee committed
238
239
240
  for (auto& exeInfoPair : transferMap)
  {
    Executor const& executor = exeInfoPair.first;
gilbertlee-amd's avatar
gilbertlee-amd committed
241
242
243
244
    ExecutorInfo& exeInfo    = exeInfoPair.second;
    ExeType const exeType    = executor.first;
    int     const exeIndex   = RemappedIndex(executor.second, IsCpuType(exeType));

Gilbert Lee's avatar
Gilbert Lee committed
245
    exeInfo.totalTime = 0.0;
gilbertlee-amd's avatar
gilbertlee-amd committed
246
    exeInfo.totalSubExecs = 0;
Gilbert Lee's avatar
Gilbert Lee committed
247
248

    // Loop over each transfer this executor is involved in
gilbertlee-amd's avatar
gilbertlee-amd committed
249
    for (Transfer* transfer : exeInfo.transfers)
Gilbert Lee's avatar
Gilbert Lee committed
250
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
251
252
253
254
255
256
      // Determine how many bytes to copy for this Transfer (use custom if pre-specified)
      transfer->numBytesActual = (transfer->numBytes ? transfer->numBytes : N * sizeof(float));

      // Allocate source memory
      transfer->srcMem.resize(transfer->numSrcs);
      for (int iSrc = 0; iSrc < transfer->numSrcs; ++iSrc)
Gilbert Lee's avatar
Gilbert Lee committed
257
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
258
259
260
        MemType const& srcType  = transfer->srcType[iSrc];
        int     const  srcIndex    = RemappedIndex(transfer->srcIndex[iSrc], IsCpuType(srcType));

Gilbert Lee's avatar
Gilbert Lee committed
261
        // Ensure executing GPU can access source memory
gilbertlee-amd's avatar
gilbertlee-amd committed
262
        if (IsGpuType(exeType) == MEM_GPU && IsGpuType(srcType) && srcIndex != exeIndex)
Gilbert Lee's avatar
Gilbert Lee committed
263
          EnablePeerAccess(exeIndex, srcIndex);
Gilbert Lee's avatar
Gilbert Lee committed
264

gilbertlee-amd's avatar
gilbertlee-amd committed
265
266
267
268
269
270
271
272
273
274
        AllocateMemory(srcType, srcIndex, transfer->numBytesActual + ev.byteOffset, (void**)&transfer->srcMem[iSrc]);
      }

      // Allocate destination memory
      transfer->dstMem.resize(transfer->numDsts);
      for (int iDst = 0; iDst < transfer->numDsts; ++iDst)
      {
        MemType const& dstType  = transfer->dstType[iDst];
        int     const  dstIndex    = RemappedIndex(transfer->dstIndex[iDst], IsCpuType(dstType));

Gilbert Lee's avatar
Gilbert Lee committed
275
        // Ensure executing GPU can access destination memory
gilbertlee-amd's avatar
gilbertlee-amd committed
276
        if (IsGpuType(exeType) == MEM_GPU && IsGpuType(dstType) && dstIndex != exeIndex)
Gilbert Lee's avatar
Gilbert Lee committed
277
278
          EnablePeerAccess(exeIndex, dstIndex);

gilbertlee-amd's avatar
gilbertlee-amd committed
279
280
        AllocateMemory(dstType, dstIndex, transfer->numBytesActual + ev.byteOffset, (void**)&transfer->dstMem[iDst]);
      }
Gilbert Lee's avatar
Gilbert Lee committed
281

gilbertlee-amd's avatar
gilbertlee-amd committed
282
      exeInfo.totalSubExecs += transfer->numSubExecs;
gilbertlee-amd's avatar
gilbertlee-amd committed
283
      transferList[transfer->transferIndex] = transfer;
Gilbert Lee's avatar
Gilbert Lee committed
284
285
    }

gilbertlee-amd's avatar
gilbertlee-amd committed
286
287
    // Prepare additional requirement for GPU-based executors
    if (IsGpuType(exeType))
Gilbert Lee's avatar
Gilbert Lee committed
288
    {
289
290
      HIP_CALL(hipSetDevice(exeIndex));

gilbertlee-amd's avatar
gilbertlee-amd committed
291
292
293
294
295
296
      // Single-stream is only supported for GFX-based executors
      int const numStreamsToUse = (exeType == EXE_GPU_DMA || !ev.useSingleStream) ? exeInfo.transfers.size() : 1;
      exeInfo.streams.resize(numStreamsToUse);
      exeInfo.startEvents.resize(numStreamsToUse);
      exeInfo.stopEvents.resize(numStreamsToUse);
      for (int i = 0; i < numStreamsToUse; ++i)
Gilbert Lee's avatar
Gilbert Lee committed
297
      {
298
299
300
301
302
303
304
305
306
307
        if (ev.cuMask.size())
        {
#if !defined(__NVCC__)
          HIP_CALL(hipExtStreamCreateWithCUMask(&exeInfo.streams[i], ev.cuMask.size(), ev.cuMask.data()));
#endif
        }
        else
        {
          HIP_CALL(hipStreamCreate(&exeInfo.streams[i]));
        }
Gilbert Lee's avatar
Gilbert Lee committed
308
309
310
        HIP_CALL(hipEventCreate(&exeInfo.startEvents[i]));
        HIP_CALL(hipEventCreate(&exeInfo.stopEvents[i]));
      }
Gilbert Lee's avatar
Gilbert Lee committed
311

gilbertlee-amd's avatar
gilbertlee-amd committed
312
      if (exeType == EXE_GPU_GFX)
Gilbert Lee's avatar
Gilbert Lee committed
313
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
314
315
        // Allocate one contiguous chunk of GPU memory for threadblock parameters
        // This allows support for executing one transfer per stream, or all transfers in a single stream
316
#if !defined(__NVCC__)
gilbertlee-amd's avatar
gilbertlee-amd committed
317
318
        AllocateMemory(MEM_GPU, exeIndex, exeInfo.totalSubExecs * sizeof(SubExecParam),
                       (void**)&exeInfo.subExecParamGpu);
319
320
321
322
#else
        AllocateMemory(MEM_CPU, exeIndex, exeInfo.totalSubExecs * sizeof(SubExecParam),
                       (void**)&exeInfo.subExecParamGpu);
#endif
Gilbert Lee's avatar
Gilbert Lee committed
323
324
325
      }
    }
  }
Gilbert Lee's avatar
Gilbert Lee committed
326

gilbertlee-amd's avatar
gilbertlee-amd committed
327
328
329
  if (verbose && !ev.outputToCsv) printf("Test %d:\n", testNum);

  // Prepare input memory and block parameters for current N
330
  bool isSrcCorrect = true;
gilbertlee-amd's avatar
gilbertlee-amd committed
331
  for (auto& exeInfoPair : transferMap)
Gilbert Lee's avatar
Gilbert Lee committed
332
  {
333
334
335
336
337
    Executor const& executor = exeInfoPair.first;
    ExecutorInfo& exeInfo    = exeInfoPair.second;
    ExeType const exeType    = executor.first;
    int     const exeIndex   = RemappedIndex(executor.second, IsCpuType(exeType));

gilbertlee-amd's avatar
gilbertlee-amd committed
338
    exeInfo.totalBytes = 0;
Gilbert Lee's avatar
Gilbert Lee committed
339

gilbertlee-amd's avatar
gilbertlee-amd committed
340
341
    int transferOffset = 0;
    for (int i = 0; i < exeInfo.transfers.size(); ++i)
Gilbert Lee's avatar
Gilbert Lee committed
342
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
343
344
      // Prepare subarrays each threadblock works on and fill src memory with patterned data
      Transfer* transfer = exeInfo.transfers[i];
gilbertlee-amd's avatar
gilbertlee-amd committed
345
      transfer->PrepareSubExecParams(ev);
346
      isSrcCorrect &= transfer->PrepareSrc(ev);
gilbertlee-amd's avatar
gilbertlee-amd committed
347
      exeInfo.totalBytes += transfer->numBytesActual;
Gilbert Lee's avatar
Gilbert Lee committed
348

gilbertlee-amd's avatar
gilbertlee-amd committed
349
      // Copy block parameters to GPU for GPU executors
gilbertlee-amd's avatar
gilbertlee-amd committed
350
      if (transfer->exeType == EXE_GPU_GFX)
Gilbert Lee's avatar
Gilbert Lee committed
351
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
352
        exeInfo.transfers[i]->subExecParamGpuPtr = exeInfo.subExecParamGpu + transferOffset;
353
        HIP_CALL(hipSetDevice(exeIndex));
gilbertlee-amd's avatar
gilbertlee-amd committed
354
355
356
        HIP_CALL(hipMemcpy(&exeInfo.subExecParamGpu[transferOffset],
                           transfer->subExecParam.data(),
                           transfer->subExecParam.size() * sizeof(SubExecParam),
gilbertlee-amd's avatar
gilbertlee-amd committed
357
                           hipMemcpyHostToDevice));
358
359
        HIP_CALL(hipDeviceSynchronize());

gilbertlee-amd's avatar
gilbertlee-amd committed
360
        transferOffset += transfer->subExecParam.size();
Gilbert Lee's avatar
Gilbert Lee committed
361
362
      }
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
363
  }
Gilbert Lee's avatar
Gilbert Lee committed
364

gilbertlee-amd's avatar
gilbertlee-amd committed
365
366
367
368
  // Launch kernels (warmup iterations are not counted)
  double totalCpuTime = 0;
  size_t numTimedIterations = 0;
  std::stack<std::thread> threads;
369
  for (int iteration = -ev.numWarmups; isSrcCorrect; iteration++)
gilbertlee-amd's avatar
gilbertlee-amd committed
370
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
371
    if (ev.numIterations > 0 && iteration    >= ev.numIterations) break;
gilbertlee-amd's avatar
gilbertlee-amd committed
372
    if (ev.numIterations < 0 && totalCpuTime > -ev.numIterations) break;
Gilbert Lee's avatar
Gilbert Lee committed
373

gilbertlee-amd's avatar
gilbertlee-amd committed
374
375
    // Pause before starting first timed iteration in interactive mode
    if (verbose && ev.useInteractive && iteration == 0)
Gilbert Lee's avatar
Gilbert Lee committed
376
    {
377
378
379
380
      printf("Memory prepared:\n");

      for (Transfer& transfer : transfers)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
381
382
383
384
385
        printf("Transfer %03d:\n", transfer.transferIndex);
        for (int iSrc = 0; iSrc < transfer.numSrcs; ++iSrc)
          printf("  SRC %0d: %p\n", iSrc, transfer.srcMem[iSrc]);
        for (int iDst = 0; iDst < transfer.numDsts; ++iDst)
          printf("  DST %0d: %p\n", iDst, transfer.dstMem[iDst]);
386
      }
gilbertlee-amd's avatar
gilbertlee-amd committed
387
      printf("Hit <Enter> to continue: ");
388
389
390
391
392
      if (scanf("%*c") != 0)
      {
        printf("[ERROR] Unexpected input\n");
        exit(1);
      }
Gilbert Lee's avatar
Gilbert Lee committed
393
394
      printf("\n");
    }
Gilbert Lee's avatar
Gilbert Lee committed
395

gilbertlee-amd's avatar
gilbertlee-amd committed
396
397
398
399
400
    // Start CPU timing for this iteration
    auto cpuStart = std::chrono::high_resolution_clock::now();

    // Execute all Transfers in parallel
    for (auto& exeInfoPair : transferMap)
401
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
402
      ExecutorInfo& exeInfo = exeInfoPair.second;
gilbertlee-amd's avatar
gilbertlee-amd committed
403
404
405
      ExeType       exeType = exeInfoPair.first.first;
      int const numTransfersToRun = (exeType == EXE_GPU_GFX && ev.useSingleStream) ? 1 : exeInfo.transfers.size();

gilbertlee-amd's avatar
gilbertlee-amd committed
406
407
      for (int i = 0; i < numTransfersToRun; ++i)
        threads.push(std::thread(RunTransfer, std::ref(ev), iteration, std::ref(exeInfo), i));
408
    }
Gilbert Lee's avatar
Gilbert Lee committed
409

gilbertlee-amd's avatar
gilbertlee-amd committed
410
411
412
413
414
415
416
    // Wait for all threads to finish
    int const numTransfers = threads.size();
    for (int i = 0; i < numTransfers; i++)
    {
      threads.top().join();
      threads.pop();
    }
Gilbert Lee's avatar
Gilbert Lee committed
417

gilbertlee-amd's avatar
gilbertlee-amd committed
418
419
420
421
422
    // Stop CPU timing for this iteration
    auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
    double deltaSec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count();

    if (iteration >= 0)
Gilbert Lee's avatar
Gilbert Lee committed
423
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
424
425
426
427
      ++numTimedIterations;
      totalCpuTime += deltaSec;
    }
  }
Gilbert Lee's avatar
Gilbert Lee committed
428

gilbertlee-amd's avatar
gilbertlee-amd committed
429
  // Pause for interactive mode
430
  if (verbose && isSrcCorrect && ev.useInteractive)
gilbertlee-amd's avatar
gilbertlee-amd committed
431
432
  {
    printf("Transfers complete. Hit <Enter> to continue: ");
433
434
435
436
437
    if (scanf("%*c") != 0)
    {
      printf("[ERROR] Unexpected input\n");
      exit(1);
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
438
439
    printf("\n");
  }
Gilbert Lee's avatar
Gilbert Lee committed
440

gilbertlee-amd's avatar
gilbertlee-amd committed
441
442
443
  // Validate that each transfer has transferred correctly
  size_t totalBytesTransferred = 0;
  int const numTransfers = transferList.size();
gilbertlee-amd's avatar
gilbertlee-amd committed
444

gilbertlee-amd's avatar
gilbertlee-amd committed
445
446
447
  for (auto transferPair : transferList)
  {
    Transfer* transfer = transferPair.second;
gilbertlee-amd's avatar
gilbertlee-amd committed
448
449
    transfer->ValidateDst(ev);
    totalBytesTransferred += transfer->numBytesActual;
gilbertlee-amd's avatar
gilbertlee-amd committed
450
  }
Gilbert Lee's avatar
Gilbert Lee committed
451

gilbertlee-amd's avatar
gilbertlee-amd committed
452
453
454
  // Report timings
  totalCpuTime = totalCpuTime / (1.0 * numTimedIterations) * 1000;
  double totalBandwidthGbs = (totalBytesTransferred / 1.0E6) / totalCpuTime;
gilbertlee-amd's avatar
gilbertlee-amd committed
455
456
  if (totalBandwidthCpu) *totalBandwidthCpu = totalBandwidthGbs;

gilbertlee-amd's avatar
gilbertlee-amd committed
457
  double maxGpuTime = 0;
Gilbert Lee's avatar
Gilbert Lee committed
458

459
  if (!isSrcCorrect) goto cleanup;
gilbertlee-amd's avatar
gilbertlee-amd committed
460
461
462
463
  if (ev.useSingleStream)
  {
    for (auto& exeInfoPair : transferMap)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
464
465
466
      ExecutorInfo  exeInfo  = exeInfoPair.second;
      ExeType const exeType  = exeInfoPair.first.first;
      int     const exeIndex = exeInfoPair.first.second;
Gilbert Lee's avatar
Gilbert Lee committed
467

gilbertlee-amd's avatar
gilbertlee-amd committed
468
469
      // Compute total time for non GPU executors
      if (exeType != EXE_GPU_GFX)
gilbertlee-amd's avatar
gilbertlee-amd committed
470
471
472
473
474
      {
        exeInfo.totalTime = 0;
        for (auto const& transfer : exeInfo.transfers)
          exeInfo.totalTime = std::max(exeInfo.totalTime, transfer->transferTime);
      }
475

gilbertlee-amd's avatar
gilbertlee-amd committed
476
477
478
      double exeDurationMsec = exeInfo.totalTime / (1.0 * numTimedIterations);
      double exeBandwidthGbs = (exeInfo.totalBytes / 1.0E9) / exeDurationMsec * 1000.0f;
      maxGpuTime = std::max(maxGpuTime, exeDurationMsec);
Gilbert Lee's avatar
Gilbert Lee committed
479

gilbertlee-amd's avatar
gilbertlee-amd committed
480
481
      if (verbose && !ev.outputToCsv)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
482
483
        printf(" Executor: %3s %02d | %7.3f GB/s | %8.3f ms | %12lu bytes\n",
               ExeTypeName[exeType], exeIndex, exeBandwidthGbs, exeDurationMsec, exeInfo.totalBytes);
Gilbert Lee's avatar
Gilbert Lee committed
484
      }
gilbertlee-amd's avatar
gilbertlee-amd committed
485
486
487

      int totalCUs = 0;
      for (auto const& transfer : exeInfo.transfers)
Gilbert Lee's avatar
Gilbert Lee committed
488
      {
Gilbert Lee's avatar
Gilbert Lee committed
489
        double transferDurationMsec = transfer->transferTime / (1.0 * numTimedIterations);
gilbertlee-amd's avatar
gilbertlee-amd committed
490
491
        double transferBandwidthGbs = (transfer->numBytesActual / 1.0E9) / transferDurationMsec * 1000.0f;
        totalCUs += transfer->numSubExecs;
gilbertlee-amd's avatar
gilbertlee-amd committed
492
493

        if (!verbose) continue;
Gilbert Lee's avatar
Gilbert Lee committed
494
495
        if (!ev.outputToCsv)
        {
gilbertlee-amd's avatar
gilbertlee-amd committed
496
          printf("     Transfer %02d  | %7.3f GB/s | %8.3f ms | %12lu bytes | %s -> %s%02d:%03d -> %s\n",
Gilbert Lee's avatar
Gilbert Lee committed
497
                 transfer->transferIndex,
gilbertlee-amd's avatar
gilbertlee-amd committed
498
499
                 transferBandwidthGbs,
                 transferDurationMsec,
gilbertlee-amd's avatar
gilbertlee-amd committed
500
501
502
503
504
                 transfer->numBytesActual,
                 transfer->SrcToStr().c_str(),
                 ExeTypeName[transfer->exeType], transfer->exeIndex,
                 transfer->numSubExecs,
                 transfer->DstToStr().c_str());
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531

          if (ev.showIterations)
          {
            std::set<std::pair<double, int>> times;
            double stdDevTime = 0;
            double stdDevBw = 0;
            for (int i = 0; i < numTimedIterations; i++)
            {
              times.insert(std::make_pair(transfer->perIterationTime[i], i+1));
              double const varTime = fabs(transferDurationMsec - transfer->perIterationTime[i]);
              stdDevTime += varTime * varTime;

              double iterBandwidthGbs = (transfer->numBytesActual / 1.0E9) / transfer->perIterationTime[i] * 1000.0f;
              double const varBw = fabs(iterBandwidthGbs - transferBandwidthGbs);
              stdDevBw += varBw * varBw;
            }
            stdDevTime = sqrt(stdDevTime / numTimedIterations);
            stdDevBw = sqrt(stdDevBw / numTimedIterations);

            for (auto t : times)
            {
              double iterDurationMsec = t.first;
              double iterBandwidthGbs = (transfer->numBytesActual / 1.0E9) / iterDurationMsec * 1000.0f;
              printf("      Iter %03d    | %7.3f GB/s | %8.3f ms |\n", t.second, iterBandwidthGbs, iterDurationMsec);
            }
            printf("      StandardDev | %7.3f GB/s | %8.3f ms |\n", stdDevBw, stdDevTime);
          }
Gilbert Lee's avatar
Gilbert Lee committed
532
533
534
        }
        else
        {
gilbertlee-amd's avatar
gilbertlee-amd committed
535
536
537
538
539
540
          printf("%d,%d,%lu,%s,%c%02d,%s,%d,%.3f,%.3f,%s,%s\n",
                 testNum, transfer->transferIndex, transfer->numBytesActual,
                 transfer->SrcToStr().c_str(),
                 MemTypeStr[transfer->exeType], transfer->exeIndex,
                 transfer->DstToStr().c_str(),
                 transfer->numSubExecs,
Gilbert Lee's avatar
Gilbert Lee committed
541
                 transferBandwidthGbs, transferDurationMsec,
gilbertlee-amd's avatar
gilbertlee-amd committed
542
543
                 PtrVectorToStr(transfer->srcMem, initOffset).c_str(),
                 PtrVectorToStr(transfer->dstMem, initOffset).c_str());
Gilbert Lee's avatar
Gilbert Lee committed
544
        }
Gilbert Lee's avatar
Gilbert Lee committed
545
      }
gilbertlee-amd's avatar
gilbertlee-amd committed
546
547
548

      if (verbose && ev.outputToCsv)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
549
        printf("%d,ALL,%lu,ALL,%c%02d,ALL,%d,%.3f,%.3f,ALL,ALL\n",
gilbertlee-amd's avatar
gilbertlee-amd committed
550
               testNum, totalBytesTransferred,
gilbertlee-amd's avatar
gilbertlee-amd committed
551
               MemTypeStr[exeType], exeIndex, totalCUs,
gilbertlee-amd's avatar
gilbertlee-amd committed
552
553
               exeBandwidthGbs, exeDurationMsec);
      }
Gilbert Lee's avatar
Gilbert Lee committed
554
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
555
556
557
558
559
560
561
  }
  else
  {
    for (auto const& transferPair : transferList)
    {
      Transfer* transfer = transferPair.second;
      double transferDurationMsec = transfer->transferTime / (1.0 * numTimedIterations);
gilbertlee-amd's avatar
gilbertlee-amd committed
562
      double transferBandwidthGbs = (transfer->numBytesActual / 1.0E9) / transferDurationMsec * 1000.0f;
gilbertlee-amd's avatar
gilbertlee-amd committed
563
564
565
566
      maxGpuTime = std::max(maxGpuTime, transferDurationMsec);
      if (!verbose) continue;
      if (!ev.outputToCsv)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
567
        printf(" Transfer %02d      | %7.3f GB/s | %8.3f ms | %12lu bytes | %s -> %s%02d:%03d -> %s\n",
gilbertlee-amd's avatar
gilbertlee-amd committed
568
569
               transfer->transferIndex,
               transferBandwidthGbs, transferDurationMsec,
gilbertlee-amd's avatar
gilbertlee-amd committed
570
571
572
573
574
               transfer->numBytesActual,
               transfer->SrcToStr().c_str(),
               ExeTypeName[transfer->exeType], transfer->exeIndex,
               transfer->numSubExecs,
               transfer->DstToStr().c_str());
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597

        if (ev.showIterations)
        {
            std::set<std::pair<double, int>> times;
            double stdDevTime = 0;
            double stdDevBw = 0;
            for (int i = 0; i < numTimedIterations; i++)
            {
              times.insert(std::make_pair(transfer->perIterationTime[i], i+1));
              double const varTime = fabs(transferDurationMsec - transfer->perIterationTime[i]);
              stdDevTime += varTime * varTime;

              double iterBandwidthGbs = (transfer->numBytesActual / 1.0E9) / transfer->perIterationTime[i] * 1000.0f;
              double const varBw = fabs(iterBandwidthGbs - transferBandwidthGbs);
              stdDevBw += varBw * varBw;
            }
            stdDevTime = sqrt(stdDevTime / numTimedIterations);
            stdDevBw = sqrt(stdDevBw / numTimedIterations);

            for (auto t : times)
            {
              double iterDurationMsec = t.first;
              double iterBandwidthGbs = (transfer->numBytesActual / 1.0E9) / iterDurationMsec * 1000.0f;
598
599
600
601
              printf("      Iter %03d    | %7.3f GB/s | %8.3f ms | CUs:", t.second, iterBandwidthGbs, iterDurationMsec);
              for (auto x : transfer->perIterationCUs[t.second - 1])
                printf(" %2d", x);
              printf("\n");
602
603
604
            }
            printf("      StandardDev | %7.3f GB/s | %8.3f ms |\n", stdDevBw, stdDevTime);
        }
gilbertlee-amd's avatar
gilbertlee-amd committed
605
606
607
      }
      else
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
608
609
610
611
612
613
        printf("%d,%d,%lu,%s,%s%02d,%s,%d,%.3f,%.3f,%s,%s\n",
               testNum, transfer->transferIndex, transfer->numBytesActual,
               transfer->SrcToStr().c_str(),
               ExeTypeName[transfer->exeType], transfer->exeIndex,
               transfer->DstToStr().c_str(),
               transfer->numSubExecs,
gilbertlee-amd's avatar
gilbertlee-amd committed
614
               transferBandwidthGbs, transferDurationMsec,
gilbertlee-amd's avatar
gilbertlee-amd committed
615
616
               PtrVectorToStr(transfer->srcMem, initOffset).c_str(),
               PtrVectorToStr(transfer->dstMem, initOffset).c_str());
gilbertlee-amd's avatar
gilbertlee-amd committed
617
618
619
      }
    }
  }
Gilbert Lee's avatar
Gilbert Lee committed
620

gilbertlee-amd's avatar
gilbertlee-amd committed
621
622
623
  // Display aggregate statistics
  if (verbose)
  {
Gilbert Lee's avatar
Gilbert Lee committed
624
    if (!ev.outputToCsv)
Gilbert Lee's avatar
Gilbert Lee committed
625
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
626
      printf(" Aggregate (CPU)  | %7.3f GB/s | %8.3f ms | %12lu bytes | Overhead: %.3f ms\n",
627
             totalBandwidthGbs, totalCpuTime, totalBytesTransferred, totalCpuTime - maxGpuTime);
Gilbert Lee's avatar
Gilbert Lee committed
628
629
630
    }
    else
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
631
      printf("%d,ALL,%lu,ALL,ALL,ALL,ALL,%.3f,%.3f,ALL,ALL\n",
632
             testNum, totalBytesTransferred, totalBandwidthGbs, totalCpuTime);
Gilbert Lee's avatar
Gilbert Lee committed
633
634
    }
  }
Gilbert Lee's avatar
Gilbert Lee committed
635

Gilbert Lee's avatar
Gilbert Lee committed
636
  // Release GPU memory
637
cleanup:
Gilbert Lee's avatar
Gilbert Lee committed
638
639
  for (auto exeInfoPair : transferMap)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
640
641
642
643
    ExecutorInfo& exeInfo  = exeInfoPair.second;
    ExeType const exeType  = exeInfoPair.first.first;
    int     const exeIndex = RemappedIndex(exeInfoPair.first.second, IsCpuType(exeType));

Gilbert Lee's avatar
Gilbert Lee committed
644
645
    for (auto& transfer : exeInfo.transfers)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
646
647
648
649
650
651
652
653
654
655
656
      for (int iSrc = 0; iSrc < transfer->numSrcs; ++iSrc)
      {
        MemType const& srcType = transfer->srcType[iSrc];
        DeallocateMemory(srcType, transfer->srcMem[iSrc], transfer->numBytesActual + ev.byteOffset);
      }
      for (int iDst = 0; iDst < transfer->numDsts; ++iDst)
      {
        MemType const& dstType = transfer->dstType[iDst];
        DeallocateMemory(dstType, transfer->dstMem[iDst], transfer->numBytesActual + ev.byteOffset);
      }
      transfer->subExecParam.clear();
Gilbert Lee's avatar
Gilbert Lee committed
657
658
    }

gilbertlee-amd's avatar
gilbertlee-amd committed
659
    if (IsGpuType(exeType))
Gilbert Lee's avatar
Gilbert Lee committed
660
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
661
662
      int const numStreams = (int)exeInfo.streams.size();
      for (int i = 0; i < numStreams; ++i)
Gilbert Lee's avatar
Gilbert Lee committed
663
      {
Gilbert Lee's avatar
Gilbert Lee committed
664
665
666
        HIP_CALL(hipEventDestroy(exeInfo.startEvents[i]));
        HIP_CALL(hipEventDestroy(exeInfo.stopEvents[i]));
        HIP_CALL(hipStreamDestroy(exeInfo.streams[i]));
Gilbert Lee's avatar
Gilbert Lee committed
667
      }
gilbertlee-amd's avatar
gilbertlee-amd committed
668
669
670

      if (exeType == EXE_GPU_GFX)
      {
671
#if !defined(__NVCC__)
gilbertlee-amd's avatar
gilbertlee-amd committed
672
        DeallocateMemory(MEM_GPU, exeInfo.subExecParamGpu);
673
674
675
#else
        DeallocateMemory(MEM_CPU, exeInfo.subExecParamGpu);
#endif
gilbertlee-amd's avatar
gilbertlee-amd committed
676
      }
Gilbert Lee's avatar
Gilbert Lee committed
677
678
679
680
681
682
    }
  }
}

void DisplayUsage(char const* cmdName)
{
Gilbert Lee's avatar
Gilbert Lee committed
683
  printf("TransferBench v%s\n", TB_VERSION);
Gilbert Lee's avatar
Gilbert Lee committed
684
685
686
687
688
689
690
691
692
693
694
695
696
  printf("========================================\n");

  if (numa_available() == -1)
  {
    printf("[ERROR] NUMA library not supported. Check to see if libnuma has been installed on this system\n");
    exit(1);
  }
  int numGpuDevices;
  HIP_CALL(hipGetDeviceCount(&numGpuDevices));
  int const numCpuDevices = numa_num_configured_nodes();

  printf("Usage: %s config <N>\n", cmdName);
  printf("  config: Either:\n");
Gilbert Lee's avatar
Gilbert Lee committed
697
  printf("          - Filename of configFile containing Transfers to execute (see example.cfg for format)\n");
gilbertlee-amd's avatar
gilbertlee-amd committed
698
699
700
  printf("          - Name of preset config:\n");
  printf("              p2p          - Peer-to-peer benchmark tests\n");
  printf("              sweep/rsweep - Sweep/random sweep across possible sets of Transfers\n");
701
702
703
704
705
  printf("                             - 3rd optional arg: # GPU SubExecs per Transfer\n");
  printf("                             - 4th optional arg: # CPU SubExecs per Transfer\n");
  printf("              scaling      - GPU SubExec scaling copy test\n");
  printf("                             - 3th optional arg: Max # of SubExecs to use\n");
  printf("                             - 4rd optional arg: GPU index to use as executor\n");
gilbertlee-amd's avatar
gilbertlee-amd committed
706
707
  printf("              a2a          - GPU All-To-All benchmark\n");
  printf("                             - 3rd optional arg: # of SubExecs to use\n");
708
  printf("              cmdline      - Read Transfers from command line arguments (after N)\n");
Gilbert Lee's avatar
Gilbert Lee committed
709
  printf("  N     : (Optional) Number of bytes to copy per Transfer.\n");
Gilbert Lee's avatar
Gilbert Lee committed
710
  printf("          If not specified, defaults to %lu bytes. Must be a multiple of 4 bytes\n",
Gilbert Lee's avatar
Gilbert Lee committed
711
         DEFAULT_BYTES_PER_TRANSFER);
Gilbert Lee's avatar
Gilbert Lee committed
712
713
714
715
716
717
718
  printf("          If 0 is specified, a range of Ns will be benchmarked\n");
  printf("          May append a suffix ('K', 'M', 'G') for kilobytes / megabytes / gigabytes\n");
  printf("\n");

  EnvVars::DisplayUsage();
}

gilbertlee-amd's avatar
gilbertlee-amd committed
719
int RemappedIndex(int const origIdx, bool const isCpuType)
Gilbert Lee's avatar
Gilbert Lee committed
720
{
721
722
  static std::vector<int> remappingCpu;
  static std::vector<int> remappingGpu;
Gilbert Lee's avatar
Gilbert Lee committed
723

724
725
726
727
728
729
730
731
  // Build CPU remapping on first use
  // Skip numa nodes that are not configured
  if (remappingCpu.empty())
  {
    for (int node = 0; node <= numa_max_node(); node++)
      if (numa_bitmask_isbitset(numa_get_mems_allowed(), node))
        remappingCpu.push_back(node);
  }
Gilbert Lee's avatar
Gilbert Lee committed
732

733
734
  // Build remappingGpu on first use
  if (remappingGpu.empty())
Gilbert Lee's avatar
Gilbert Lee committed
735
736
737
  {
    int numGpuDevices;
    HIP_CALL(hipGetDeviceCount(&numGpuDevices));
738
    remappingGpu.resize(numGpuDevices);
Gilbert Lee's avatar
Gilbert Lee committed
739
740
741
742

    int const usePcieIndexing = getenv("USE_PCIE_INDEX") ? atoi(getenv("USE_PCIE_INDEX")) : 0;
    if (!usePcieIndexing)
    {
743
      // For HIP-based indexing no remappingGpu is necessary
Gilbert Lee's avatar
Gilbert Lee committed
744
      for (int i = 0; i < numGpuDevices; ++i)
745
        remappingGpu[i] = i;
Gilbert Lee's avatar
Gilbert Lee committed
746
747
748
749
750
751
752
753
754
755
756
757
758
759
    }
    else
    {
      // Collect PCIe address for each GPU
      std::vector<std::pair<std::string, int>> mapping;
      char pciBusId[20];
      for (int i = 0; i < numGpuDevices; ++i)
      {
        HIP_CALL(hipDeviceGetPCIBusId(pciBusId, 20, i));
        mapping.push_back(std::make_pair(pciBusId, i));
      }
      // Sort GPUs by PCIe address then use that as mapping
      std::sort(mapping.begin(), mapping.end());
      for (int i = 0; i < numGpuDevices; ++i)
760
        remappingGpu[i] = mapping[i].second;
Gilbert Lee's avatar
Gilbert Lee committed
761
762
    }
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
763
  return isCpuType ? remappingCpu[origIdx] : remappingGpu[origIdx];
Gilbert Lee's avatar
Gilbert Lee committed
764
765
766
767
}

void DisplayTopology(bool const outputToCsv)
{
768

769
  int numCpuDevices = numa_num_configured_nodes();
Gilbert Lee's avatar
Gilbert Lee committed
770
771
772
773
774
  int numGpuDevices;
  HIP_CALL(hipGetDeviceCount(&numGpuDevices));

  if (outputToCsv)
  {
775
    printf("NumCpus,%d\n", numCpuDevices);
Gilbert Lee's avatar
Gilbert Lee committed
776
    printf("NumGpus,%d\n", numGpuDevices);
777
778
779
  }
  else
  {
780
781
    printf("\nDetected topology: %d configured CPU NUMA node(s) [%d total]   %d GPU device(s)\n",
           numa_num_configured_nodes(), numa_max_node() + 1, numGpuDevices);
782
783
784
785
786
787
788
789
  }

  // Print out detected CPU topology
  if (outputToCsv)
  {
    printf("NUMA");
    for (int j = 0; j < numCpuDevices; j++)
      printf(",NUMA%02d", j);
790
    printf(",# CPUs,ClosestGPUs,ActualNode\n");
791
792
793
  }
  else
  {
794
    printf("            |");
795
    for (int j = 0; j < numCpuDevices; j++)
796
797
798
799
      printf("NUMA %02d|", j);
    printf(" #Cpus | Closest GPU(s)\n");

    printf("------------+");
800
    for (int j = 0; j <= numCpuDevices; j++)
801
802
      printf("-------+");
    printf("---------------\n");
803
804
805
806
  }

  for (int i = 0; i < numCpuDevices; i++)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
807
    int nodeI = RemappedIndex(i, true);
808
    printf("NUMA %02d (%02d)%s", i, nodeI, outputToCsv ? "," : "|");
809
810
    for (int j = 0; j < numCpuDevices; j++)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
811
      int nodeJ = RemappedIndex(j, true);
812
      int numaDist = numa_distance(nodeI, nodeJ);
813
      if (outputToCsv)
gilbertlee-amd's avatar
gilbertlee-amd committed
814
        printf("%d,", numaDist);
815
      else
816
        printf(" %5d |", numaDist);
817
818
819
820
    }

    int numCpus = 0;
    for (int j = 0; j < numa_num_configured_cpus(); j++)
821
      if (numa_node_of_cpu(j) == nodeI) numCpus++;
822
823
824
    if (outputToCsv)
      printf("%d,", numCpus);
    else
825
      printf(" %5d | ", numCpus);
826

827
#if !defined(__NVCC__)
828
829
830
    bool isFirst = true;
    for (int j = 0; j < numGpuDevices; j++)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
831
      if (GetClosestNumaNode(RemappedIndex(j, false)) == i)
832
833
      {
        if (isFirst) isFirst = false;
gilbertlee-amd's avatar
gilbertlee-amd committed
834
835
        else printf(",");
        printf("%d", j);
836
837
      }
    }
838
#endif
839
840
841
842
    printf("\n");
  }
  printf("\n");

843
844
845
846
847
#if defined(__NVCC__)
  // No further topology detection done for NVIDIA platforms
  return;
#endif

848
849
850
  // Print out detected GPU topology
  if (outputToCsv)
  {
Gilbert Lee's avatar
Gilbert Lee committed
851
852
853
854
855
856
857
    printf("GPU");
    for (int j = 0; j < numGpuDevices; j++)
      printf(",GPU %02d", j);
    printf(",PCIe Bus ID,ClosestNUMA\n");
  }
  else
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
858
859
860
861
862
863
864
865
866
867
    printf("        |");
    for (int j = 0; j < numGpuDevices; j++)
    {
      hipDeviceProp_t prop;
      HIP_CALL(hipGetDeviceProperties(&prop, j));
      std::string fullName = prop.gcnArchName;
      std::string archName = fullName.substr(0, fullName.find(':'));
      printf(" %6s |", archName.c_str());
    }
    printf("\n");
Gilbert Lee's avatar
Gilbert Lee committed
868
869
870
    printf("        |");
    for (int j = 0; j < numGpuDevices; j++)
      printf(" GPU %02d |", j);
gilbertlee-amd's avatar
gilbertlee-amd committed
871
    printf(" PCIe Bus ID  | #CUs | Closest NUMA\n");
Gilbert Lee's avatar
Gilbert Lee committed
872
873
    for (int j = 0; j <= numGpuDevices; j++)
      printf("--------+");
gilbertlee-amd's avatar
gilbertlee-amd committed
874
    printf("--------------+------+-------------\n");
Gilbert Lee's avatar
Gilbert Lee committed
875
876
  }

877
#if !defined(__NVCC__)
Gilbert Lee's avatar
Gilbert Lee committed
878
879
880
  char pciBusId[20];
  for (int i = 0; i < numGpuDevices; i++)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
881
    int const deviceIdx = RemappedIndex(i, false);
Gilbert Lee's avatar
Gilbert Lee committed
882
883
884
885
886
887
888
889
890
891
892
893
894
    printf("%sGPU %02d%s", outputToCsv ? "" : " ", i, outputToCsv ? "," : " |");
    for (int j = 0; j < numGpuDevices; j++)
    {
      if (i == j)
      {
        if (outputToCsv)
          printf("-,");
        else
          printf("    -   |");
      }
      else
      {
        uint32_t linkType, hopCount;
gilbertlee-amd's avatar
gilbertlee-amd committed
895
896
        HIP_CALL(hipExtGetLinkTypeAndHopCount(deviceIdx,
                                              RemappedIndex(j, false),
Gilbert Lee's avatar
Gilbert Lee committed
897
898
899
900
901
902
903
904
905
906
907
                                              &linkType, &hopCount));
        printf("%s%s-%d%s",
               outputToCsv ? "" : " ",
               linkType == HSA_AMD_LINK_INFO_TYPE_HYPERTRANSPORT ? "  HT" :
               linkType == HSA_AMD_LINK_INFO_TYPE_QPI            ? " QPI" :
               linkType == HSA_AMD_LINK_INFO_TYPE_PCIE           ? "PCIE" :
               linkType == HSA_AMD_LINK_INFO_TYPE_INFINBAND      ? "INFB" :
               linkType == HSA_AMD_LINK_INFO_TYPE_XGMI           ? "XGMI" : "????",
               hopCount, outputToCsv ? "," : " |");
      }
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
908
909
910
911
912
    HIP_CALL(hipDeviceGetPCIBusId(pciBusId, 20, deviceIdx));

    int numDeviceCUs = 0;
    HIP_CALL(hipDeviceGetAttribute(&numDeviceCUs, hipDeviceAttributeMultiprocessorCount, deviceIdx));

Gilbert Lee's avatar
Gilbert Lee committed
913
    if (outputToCsv)
gilbertlee-amd's avatar
gilbertlee-amd committed
914
      printf("%s,%d,%d\n", pciBusId, numDeviceCUs, GetClosestNumaNode(deviceIdx));
Gilbert Lee's avatar
Gilbert Lee committed
915
    else
gilbertlee-amd's avatar
gilbertlee-amd committed
916
      printf(" %11s | %4d | %d\n", pciBusId, numDeviceCUs, GetClosestNumaNode(deviceIdx));
Gilbert Lee's avatar
Gilbert Lee committed
917
  }
918
#endif
Gilbert Lee's avatar
Gilbert Lee committed
919
920
}

gilbertlee-amd's avatar
gilbertlee-amd committed
921
922
void ParseMemType(std::string const& token, int const numCpus, int const numGpus,
                  std::vector<MemType>& memTypes, std::vector<int>& memIndices)
Gilbert Lee's avatar
Gilbert Lee committed
923
924
{
  char typeChar;
gilbertlee-amd's avatar
gilbertlee-amd committed
925
926
  int offset = 0, devIndex, inc;
  bool found = false;
Gilbert Lee's avatar
Gilbert Lee committed
927

gilbertlee-amd's avatar
gilbertlee-amd committed
928
929
930
  memTypes.clear();
  memIndices.clear();
  while (sscanf(token.c_str() + offset, " %c %d%n", &typeChar, &devIndex, &inc) == 2)
Gilbert Lee's avatar
Gilbert Lee committed
931
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
932
933
934
935
    offset += inc;
    MemType memType = CharToMemType(typeChar);

    if (IsCpuType(memType) && (devIndex < 0 || devIndex >= numCpus))
Gilbert Lee's avatar
Gilbert Lee committed
936
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
937
      printf("[ERROR] CPU index must be between 0 and %d (instead of %d)\n", numCpus-1, devIndex);
Gilbert Lee's avatar
Gilbert Lee committed
938
939
      exit(1);
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
940
    if (IsGpuType(memType) && (devIndex < 0 || devIndex >= numGpus))
Gilbert Lee's avatar
Gilbert Lee committed
941
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
942
      printf("[ERROR] GPU index must be between 0 and %d (instead of %d)\n", numGpus-1, devIndex);
Gilbert Lee's avatar
Gilbert Lee committed
943
944
      exit(1);
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
945
946
947
948
949
950
951
952
953
954
955
956
957
958
959
960
961
962
963
964
965
966
967
968
969
970
971
972
973
974
975
976
977
978
979
980

    found = true;
    if (memType != MEM_NULL)
    {
      memTypes.push_back(memType);
      memIndices.push_back(devIndex);
    }
  }
  if (!found)
  {
    printf("[ERROR] Unable to parse memory type token %s.  Expected one of %s followed by an index\n",
           token.c_str(), MemTypeStr);
    exit(1);
  }
}

void ParseExeType(std::string const& token, int const numCpus, int const numGpus,
                  ExeType &exeType, int& exeIndex)
{
  char typeChar;
  if (sscanf(token.c_str(), " %c%d", &typeChar, &exeIndex) != 2)
  {
    printf("[ERROR] Unable to parse valid executor token (%s).  Exepected one of %s followed by an index\n",
           token.c_str(), ExeTypeStr);
    exit(1);
  }
  exeType = CharToExeType(typeChar);

  if (IsCpuType(exeType) && (exeIndex < 0 || exeIndex >= numCpus))
  {
    printf("[ERROR] CPU index must be between 0 and %d (instead of %d)\n", numCpus-1, exeIndex);
    exit(1);
  }
  if (IsGpuType(exeType) && (exeIndex < 0 || exeIndex >= numGpus))
  {
    printf("[ERROR] GPU index must be between 0 and %d (instead of %d)\n", numGpus-1, exeIndex);
Gilbert Lee's avatar
Gilbert Lee committed
981
982
983
984
    exit(1);
  }
}

Gilbert Lee's avatar
Gilbert Lee committed
985
// Helper function to parse a list of Transfer definitions
Gilbert Lee's avatar
Gilbert Lee committed
986
void ParseTransfers(char* line, int numCpus, int numGpus, std::vector<Transfer>& transfers)
Gilbert Lee's avatar
Gilbert Lee committed
987
988
989
990
991
{
  // Replace any round brackets or '->' with spaces,
  for (int i = 1; line[i]; i++)
    if (line[i] == '(' || line[i] == ')' || line[i] == '-' || line[i] == '>' ) line[i] = ' ';

Gilbert Lee's avatar
Gilbert Lee committed
992
  transfers.clear();
Gilbert Lee's avatar
Gilbert Lee committed
993

Gilbert Lee's avatar
Gilbert Lee committed
994
  int numTransfers = 0;
Gilbert Lee's avatar
Gilbert Lee committed
995
  std::istringstream iss(line);
Gilbert Lee's avatar
Gilbert Lee committed
996
  iss >> numTransfers;
Gilbert Lee's avatar
Gilbert Lee committed
997
998
999
1000
1001
  if (iss.fail()) return;

  std::string exeMem;
  std::string srcMem;
  std::string dstMem;
Gilbert Lee's avatar
Gilbert Lee committed
1002

gilbertlee-amd's avatar
gilbertlee-amd committed
1003
  // If numTransfers < 0, read 5-tuple (srcMem, exeMem, dstMem, #CUs, #Bytes)
Gilbert Lee's avatar
Gilbert Lee committed
1004
  // otherwise read triples (srcMem, exeMem, dstMem)
gilbertlee-amd's avatar
gilbertlee-amd committed
1005
  bool const advancedMode = (numTransfers < 0);
Gilbert Lee's avatar
Gilbert Lee committed
1006
1007
  numTransfers = abs(numTransfers);

gilbertlee-amd's avatar
gilbertlee-amd committed
1008
  int numSubExecs;
gilbertlee-amd's avatar
gilbertlee-amd committed
1009
  if (!advancedMode)
Gilbert Lee's avatar
Gilbert Lee committed
1010
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1011
1012
    iss >> numSubExecs;
    if (numSubExecs <= 0 || iss.fail())
Gilbert Lee's avatar
Gilbert Lee committed
1013
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1014
      printf("Parsing error: Number of blocks to use (%d) must be greater than 0\n", numSubExecs);
Gilbert Lee's avatar
Gilbert Lee committed
1015
1016
1017
1018
      exit(1);
    }
  }

gilbertlee-amd's avatar
gilbertlee-amd committed
1019
  size_t numBytes = 0;
Gilbert Lee's avatar
Gilbert Lee committed
1020
1021
1022
  for (int i = 0; i < numTransfers; i++)
  {
    Transfer transfer;
gilbertlee-amd's avatar
gilbertlee-amd committed
1023
    transfer.numBytes = 0;
gilbertlee-amd's avatar
gilbertlee-amd committed
1024
    transfer.numBytesActual = 0;
gilbertlee-amd's avatar
gilbertlee-amd committed
1025
    if (!advancedMode)
Gilbert Lee's avatar
Gilbert Lee committed
1026
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1027
1028
1029
1030
1031
1032
1033
1034
1035
1036
      iss >> srcMem >> exeMem >> dstMem;
      if (iss.fail())
      {
        printf("Parsing error: Unable to read valid Transfer %d (SRC EXE DST) triplet\n", i+1);
        exit(1);
      }
    }
    else
    {
      std::string numBytesToken;
gilbertlee-amd's avatar
gilbertlee-amd committed
1037
      iss >> srcMem >> exeMem >> dstMem >> numSubExecs >> numBytesToken;
gilbertlee-amd's avatar
gilbertlee-amd committed
1038
1039
1040
1041
1042
1043
1044
1045
1046
1047
1048
      if (iss.fail())
      {
        printf("Parsing error: Unable to read valid Transfer %d (SRC EXE DST #CU #Bytes) tuple\n", i+1);
        exit(1);
      }
      if (sscanf(numBytesToken.c_str(), "%lu", &numBytes) != 1)
      {
        printf("Parsing error: '%s' is not a valid expression of numBytes for Transfer %d\n", numBytesToken.c_str(), i+1);
        exit(1);
      }
      char units = numBytesToken.back();
gilbertlee-amd's avatar
gilbertlee-amd committed
1049
      switch (toupper(units))
gilbertlee-amd's avatar
gilbertlee-amd committed
1050
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
1051
1052
1053
      case 'K': numBytes *= 1024; break;
      case 'M': numBytes *= 1024*1024; break;
      case 'G': numBytes *= 1024*1024*1024; break;
gilbertlee-amd's avatar
gilbertlee-amd committed
1054
      }
Gilbert Lee's avatar
Gilbert Lee committed
1055
    }
Gilbert Lee's avatar
Gilbert Lee committed
1056

gilbertlee-amd's avatar
gilbertlee-amd committed
1057
1058
1059
1060
1061
1062
1063
1064
1065
1066
1067
1068
1069
1070
1071
1072
1073
1074
1075
    ParseMemType(srcMem, numCpus, numGpus, transfer.srcType, transfer.srcIndex);
    ParseMemType(dstMem, numCpus, numGpus, transfer.dstType, transfer.dstIndex);
    ParseExeType(exeMem, numCpus, numGpus, transfer.exeType, transfer.exeIndex);

    transfer.numSrcs = (int)transfer.srcType.size();
    transfer.numDsts = (int)transfer.dstType.size();
    if (transfer.numSrcs == 0 && transfer.numDsts == 0)
    {
      printf("[ERROR] Transfer must have at least one src or dst\n");
      exit(1);
    }

    if (transfer.exeType == EXE_GPU_DMA && (transfer.numSrcs > 1 || transfer.numDsts > 1))
    {
      printf("[ERROR] GPU DMA executor can only be used for single source / single dst Transfers\n");
      exit(1);
    }

    transfer.numSubExecs = numSubExecs;
gilbertlee-amd's avatar
gilbertlee-amd committed
1076
    transfer.numBytes = numBytes;
Gilbert Lee's avatar
Gilbert Lee committed
1077
    transfers.push_back(transfer);
Gilbert Lee's avatar
Gilbert Lee committed
1078
1079
1080
1081
1082
1083
1084
1085
1086
1087
1088
1089
1090
  }
}

void EnablePeerAccess(int const deviceId, int const peerDeviceId)
{
  int canAccess;
  HIP_CALL(hipDeviceCanAccessPeer(&canAccess, deviceId, peerDeviceId));
  if (!canAccess)
  {
    printf("[ERROR] Unable to enable peer access from GPU devices %d to %d\n", peerDeviceId, deviceId);
    exit(1);
  }
  HIP_CALL(hipSetDevice(deviceId));
Gilbert Lee's avatar
Gilbert Lee committed
1091
1092
1093
1094
1095
1096
1097
  hipError_t error = hipDeviceEnablePeerAccess(peerDeviceId, 0);
  if (error != hipSuccess && error != hipErrorPeerAccessAlreadyEnabled)
  {
    printf("[ERROR] Unable to enable peer to peer access from %d to %d (%s)\n",
           deviceId, peerDeviceId, hipGetErrorString(error));
    exit(1);
  }
Gilbert Lee's avatar
Gilbert Lee committed
1098
1099
1100
1101
1102
1103
1104
1105
1106
}

void AllocateMemory(MemType memType, int devIndex, size_t numBytes, void** memPtr)
{
  if (numBytes == 0)
  {
    printf("[ERROR] Unable to allocate 0 bytes\n");
    exit(1);
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
1107
  *memPtr = nullptr;
gilbertlee-amd's avatar
gilbertlee-amd committed
1108
  if (IsCpuType(memType))
Gilbert Lee's avatar
Gilbert Lee committed
1109
1110
  {
    // Set numa policy prior to call to hipHostMalloc
1111
    numa_set_preferred(devIndex);
Gilbert Lee's avatar
Gilbert Lee committed
1112
1113
1114
1115

    // Allocate host-pinned memory (should respect NUMA mem policy)
    if (memType == MEM_CPU_FINE)
    {
1116
1117
1118
1119
#if defined (__NVCC__)
      printf("[ERROR] Fine-grained CPU memory not supported on NVIDIA platform\n");
      exit(1);
#else
Gilbert Lee's avatar
Gilbert Lee committed
1120
      HIP_CALL(hipHostMalloc((void **)memPtr, numBytes, hipHostMallocNumaUser));
1121
#endif
Gilbert Lee's avatar
Gilbert Lee committed
1122
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
1123
    else if (memType == MEM_CPU)
Gilbert Lee's avatar
Gilbert Lee committed
1124
    {
1125
1126
1127
#if defined (__NVCC__)
      if (hipHostMalloc((void **)memPtr, numBytes, 0) != hipSuccess)
#else
1128
      if (hipHostMalloc((void **)memPtr, numBytes, hipHostMallocNumaUser | hipHostMallocNonCoherent) != hipSuccess)
1129
#endif
1130
1131
1132
1133
      {
        printf("[ERROR] Unable to allocate non-coherent host memory on NUMA node %d\n", devIndex);
        exit(1);
      }
Gilbert Lee's avatar
Gilbert Lee committed
1134
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
1135
1136
1137
1138
    else if (memType == MEM_CPU_UNPINNED)
    {
      *memPtr = numa_alloc_onnode(numBytes, devIndex);
    }
Gilbert Lee's avatar
Gilbert Lee committed
1139
1140

    // Check that the allocated pages are actually on the correct NUMA node
gilbertlee-amd's avatar
gilbertlee-amd committed
1141
1142
    memset(*memPtr, 0, numBytes);
    CheckPages((char*)*memPtr, numBytes, devIndex);
Gilbert Lee's avatar
Gilbert Lee committed
1143
1144

    // Reset to default numa mem policy
1145
    numa_set_preferred(-1);
Gilbert Lee's avatar
Gilbert Lee committed
1146
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
1147
  else if (IsGpuType(memType))
Gilbert Lee's avatar
Gilbert Lee committed
1148
  {
1149
1150
1151
1152
1153
1154
1155
1156
    if (memType == MEM_GPU)
    {
      // Allocate GPU memory on appropriate device
      HIP_CALL(hipSetDevice(devIndex));
      HIP_CALL(hipMalloc((void**)memPtr, numBytes));
    }
    else if (memType == MEM_GPU_FINE)
    {
1157
#if defined (__NVCC__)
1158
1159
      printf("[ERROR] Fine-grained GPU memory not supported on NVIDIA platform\n");
      exit(1);
1160
#else
1161
1162
      HIP_CALL(hipSetDevice(devIndex));

gilbertlee-amd's avatar
gilbertlee-amd committed
1163
1164
1165
1166
1167
1168
      // NOTE: hipDeviceMallocFinegrained will be replaced by hipDeviceMallocUncached eventually
      //       Until then, this workaround is required
      hipDeviceProp_t prop;
      HIP_CALL(hipGetDeviceProperties(&prop, 0));
      int flag = (prop.gcnArch / 10 == 94) ? 0x3 : hipDeviceMallocFinegrained;
      HIP_CALL(hipExtMallocWithFlags((void**)memPtr, numBytes, flag));
1169
#endif
1170
1171
    }
    HIP_CALL(hipMemset(*memPtr, 0, numBytes));
gilbertlee-amd's avatar
gilbertlee-amd committed
1172
    HIP_CALL(hipDeviceSynchronize());
Gilbert Lee's avatar
Gilbert Lee committed
1173
1174
1175
1176
1177
1178
1179
1180
  }
  else
  {
    printf("[ERROR] Unsupported memory type %d\n", memType);
    exit(1);
  }
}

gilbertlee-amd's avatar
gilbertlee-amd committed
1181
void DeallocateMemory(MemType memType, void* memPtr, size_t const bytes)
Gilbert Lee's avatar
Gilbert Lee committed
1182
1183
1184
{
  if (memType == MEM_CPU || memType == MEM_CPU_FINE)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1185
1186
1187
1188
1189
    if (memPtr == nullptr)
    {
      printf("[ERROR] Attempting to free null CPU pointer for %lu bytes.  Skipping hipHostFree\n", bytes);
      return;
    }
Gilbert Lee's avatar
Gilbert Lee committed
1190
1191
    HIP_CALL(hipHostFree(memPtr));
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
1192
1193
  else if (memType == MEM_CPU_UNPINNED)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1194
1195
1196
1197
1198
    if (memPtr == nullptr)
    {
      printf("[ERROR] Attempting to free null unpinned CPU pointer for %lu bytes.  Skipping numa_free\n", bytes);
      return;
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
1199
1200
    numa_free(memPtr, bytes);
  }
Gilbert Lee's avatar
Gilbert Lee committed
1201
1202
  else if (memType == MEM_GPU || memType == MEM_GPU_FINE)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1203
1204
1205
1206
1207
    if (memPtr == nullptr)
    {
      printf("[ERROR] Attempting to free null GPU pointer for %lu bytes. Skipping hipFree\n", bytes);
      return;
    }
Gilbert Lee's avatar
Gilbert Lee committed
1208
1209
1210
1211
1212
1213
1214
1215
1216
1217
1218
1219
1220
1221
1222
1223
1224
1225
1226
1227
1228
1229
1230
1231
1232
1233
1234
1235
1236
1237
1238
1239
1240
1241
1242
1243
1244
1245
1246
1247
1248
1249
    HIP_CALL(hipFree(memPtr));
  }
}

void CheckPages(char* array, size_t numBytes, int targetId)
{
  unsigned long const pageSize = getpagesize();
  unsigned long const numPages = (numBytes + pageSize - 1) / pageSize;

  std::vector<void *> pages(numPages);
  std::vector<int> status(numPages);

  pages[0] = array;
  for (int i = 1; i < numPages; i++)
  {
    pages[i] = (char*)pages[i-1] + pageSize;
  }

  long const retCode = move_pages(0, numPages, pages.data(), NULL, status.data(), 0);
  if (retCode)
  {
    printf("[ERROR] Unable to collect page info\n");
    exit(1);
  }

  size_t mistakeCount = 0;
  for (int i = 0; i < numPages; i++)
  {
    if (status[i] < 0)
    {
      printf("[ERROR] Unexpected page status %d for page %d\n", status[i], i);
      exit(1);
    }
    if (status[i] != targetId) mistakeCount++;
  }
  if (mistakeCount > 0)
  {
    printf("[ERROR] %lu out of %lu pages for memory allocation were not on NUMA node %d\n", mistakeCount, numPages, targetId);
    exit(1);
  }
}

1250
1251
1252
1253
1254
1255
1256
1257
1258
uint32_t GetId(uint32_t hwId)
{
  // Based on instinct-mi200-cdna2-instruction-set-architecture.pdf
  int const shId = (hwId >> 12) & 1;
  int const cuId = (hwId >>  8) & 7;
  int const seId = (hwId >> 13) & 3;
  return (shId << 5) + (cuId << 2) + seId;
}

1259
void RunTransfer(EnvVars const& ev, int const iteration,
Gilbert Lee's avatar
Gilbert Lee committed
1260
                 ExecutorInfo& exeInfo, int const transferIdx)
Gilbert Lee's avatar
Gilbert Lee committed
1261
{
gilbertlee-amd's avatar
gilbertlee-amd committed
1262
  Transfer* transfer = exeInfo.transfers[transferIdx];
Gilbert Lee's avatar
Gilbert Lee committed
1263

gilbertlee-amd's avatar
gilbertlee-amd committed
1264
  if (transfer->exeType == EXE_GPU_GFX)
Gilbert Lee's avatar
Gilbert Lee committed
1265
1266
  {
    // Switch to executing GPU
gilbertlee-amd's avatar
gilbertlee-amd committed
1267
    int const exeIndex = RemappedIndex(transfer->exeIndex, false);
Gilbert Lee's avatar
Gilbert Lee committed
1268
1269
    HIP_CALL(hipSetDevice(exeIndex));

Gilbert Lee's avatar
Gilbert Lee committed
1270
1271
1272
    hipStream_t& stream     = exeInfo.streams[transferIdx];
    hipEvent_t&  startEvent = exeInfo.startEvents[transferIdx];
    hipEvent_t&  stopEvent  = exeInfo.stopEvents[transferIdx];
Gilbert Lee's avatar
Gilbert Lee committed
1273

gilbertlee-amd's avatar
gilbertlee-amd committed
1274
1275
1276
1277
    // Figure out how many threadblocks to use.
    // In single stream mode, all the threadblocks for this GPU are launched
    // Otherwise, just launch the threadblocks associated with this single Transfer
    int const numBlocksToRun = ev.useSingleStream ? exeInfo.totalSubExecs : transfer->numSubExecs;
1278
1279
1280
1281
1282
#if defined(__NVCC__)
    HIP_CALL(hipEventRecord(startEvent, stream));
    GpuKernelTable[ev.gpuKernel]<<<numBlocksToRun, BLOCKSIZE, ev.sharedMemBytes, stream>>>(transfer->subExecParamGpuPtr);
    HIP_CALL(hipEventRecord(stopEvent, stream));
#else
gilbertlee-amd's avatar
gilbertlee-amd committed
1283
1284
1285
1286
1287
1288
    hipExtLaunchKernelGGL(GpuKernelTable[ev.gpuKernel],
                          dim3(numBlocksToRun, 1, 1),
                          dim3(BLOCKSIZE, 1, 1),
                          ev.sharedMemBytes, stream,
                          startEvent, stopEvent,
                          0, transfer->subExecParamGpuPtr);
1289
#endif
Gilbert Lee's avatar
Gilbert Lee committed
1290
1291
    // Synchronize per iteration, unless in single sync mode, in which case
    // synchronize during last warmup / last actual iteration
Gilbert Lee's avatar
Gilbert Lee committed
1292
    HIP_CALL(hipStreamSynchronize(stream));
Gilbert Lee's avatar
Gilbert Lee committed
1293
1294
1295
1296

    if (iteration >= 0)
    {
      // Record GPU timing
Gilbert Lee's avatar
Gilbert Lee committed
1297
1298
      float gpuDeltaMsec;
      HIP_CALL(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent));
Gilbert Lee's avatar
Gilbert Lee committed
1299

Gilbert Lee's avatar
Gilbert Lee committed
1300
1301
      if (ev.useSingleStream)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
1302
        // Figure out individual timings for Transfers that were all launched together
gilbertlee-amd's avatar
gilbertlee-amd committed
1303
        for (Transfer* currTransfer : exeInfo.transfers)
Gilbert Lee's avatar
Gilbert Lee committed
1304
        {
gilbertlee-amd's avatar
gilbertlee-amd committed
1305
1306
1307
          long long minStartCycle = currTransfer->subExecParamGpuPtr[0].startCycle;
          long long maxStopCycle  = currTransfer->subExecParamGpuPtr[0].stopCycle;
          for (int i = 1; i < currTransfer->numSubExecs; i++)
Gilbert Lee's avatar
Gilbert Lee committed
1308
          {
gilbertlee-amd's avatar
gilbertlee-amd committed
1309
1310
            minStartCycle = std::min(minStartCycle, currTransfer->subExecParamGpuPtr[i].startCycle);
            maxStopCycle  = std::max(maxStopCycle,  currTransfer->subExecParamGpuPtr[i].stopCycle);
Gilbert Lee's avatar
Gilbert Lee committed
1311
          }
Gilbert Lee's avatar
Gilbert Lee committed
1312
1313
          int const wallClockRate = GetWallClockRate(exeIndex);
          double iterationTimeMs = (maxStopCycle - minStartCycle) / (double)(wallClockRate);
gilbertlee-amd's avatar
gilbertlee-amd committed
1314
          currTransfer->transferTime += iterationTimeMs;
1315
          if (ev.showIterations)
1316
          {
1317
            currTransfer->perIterationTime.push_back(iterationTimeMs);
1318
1319
1320
1321
1322
            std::set<int> CUs;
            for (int i = 0; i < currTransfer->numSubExecs; i++)
              CUs.insert(GetId(currTransfer->subExecParamGpuPtr[i].hwId));
            currTransfer->perIterationCUs.push_back(CUs);
          }
Gilbert Lee's avatar
Gilbert Lee committed
1323
        }
Gilbert Lee's avatar
Gilbert Lee committed
1324
1325
1326
1327
        exeInfo.totalTime += gpuDeltaMsec;
      }
      else
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
1328
        transfer->transferTime += gpuDeltaMsec;
1329
        if (ev.showIterations)
1330
        {
1331
          transfer->perIterationTime.push_back(gpuDeltaMsec);
1332
1333
1334
1335
1336
          std::set<int> CUs;
          for (int i = 0; i < transfer->numSubExecs; i++)
            CUs.insert(GetId(transfer->subExecParamGpuPtr[i].hwId));
          transfer->perIterationCUs.push_back(CUs);
        }
Gilbert Lee's avatar
Gilbert Lee committed
1337
1338
1339
      }
    }
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
1340
1341
1342
1343
1344
1345
1346
1347
1348
1349
1350
1351
1352
1353
1354
1355
1356
1357
1358
1359
1360
1361
1362
1363
1364
1365
1366
1367
1368
1369
1370
  else if (transfer->exeType == EXE_GPU_DMA)
  {
    // Switch to executing GPU
    int const exeIndex = RemappedIndex(transfer->exeIndex, false);
    HIP_CALL(hipSetDevice(exeIndex));

    hipStream_t& stream     = exeInfo.streams[transferIdx];
    hipEvent_t&  startEvent = exeInfo.startEvents[transferIdx];
    hipEvent_t&  stopEvent  = exeInfo.stopEvents[transferIdx];

    HIP_CALL(hipEventRecord(startEvent, stream));
    if (transfer->numSrcs == 0 && transfer->numDsts == 1)
    {
      HIP_CALL(hipMemsetAsync(transfer->dstMem[0],
                              MEMSET_CHAR, transfer->numBytesActual, stream));
    }
    else if (transfer->numSrcs == 1 && transfer->numDsts == 1)
    {
      HIP_CALL(hipMemcpyAsync(transfer->dstMem[0], transfer->srcMem[0],
                              transfer->numBytesActual, hipMemcpyDefault,
                              stream));
    }
    HIP_CALL(hipEventRecord(stopEvent, stream));
    HIP_CALL(hipStreamSynchronize(stream));

    if (iteration >= 0)
    {
      // Record GPU timing
      float gpuDeltaMsec;
      HIP_CALL(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent));
      transfer->transferTime += gpuDeltaMsec;
1371
1372
      if (ev.showIterations)
        transfer->perIterationTime.push_back(gpuDeltaMsec);
gilbertlee-amd's avatar
gilbertlee-amd committed
1373
1374
1375
    }
  }
  else if (transfer->exeType == EXE_CPU) // CPU execution agent
Gilbert Lee's avatar
Gilbert Lee committed
1376
1377
  {
    // Force this thread and all child threads onto correct NUMA node
gilbertlee-amd's avatar
gilbertlee-amd committed
1378
    int const exeIndex = RemappedIndex(transfer->exeIndex, true);
1379
    if (numa_run_on_node(exeIndex))
Gilbert Lee's avatar
Gilbert Lee committed
1380
    {
1381
      printf("[ERROR] Unable to set CPU to NUMA node %d\n", exeIndex);
Gilbert Lee's avatar
Gilbert Lee committed
1382
1383
1384
1385
1386
1387
1388
      exit(1);
    }

    std::vector<std::thread> childThreads;

    auto cpuStart = std::chrono::high_resolution_clock::now();

gilbertlee-amd's avatar
gilbertlee-amd committed
1389
1390
1391
    // Launch each subExecutor in child-threads to perform memcopies
    for (int i = 0; i < transfer->numSubExecs; ++i)
      childThreads.push_back(std::thread(CpuReduceKernel, std::ref(transfer->subExecParam[i])));
Gilbert Lee's avatar
Gilbert Lee committed
1392
1393

    // Wait for child-threads to finish
gilbertlee-amd's avatar
gilbertlee-amd committed
1394
    for (int i = 0; i < transfer->numSubExecs; ++i)
Gilbert Lee's avatar
Gilbert Lee committed
1395
1396
1397
1398
1399
1400
      childThreads[i].join();

    auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;

    // Record time if not a warmup iteration
    if (iteration >= 0)
1401
1402
1403
1404
1405
1406
    {
      double const delta = (std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() * 1000.0);
      transfer->transferTime += delta;
      if (ev.showIterations)
        transfer->perIterationTime.push_back(delta);
    }
Gilbert Lee's avatar
Gilbert Lee committed
1407
1408
1409
  }
}

gilbertlee-amd's avatar
gilbertlee-amd committed
1410
void RunPeerToPeerBenchmarks(EnvVars const& ev, size_t N)
Gilbert Lee's avatar
Gilbert Lee committed
1411
{
gilbertlee-amd's avatar
gilbertlee-amd committed
1412
1413
  ev.DisplayP2PBenchmarkEnvVars();

1414
1415
1416
  char const separator = ev.outputToCsv ? ',' : ' ';
  printf("Bytes Per Direction%c%lu\n", separator, N * sizeof(float));

Gilbert Lee's avatar
Gilbert Lee committed
1417
  // Collect the number of available CPUs/GPUs on this machine
gilbertlee-amd's avatar
gilbertlee-amd committed
1418
1419
  int const numCpus    = ev.numCpuDevices;
  int const numGpus    = ev.numGpuDevices;
Gilbert Lee's avatar
Gilbert Lee committed
1420
1421
1422
1423
1424
1425
1426
1427
1428
1429
  int const numDevices = numCpus + numGpus;

  // Enable peer to peer for each GPU
  for (int i = 0; i < numGpus; i++)
    for (int j = 0; j < numGpus; j++)
      if (i != j) EnablePeerAccess(i, j);

  // Perform unidirectional / bidirectional
  for (int isBidirectional = 0; isBidirectional <= 1; isBidirectional++)
  {
1430
1431
1432
    if (ev.p2pMode == 1 && isBidirectional == 1 ||
        ev.p2pMode == 2 && isBidirectional == 0) continue;

1433
1434
1435
1436
1437
    printf("%sdirectional copy peak bandwidth GB/s [%s read / %s write] (GPU-Executor: %s)\n", isBidirectional ? "Bi" : "Uni",
           ev.useRemoteRead ? "Remote" : "Local",
           ev.useRemoteRead ? "Local" : "Remote",
           ev.useDmaCopy    ? "DMA"   : "GFX");

Gilbert Lee's avatar
Gilbert Lee committed
1438
    // Print header
1439
    if (isBidirectional)
Gilbert Lee's avatar
Gilbert Lee committed
1440
    {
1441
1442
1443
1444
1445
1446
      printf("%12s", "SRC\\DST");
    }
    else
    {
      if (ev.useRemoteRead)
        printf("%12s", "SRC\\EXE+DST");
1447
      else
1448
1449
1450
1451
1452
1453
1454
1455
1456
1457
1458
1459
        printf("%12s", "SRC+EXE\\DST");
    }
    if (ev.outputToCsv) printf(",");
    for (int i = 0; i < numCpus; i++)
    {
      printf("%7s %02d", "CPU", i);
      if (ev.outputToCsv) printf(",");
    }
    for (int i = 0; i < numGpus; i++)
    {
      printf("%7s %02d", "GPU", i);
      if (ev.outputToCsv) printf(",");
Gilbert Lee's avatar
Gilbert Lee committed
1460
    }
1461
1462
1463
    printf("\n");

    ExeType const gpuExeType = ev.useDmaCopy ? EXE_GPU_DMA : EXE_GPU_GFX;
Gilbert Lee's avatar
Gilbert Lee committed
1464
1465
1466
    // Loop over all possible src/dst pairs
    for (int src = 0; src < numDevices; src++)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1467
1468
1469
      MemType const srcType  = (src < numCpus ? MEM_CPU : MEM_GPU);
      int     const srcIndex = (srcType == MEM_CPU ? src : src - numCpus);

1470
1471
1472
1473
      std::vector<std::vector<double>> avgBandwidth(isBidirectional + 1);
      std::vector<std::vector<double>> minBandwidth(isBidirectional + 1);
      std::vector<std::vector<double>> maxBandwidth(isBidirectional + 1);
      std::vector<std::vector<double>> stdDev(isBidirectional + 1);
gilbertlee-amd's avatar
gilbertlee-amd committed
1474

Gilbert Lee's avatar
Gilbert Lee committed
1475
1476
      for (int dst = 0; dst < numDevices; dst++)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
1477
1478
1479
        MemType const dstType  = (dst < numCpus ? MEM_CPU : MEM_GPU);
        int     const dstIndex = (dstType == MEM_CPU ? dst : dst - numCpus);

1480
1481
1482
1483
1484
1485
1486
1487
1488
1489
1490
1491
1492
1493
1494
1495
1496
1497
1498
1499
1500
1501
1502
1503
1504
1505
1506
1507
1508
1509
1510
1511
1512
1513
1514
1515
1516
1517
1518
1519
1520
1521
1522
1523
1524
1525
1526
1527
1528
1529
1530
1531
1532
1533
1534
1535
1536
1537
1538
1539
1540
1541
1542
1543
1544
1545
1546
1547
1548
1549
1550
1551
1552
1553
1554
1555
1556
1557
1558
1559
1560
1561
1562
1563
        // Prepare Transfers
        std::vector<Transfer> transfers(isBidirectional + 1);

        // SRC -> DST
        transfers[0].numBytes = N * sizeof(float);
        transfers[0].srcType.push_back(srcType);
        transfers[0].dstType.push_back(dstType);
        transfers[0].srcIndex.push_back(srcIndex);
        transfers[0].dstIndex.push_back(dstIndex);
        transfers[0].numSrcs = transfers[0].numDsts = 1;
        transfers[0].exeType = IsGpuType(ev.useRemoteRead ? dstType : srcType) ? gpuExeType : EXE_CPU;
        transfers[0].exeIndex = (ev.useRemoteRead ? dstIndex : srcIndex);
        transfers[0].numSubExecs = IsGpuType(transfers[0].exeType) ? ev.numGpuSubExecs : ev.numCpuSubExecs;

        // DST -> SRC
        if (isBidirectional)
        {
          transfers[1].numBytes = N * sizeof(float);
          transfers[1].numSrcs = transfers[1].numDsts = 1;
          transfers[1].srcType.push_back(dstType);
          transfers[1].dstType.push_back(srcType);
          transfers[1].srcIndex.push_back(dstIndex);
          transfers[1].dstIndex.push_back(srcIndex);
          transfers[1].exeType = IsGpuType(ev.useRemoteRead ? srcType : dstType) ? gpuExeType : EXE_CPU;
          transfers[1].exeIndex = (ev.useRemoteRead ? srcIndex : dstIndex);
          transfers[1].numSubExecs = IsGpuType(transfers[1].exeType) ? ev.numGpuSubExecs : ev.numCpuSubExecs;
        }

        bool skipTest = false;

        // Abort if executing on NUMA node with no CPUs
        for (int i = 0; i <= isBidirectional; i++)
        {
          if (transfers[i].exeType == EXE_CPU && ev.numCpusPerNuma[transfers[i].exeIndex] == 0)
          {
            skipTest = true;
            break;
          }

#if defined(__NVCC__)
          // NVIDIA platform cannot access GPU memory directly from CPU executors
          if (transfers[i].exeType == EXE_CPU && (IsGpuType(srcType) || IsGpuType(dstType)))
          {
            skipTest = true;
            break;
          }
#endif
        }

        if (isBidirectional && srcType == dstType && srcIndex == dstIndex) skipTest = true;

        if (!skipTest)
        {
          ExecuteTransfers(ev, 0, N, transfers, false);

          for (int dir = 0; dir <= isBidirectional; dir++)
          {
            double const avgTime = transfers[dir].transferTime / ev.numIterations;
            double const avgBw   = (transfers[dir].numBytesActual / 1.0E9) / avgTime * 1000.0f;
            avgBandwidth[dir].push_back(avgBw);

            if (ev.showIterations)
            {
              double minTime = transfers[dir].perIterationTime[0];
              double maxTime = transfers[dir].perIterationTime[0];
              double varSum  = 0;
              for (int i = 0; i < transfers[dir].perIterationTime.size(); i++)
              {
                minTime = std::min(minTime, transfers[dir].perIterationTime[i]);
                maxTime = std::max(maxTime, transfers[dir].perIterationTime[i]);
                double const bw  = (transfers[dir].numBytesActual / 1.0E9) / transfers[dir].perIterationTime[i] * 1000.0f;
                double const delta = (avgBw - bw);
                varSum += delta * delta;
              }
              double const minBw = (transfers[dir].numBytesActual / 1.0E9) / maxTime * 1000.0f;
              double const maxBw = (transfers[dir].numBytesActual / 1.0E9) / minTime * 1000.0f;
              double const stdev = sqrt(varSum / transfers[dir].perIterationTime.size());
              minBandwidth[dir].push_back(minBw);
              maxBandwidth[dir].push_back(maxBw);
              stdDev[dir].push_back(stdev);
            }
          }
        }
        else
Gilbert Lee's avatar
Gilbert Lee committed
1564
        {
1565
1566
1567
1568
1569
1570
1571
1572
1573
1574
1575
1576
1577
1578
1579
1580
1581
1582
1583
1584
          for (int dir = 0; dir <= isBidirectional; dir++)
          {
            avgBandwidth[dir].push_back(0);
            minBandwidth[dir].push_back(0);
            maxBandwidth[dir].push_back(0);
            stdDev[dir].push_back(-1.0);
          }
        }
      }

      for (int dir = 0; dir <= isBidirectional; dir++)
      {
        printf("%5s %02d %3s", (srcType == MEM_CPU) ? "CPU" : "GPU", srcIndex, dir ? "<- " : " ->");
        if (ev.outputToCsv) printf(",");

        for (int dst = 0; dst < numDevices; dst++)
        {
          double const avgBw = avgBandwidth[dir][dst];

          if (avgBw == 0.0)
Gilbert Lee's avatar
Gilbert Lee committed
1585
1586
            printf("%10s", "N/A");
          else
1587
1588
            printf("%10.2f", avgBw);
          if (ev.outputToCsv) printf(",");
Gilbert Lee's avatar
Gilbert Lee committed
1589
        }
1590
1591
1592
        printf("\n");

        if (ev.showIterations)
Gilbert Lee's avatar
Gilbert Lee committed
1593
        {
1594
1595
1596
1597
1598
1599
1600
1601
1602
1603
1604
1605
1606
1607
1608
1609
1610
1611
1612
1613
1614
1615
1616
1617
1618
1619
1620
1621
1622
1623
1624
1625
1626
1627
1628
1629
1630
1631
1632
1633
1634
          // minBw
          printf("%5s %02d %3s", (srcType == MEM_CPU) ? "CPU" : "GPU", srcIndex, "min");
          if (ev.outputToCsv) printf(",");
          for (int i = 0; i < numDevices; i++)
          {
            double const minBw = minBandwidth[dir][i];
            if (minBw == 0.0)
              printf("%10s", "N/A");
            else
              printf("%10.2f", minBw);
            if (ev.outputToCsv) printf(",");
          }
          printf("\n");

          // maxBw
          printf("%5s %02d %3s", (srcType == MEM_CPU) ? "CPU" : "GPU", srcIndex, "max");
          if (ev.outputToCsv) printf(",");
          for (int i = 0; i < numDevices; i++)
          {
            double const maxBw = maxBandwidth[dir][i];
            if (maxBw == 0.0)
              printf("%10s", "N/A");
            else
              printf("%10.2f", maxBw);
            if (ev.outputToCsv) printf(",");
          }
          printf("\n");

          // stddev
          printf("%5s %02d %3s", (srcType == MEM_CPU) ? "CPU" : "GPU", srcIndex, " sd");
          if (ev.outputToCsv) printf(",");
          for (int i = 0; i < numDevices; i++)
          {
            double const sd = stdDev[dir][i];
            if (sd == -1.0)
              printf("%10s", "N/A");
            else
              printf("%10.2f", sd);
            if (ev.outputToCsv) printf(",");
          }
          printf("\n");
Gilbert Lee's avatar
Gilbert Lee committed
1635
1636
1637
        }
        fflush(stdout);
      }
1638
1639
1640
1641
1642
1643
1644
1645
1646
1647
1648
1649
1650
1651
1652
1653

      if (isBidirectional)
      {
        printf("%5s %02d %3s", (srcType == MEM_CPU) ? "CPU" : "GPU", srcIndex, "<->");
        if (ev.outputToCsv) printf(",");
        for (int dst = 0; dst < numDevices; dst++)
        {
          double const sumBw = avgBandwidth[0][dst] + avgBandwidth[1][dst];
          if (sumBw == 0.0)
            printf("%10s", "N/A");
          else
            printf("%10.2f", sumBw);
          if (ev.outputToCsv) printf(",");
        }
        if (src < numDevices - 1) printf("\n\n");
      }
Gilbert Lee's avatar
Gilbert Lee committed
1654
    }
1655
    printf("\n");
Gilbert Lee's avatar
Gilbert Lee committed
1656
1657
1658
  }
}

1659
1660
1661
1662
1663
1664
1665
1666
1667
1668
1669
1670
1671
1672
1673
1674
1675
1676
1677
1678
1679
1680
1681
1682
1683
1684
1685
1686
1687
1688
1689
1690
1691
1692
1693
1694
1695
1696
1697
1698
1699
1700
1701
1702
1703
1704
1705
1706
1707
1708
1709
1710
1711
1712
1713
1714
1715
1716
1717
1718
1719
1720
1721
1722
1723
1724
1725
1726
1727
1728
void RunScalingBenchmark(EnvVars const& ev, size_t N, int const exeIndex, int const maxSubExecs)
{
  ev.DisplayEnvVars();

  // Collect the number of available CPUs/GPUs on this machine
  int const numCpus    = ev.numCpuDevices;
  int const numGpus    = ev.numGpuDevices;
  int const numDevices = numCpus + numGpus;

  // Enable peer to peer for each GPU
  for (int i = 0; i < numGpus; i++)
    for (int j = 0; j < numGpus; j++)
      if (i != j) EnablePeerAccess(i, j);

  char separator = (ev.outputToCsv ? ',' : ' ');

  std::vector<Transfer> transfers(1);
  transfers[0].numBytes = N * sizeof(float);
  transfers[0].numSrcs  = 1;
  transfers[0].numDsts  = 1;
  transfers[0].exeType  = EXE_GPU_GFX;
  transfers[0].exeIndex = exeIndex;
  transfers[0].srcType.resize(1, MEM_GPU);
  transfers[0].dstType.resize(1, MEM_GPU);
  transfers[0].srcIndex.resize(1);
  transfers[0].dstIndex.resize(1);

  printf("GPU-GFX Scaling benchmark:\n");
  printf("==========================\n");
  printf("- Copying %lu bytes from GPU %d to other devices\n", transfers[0].numBytes, exeIndex);
  printf("- All numbers reported as GB/sec\n\n");

  printf("NumCUs");
  for (int i = 0; i < numDevices; i++)
    printf("%c  %s%02d     ", separator, i < numCpus ? "CPU" : "GPU", i < numCpus ? i : i - numCpus);
  printf("\n");

  std::vector<std::pair<double, int>> bestResult(numDevices);
  for (int numSubExec = 1; numSubExec <= maxSubExecs; numSubExec++)
  {
    transfers[0].numSubExecs = numSubExec;
    printf("%4d  ", numSubExec);

    for (int i = 0; i < numDevices; i++)
    {
      transfers[0].dstType[0]  = i < numCpus ? MEM_CPU : MEM_GPU;
      transfers[0].dstIndex[0] = i < numCpus ? i : i - numCpus;

      ExecuteTransfers(ev, 0, N, transfers, false);
      double transferDurationMsec = transfers[0].transferTime / (1.0 * ev.numIterations);
      double transferBandwidthGbs = (transfers[0].numBytesActual / 1.0E9) / transferDurationMsec * 1000.0f;
      printf("%c%7.2f     ", separator, transferBandwidthGbs);

      if (transferBandwidthGbs > bestResult[i].first)
      {
        bestResult[i].first  = transferBandwidthGbs;
        bestResult[i].second = numSubExec;
      }
    }
    printf("\n");
  }

  printf(" Best ");
  for (int i = 0; i < numDevices; i++)
  {
    printf("%c%7.2f(%3d)", separator, bestResult[i].first, bestResult[i].second);
  }
  printf("\n");
}

gilbertlee-amd's avatar
gilbertlee-amd committed
1729
1730
1731
1732
1733
1734
1735
1736
1737
1738
1739
1740
1741
1742
1743
1744
1745
1746
1747
1748
1749
1750
1751
1752
1753
1754
1755
1756
1757
1758
1759
1760
1761
1762
1763
1764
1765
1766
1767
1768
1769
1770
1771
1772
1773
1774
1775
1776
1777
1778
1779
1780
1781
1782
1783
1784
1785
1786
1787
1788
1789
1790
1791
1792
1793
1794
1795
void RunAllToAllBenchmark(EnvVars const& ev, size_t const numBytesPerTransfer, int const numSubExecs)
{
  ev.DisplayEnvVars();

  // Collect the number of GPU devices to use
  int const numGpus = ev.numGpuDevices;

  // Enable peer to peer for each GPU
  for (int i = 0; i < numGpus; i++)
    for (int j = 0; j < numGpus; j++)
      if (i != j) EnablePeerAccess(i, j);

  char separator = (ev.outputToCsv ? ',' : ' ');

  Transfer transfer;
  transfer.numBytes    = numBytesPerTransfer;
  transfer.numSubExecs = numSubExecs;
  transfer.numSrcs     = 1;
  transfer.numDsts     = 1;
  transfer.exeType     = EXE_GPU_GFX;
  transfer.srcType.resize(1, MEM_GPU);
  transfer.dstType.resize(1, MEM_GPU);
  transfer.srcIndex.resize(1);
  transfer.dstIndex.resize(1);

  std::vector<Transfer> transfers;
  for (int i = 0; i < numGpus; i++)
  {
    transfer.srcIndex[0] = i;
    transfer.exeIndex    = i;
    for (int j = 0; j < numGpus; j++)
    {
      transfer.dstIndex[0] = j;
      transfers.push_back(transfer);
    }
  }

  printf("GPU-GFX All-To-All benchmark:\n");
  printf("==========================\n");
  printf("- Copying %lu bytes between every pair of GPUs using %d CUs\n", numBytesPerTransfer, numSubExecs);
  printf("- All numbers reported as GB/sec\n\n");

  double totalBandwidthCpu = 0;
  ExecuteTransfers(ev, 0, numBytesPerTransfer / sizeof(float), transfers, true, &totalBandwidthCpu);

  printf("\nSummary:\n");
  printf("==========================================================\n");
  printf("SRC\\DST");
  for (int dst = 0; dst < numGpus; dst++)
    printf("%cGPU %02d   ", separator, dst);
  printf("\n");

  for (int src = 0; src < numGpus; src++)
  {
    printf("GPU %02d", src);
    for (int dst = 0; dst < numGpus; dst++)
    {
      Transfer const& transfer = transfers[src * numGpus + dst];
      double transferDurationMsec = transfer.transferTime / (1.0 * ev.numIterations);
      double transferBandwidthGbs = (transfer.numBytesActual / 1.0E9) / transferDurationMsec * 1000.0f;
      printf("%c%7.2f  ", separator, transferBandwidthGbs);
    }
    printf("\n");
  }
  printf("Aggregate bandwidth (CPU Timed): %7.2f\n", totalBandwidthCpu);
}

gilbertlee-amd's avatar
gilbertlee-amd committed
1796
void Transfer::PrepareSubExecParams(EnvVars const& ev)
Gilbert Lee's avatar
Gilbert Lee committed
1797
{
gilbertlee-amd's avatar
gilbertlee-amd committed
1798
1799
1800
1801
1802
1803
1804
  // Each subExecutor needs to know src/dst pointers and how many elements to transfer
  // Figure out the sub-array each subExecutor works on for this Transfer
  // - Partition N as evenly as possible, but try to keep subarray sizes as multiples of BLOCK_BYTES bytes,
  //   except the very last one, for alignment reasons
  size_t const N              = this->numBytesActual / sizeof(float);
  int    const initOffset     = ev.byteOffset / sizeof(float);
  int    const targetMultiple = ev.blockBytes / sizeof(float);
Gilbert Lee's avatar
Gilbert Lee committed
1805

gilbertlee-amd's avatar
gilbertlee-amd committed
1806
  // In some cases, there may not be enough data for all subExectors
1807
  int const maxSubExecToUse = std::min((size_t)(N + targetMultiple - 1) / targetMultiple, (size_t)this->numSubExecs);
gilbertlee-amd's avatar
gilbertlee-amd committed
1808
1809
  this->subExecParam.clear();
  this->subExecParam.resize(this->numSubExecs);
Gilbert Lee's avatar
Gilbert Lee committed
1810
1811

  size_t assigned = 0;
gilbertlee-amd's avatar
gilbertlee-amd committed
1812
1813
1814
1815
1816
1817
1818
1819
1820
1821
1822
1823
1824
1825
1826
1827
1828
1829
1830
1831
  for (int i = 0; i < this->numSubExecs; ++i)
  {
    int    const subExecLeft = std::max(0, maxSubExecToUse - i);
    size_t const leftover    = N - assigned;
    size_t const roundedN    = (leftover + targetMultiple - 1) / targetMultiple;

    SubExecParam& p = this->subExecParam[i];
    p.N             = subExecLeft ? std::min(leftover, ((roundedN / subExecLeft) * targetMultiple)) : 0;
    p.numSrcs       = this->numSrcs;
    p.numDsts       = this->numDsts;
    for (int iSrc = 0; iSrc < this->numSrcs; ++iSrc)
      p.src[iSrc] = this->srcMem[iSrc] + assigned + initOffset;
    for (int iDst = 0; iDst < this->numDsts; ++iDst)
      p.dst[iDst] = this->dstMem[iDst] + assigned + initOffset;

    if (ev.enableDebug)
    {
      printf("Transfer %02d SE:%02d: %10lu floats: %10lu to %10lu\n",
             this->transferIndex, i, p.N, assigned, assigned + p.N);
    }
Gilbert Lee's avatar
Gilbert Lee committed
1832

gilbertlee-amd's avatar
gilbertlee-amd committed
1833
1834
1835
    p.startCycle = 0;
    p.stopCycle  = 0;
    assigned += p.N;
Gilbert Lee's avatar
Gilbert Lee committed
1836
1837
  }

Gilbert Lee's avatar
Gilbert Lee committed
1838
  this->transferTime = 0.0;
1839
  this->perIterationTime.clear();
Gilbert Lee's avatar
Gilbert Lee committed
1840
1841
}

gilbertlee-amd's avatar
gilbertlee-amd committed
1842
1843
1844
1845
1846
1847
1848
1849
1850
1851
1852
1853
1854
1855
void Transfer::PrepareReference(EnvVars const& ev, std::vector<float>& buffer, int bufferIdx)
{
  size_t N = buffer.size();
  if (bufferIdx >= 0)
  {
    size_t patternLen = ev.fillPattern.size();
    if (patternLen > 0)
    {
      for (size_t i = 0; i < N; ++i)
        buffer[i] = ev.fillPattern[i % patternLen];
    }
    else
    {
      for (size_t i = 0; i < N; ++i)
1856
        buffer[i] = PrepSrcValue(bufferIdx, i);
gilbertlee-amd's avatar
gilbertlee-amd committed
1857
1858
1859
1860
1861
1862
1863
1864
1865
1866
1867
1868
1869
1870
1871
1872
1873
1874
1875
1876
1877
1878
1879
1880
1881
1882
1883
1884
1885
    }
  }
  else // Destination buffer
  {
    if (this->numSrcs == 0)
    {
      // Note: 0x75757575 = 13323083.0
      memset(buffer.data(), MEMSET_CHAR, N * sizeof(float));
    }
    else
    {
      PrepareReference(ev, buffer, 0);

      if (this->numSrcs > 1)
      {
        std::vector<float> temp(N);
        for (int srcIdx = 1; srcIdx < this->numSrcs; ++srcIdx)
        {
          PrepareReference(ev, temp, srcIdx);
          for (int i = 0; i < N; ++i)
          {
            buffer[i] += temp[i];
          }
        }
      }
    }
  }
}

1886
bool Transfer::PrepareSrc(EnvVars const& ev)
gilbertlee-amd's avatar
gilbertlee-amd committed
1887
{
1888
  if (this->numSrcs == 0) return true;
gilbertlee-amd's avatar
gilbertlee-amd committed
1889
1890
1891
1892
1893
1894
  size_t const N = this->numBytesActual / sizeof(float);
  int const initOffset = ev.byteOffset / sizeof(float);

  std::vector<float> reference(N);
  for (int srcIdx = 0; srcIdx < this->numSrcs; ++srcIdx)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1895
    float* srcPtr = this->srcMem[srcIdx] + initOffset;
1896
    PrepareReference(ev, reference, srcIdx);
gilbertlee-amd's avatar
gilbertlee-amd committed
1897
1898
1899

    // Initialize source memory array with reference pattern
    if (IsGpuType(this->srcType[srcIdx]))
1900
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1901
1902
1903
1904
1905
1906
      int const deviceIdx = RemappedIndex(this->srcIndex[srcIdx], false);
      HIP_CALL(hipSetDevice(deviceIdx));
      if (ev.usePrepSrcKernel)
        PrepSrcDataKernel<<<32, BLOCKSIZE>>>(srcPtr, N, srcIdx);
      else
        HIP_CALL(hipMemcpy(srcPtr, reference.data(), this->numBytesActual, hipMemcpyDefault));
1907
1908
      HIP_CALL(hipDeviceSynchronize());
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
1909
    else if (IsCpuType(this->srcType[srcIdx]))
1910
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1911
      memcpy(srcPtr, reference.data(), this->numBytesActual);
1912
    }
1913
1914

    // Perform check just to make sure that data has been copied properly
gilbertlee-amd's avatar
gilbertlee-amd committed
1915
    float* srcCheckPtr = srcPtr;
1916
    std::vector<float> srcCopy(N);
gilbertlee-amd's avatar
gilbertlee-amd committed
1917
1918
1919
1920
1921
1922
1923
1924
1925
    if (IsGpuType(this->srcType[srcIdx]))
    {
      if (!ev.validateDirect)
      {
        HIP_CALL(hipMemcpy(srcCopy.data(), srcPtr, this->numBytesActual, hipMemcpyDefault));
        HIP_CALL(hipDeviceSynchronize());
        srcCheckPtr = srcCopy.data();
      }
    }
1926
1927
1928

    for (size_t i = 0; i < N; ++i)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1929
      if (reference[i] != srcCheckPtr[i])
1930
1931
      {
        printf("\n[ERROR] Unexpected mismatch at index %lu of source array %d:\n", i, srcIdx);
1932
1933
1934
#if !defined(__NVCC__)
        float const val = this->srcMem[srcIdx][initOffset + i];
        printf("[ERROR] SRC %02d   value: %10.5f [%08X] Direct: %10.5f [%08X]\n",
gilbertlee-amd's avatar
gilbertlee-amd committed
1935
               srcIdx, srcCheckPtr[i], *(unsigned int*)&srcCheckPtr[i], val, *(unsigned int*)&val);
1936
#else
gilbertlee-amd's avatar
gilbertlee-amd committed
1937
        printf("[ERROR] SRC %02d   value: %10.5f [%08X]\n", srcIdx, srcCheckPtr[i], *(unsigned int*)&srcCheckPtr[i]);
1938
#endif
1939
1940
1941
1942
1943
1944
1945
1946
1947
        printf("[ERROR] EXPECTED value: %10.5f [%08X]\n", reference[i], *(unsigned int*)&reference[i]);
        printf("[ERROR] Failed Transfer details: #%d: %s -> [%c%d:%d] -> %s\n",
               this->transferIndex,
               this->SrcToStr().c_str(),
               ExeTypeStr[this->exeType], this->exeIndex,
               this->numSubExecs,
               this->DstToStr().c_str());
        if (!ev.continueOnError)
          exit(1);
1948
        return false;
1949
1950
      }
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
1951
  }
1952
  return true;
gilbertlee-amd's avatar
gilbertlee-amd committed
1953
1954
1955
1956
1957
1958
1959
1960
1961
1962
1963
1964
1965
1966
1967
}

void Transfer::ValidateDst(EnvVars const& ev)
{
  if (this->numDsts == 0) return;
  size_t const N = this->numBytesActual / sizeof(float);
  int const initOffset = ev.byteOffset / sizeof(float);

  std::vector<float> reference(N);
  PrepareReference(ev, reference, -1);

  std::vector<float> hostBuffer(N);
  for (int dstIdx = 0; dstIdx < this->numDsts; ++dstIdx)
  {
    float* output;
1968
    if (IsCpuType(this->dstType[dstIdx]) || ev.validateDirect)
gilbertlee-amd's avatar
gilbertlee-amd committed
1969
1970
1971
1972
1973
    {
      output = this->dstMem[dstIdx] + initOffset;
    }
    else
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1974
1975
      int const deviceIdx = RemappedIndex(this->dstIndex[dstIdx], false);
      HIP_CALL(hipSetDevice(deviceIdx));
gilbertlee-amd's avatar
gilbertlee-amd committed
1976
      HIP_CALL(hipMemcpy(hostBuffer.data(), this->dstMem[dstIdx] + initOffset, this->numBytesActual, hipMemcpyDefault));
gilbertlee-amd's avatar
gilbertlee-amd committed
1977
      HIP_CALL(hipDeviceSynchronize());
gilbertlee-amd's avatar
gilbertlee-amd committed
1978
1979
1980
1981
1982
1983
1984
      output = hostBuffer.data();
    }

    for (size_t i = 0; i < N; ++i)
    {
      if (reference[i] != output[i])
      {
1985
1986
1987
1988
1989
        printf("\n[ERROR] Unexpected mismatch at index %lu of destination array %d:\n", i, dstIdx);
        for (int srcIdx = 0; srcIdx < this->numSrcs; ++srcIdx)
        {
          float srcVal;
          HIP_CALL(hipMemcpy(&srcVal, this->srcMem[srcIdx] + initOffset + i, sizeof(float), hipMemcpyDefault));
1990
1991
1992
1993
1994
#if !defined(__NVCC__)
          float val = this->srcMem[srcIdx][initOffset + i];
          printf("[ERROR] SRC %02dD  value: %10.5f [%08X] Direct: %10.5f [%08X]\n",
                 srcIdx, srcVal, *(unsigned int*)&srcVal, val, *(unsigned int*)&val);
#else
1995
          printf("[ERROR] SRC %02d   value: %10.5f [%08X]\n", srcIdx, srcVal, *(unsigned int*)&srcVal);
1996
#endif
1997
        }
1998
        printf("[ERROR] EXPECTED value: %10.5f [%08X]\n", reference[i], *(unsigned int*)&reference[i]);
1999
2000
2001
2002
2003
#if !defined(__NVCC__)
        float dstVal = this->dstMem[dstIdx][initOffset + i];
        printf("[ERROR] DST %02d   value: %10.5f [%08X] Direct: %10.5f [%08X]\n",
               dstIdx, output[i], *(unsigned int*)&output[i], dstVal, *(unsigned int*)&dstVal);
#else
2004
        printf("[ERROR] DST %02d   value: %10.5f [%08X]\n", dstIdx, output[i], *(unsigned int*)&output[i]);
2005
#endif
gilbertlee-amd's avatar
gilbertlee-amd committed
2006
2007
2008
2009
2010
2011
        printf("[ERROR] Failed Transfer details: #%d: %s -> [%c%d:%d] -> %s\n",
               this->transferIndex,
               this->SrcToStr().c_str(),
               ExeTypeStr[this->exeType], this->exeIndex,
               this->numSubExecs,
               this->DstToStr().c_str());
2012
2013
        if (!ev.continueOnError)
          exit(1);
2014
2015
        else
          break;
gilbertlee-amd's avatar
gilbertlee-amd committed
2016
2017
2018
2019
2020
2021
2022
2023
2024
2025
2026
2027
2028
2029
2030
2031
2032
2033
2034
2035
2036
2037
2038
      }
    }
  }
}

std::string Transfer::SrcToStr() const
{
  if (numSrcs == 0) return "N";
  std::stringstream ss;
  for (int i = 0; i < numSrcs; ++i)
    ss << MemTypeStr[srcType[i]] << srcIndex[i];
  return ss.str();
}

std::string Transfer::DstToStr() const
{
  if (numDsts == 0) return "N";
  std::stringstream ss;
  for (int i = 0; i < numDsts; ++i)
    ss << MemTypeStr[dstType[i]] << dstIndex[i];
  return ss.str();
}

Gilbert Lee's avatar
Gilbert Lee committed
2039
2040
2041
2042
2043
2044
2045
2046
2047
2048
2049
2050
2051
// NOTE: This is a stop-gap solution until HIP provides wallclock values
int GetWallClockRate(int deviceId)
{
  static std::vector<int> wallClockPerDeviceMhz;

  if (wallClockPerDeviceMhz.size() == 0)
  {
    int numGpuDevices;
    HIP_CALL(hipGetDeviceCount(&numGpuDevices));
    wallClockPerDeviceMhz.resize(numGpuDevices);

    for (int i = 0; i < numGpuDevices; i++)
    {
2052
2053
2054
2055
2056
2057
#if defined(__NVCC__)
      int value = 1410000;
      //HIP_CALL(hipDeviceGetAttribute(&value, hipDeviceAttributeClockRate, i));
      //value *= 1000;
#else
      hipDeviceProp_t prop;
Gilbert Lee's avatar
Gilbert Lee committed
2058
2059
2060
2061
2062
      HIP_CALL(hipGetDeviceProperties(&prop, i));
      int value = 25000;
      switch (prop.gcnArch)
      {
      case 906: case 910: value = 25000; break;
gilbertlee-amd's avatar
gilbertlee-amd committed
2063
      case 940: case 941: case 942: value = 100000; break;
Gilbert Lee's avatar
Gilbert Lee committed
2064
2065
2066
      default:
        printf("Unrecognized GCN arch %d\n", prop.gcnArch);
      }
2067
#endif
Gilbert Lee's avatar
Gilbert Lee committed
2068
2069
2070
2071
2072
      wallClockPerDeviceMhz[i] = value;
    }
  }
  return wallClockPerDeviceMhz[deviceId];
}
Gilbert Lee's avatar
Gilbert Lee committed
2073

gilbertlee-amd's avatar
gilbertlee-amd committed
2074
void RunSweepPreset(EnvVars const& ev, size_t const numBytesPerTransfer, int const numGpuSubExecs, int const numCpuSubExecs, bool const isRandom)
Gilbert Lee's avatar
Gilbert Lee committed
2075
2076
2077
2078
{
  ev.DisplaySweepEnvVars();

  // Compute how many possible Transfers are permitted (unique SRC/EXE/DST triplets)
gilbertlee-amd's avatar
gilbertlee-amd committed
2079
  std::vector<std::pair<ExeType, int>> exeList;
Gilbert Lee's avatar
Gilbert Lee committed
2080
2081
  for (auto exe : ev.sweepExe)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
2082
2083
    ExeType const exeType = CharToExeType(exe);
    if (IsGpuType(exeType))
Gilbert Lee's avatar
Gilbert Lee committed
2084
    {
2085
      for (int exeIndex = 0; exeIndex < ev.numGpuDevices; ++exeIndex)
gilbertlee-amd's avatar
gilbertlee-amd committed
2086
        exeList.push_back(std::make_pair(exeType, exeIndex));
Gilbert Lee's avatar
Gilbert Lee committed
2087
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
2088
    else if (IsCpuType(exeType))
Gilbert Lee's avatar
Gilbert Lee committed
2089
    {
2090
2091
2092
2093
      for (int exeIndex = 0; exeIndex < ev.numCpuDevices; ++exeIndex)
      {
        // Skip NUMA nodes that have no CPUs (e.g. CXL)
        if (ev.numCpusPerNuma[exeIndex] == 0) continue;
gilbertlee-amd's avatar
gilbertlee-amd committed
2094
        exeList.push_back(std::make_pair(exeType, exeIndex));
2095
      }
Gilbert Lee's avatar
Gilbert Lee committed
2096
2097
    }
  }
2098
  int numExes = exeList.size();
Gilbert Lee's avatar
Gilbert Lee committed
2099
2100
2101
2102

  std::vector<std::pair<MemType, int>> srcList;
  for (auto src : ev.sweepSrc)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
2103
2104
    MemType const srcType = CharToMemType(src);
    int const numDevices = IsGpuType(srcType) ? ev.numGpuDevices : ev.numCpuDevices;
2105

Gilbert Lee's avatar
Gilbert Lee committed
2106
    for (int srcIndex = 0; srcIndex < numDevices; ++srcIndex)
gilbertlee-amd's avatar
gilbertlee-amd committed
2107
      srcList.push_back(std::make_pair(srcType, srcIndex));
Gilbert Lee's avatar
Gilbert Lee committed
2108
2109
2110
2111
2112
2113
2114
  }
  int numSrcs = srcList.size();


  std::vector<std::pair<MemType, int>> dstList;
  for (auto dst : ev.sweepDst)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
2115
2116
    MemType const dstType = CharToMemType(dst);
    int const numDevices = IsGpuType(dstType) ? ev.numGpuDevices : ev.numCpuDevices;
Gilbert Lee's avatar
Gilbert Lee committed
2117
2118

    for (int dstIndex = 0; dstIndex < numDevices; ++dstIndex)
gilbertlee-amd's avatar
gilbertlee-amd committed
2119
      dstList.push_back(std::make_pair(dstType, dstIndex));
Gilbert Lee's avatar
Gilbert Lee committed
2120
2121
2122
  }
  int numDsts = dstList.size();

2123
2124
  // Build array of possibilities, respecting any additional restrictions (e.g. XGMI hop count)
  struct TransferInfo
Gilbert Lee's avatar
Gilbert Lee committed
2125
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
2126
2127
2128
    MemType srcType; int srcIndex;
    ExeType exeType; int exeIndex;
    MemType dstType; int dstIndex;
2129
2130
2131
2132
2133
2134
2135
2136
  };

  // If either XGMI minimum is non-zero, or XGMI maximum is specified and non-zero then both links must be XGMI
  bool const useXgmiOnly = (ev.sweepXgmiMin > 0 || ev.sweepXgmiMax > 0);

  std::vector<TransferInfo> possibleTransfers;
  TransferInfo tinfo;
  for (int i = 0; i < numExes; ++i)
Gilbert Lee's avatar
Gilbert Lee committed
2137
  {
2138
2139
    // Skip CPU executors if XGMI link must be used
    if (useXgmiOnly && !IsGpuType(exeList[i].first)) continue;
gilbertlee-amd's avatar
gilbertlee-amd committed
2140
2141
    tinfo.exeType  = exeList[i].first;
    tinfo.exeIndex = exeList[i].second;
2142

gilbertlee-amd's avatar
gilbertlee-amd committed
2143
    bool isXgmiSrc  = false;
2144
2145
2146
2147
2148
2149
2150
    int  numHopsSrc = 0;
    for (int j = 0; j < numSrcs; ++j)
    {
      if (IsGpuType(exeList[i].first) && IsGpuType(srcList[j].first))
      {
        if (exeList[i].second != srcList[j].second)
        {
2151
2152
2153
#if defined(__NVCC__)
          isXgmiSrc = false;
#else
2154
          uint32_t exeToSrcLinkType, exeToSrcHopCount;
gilbertlee-amd's avatar
gilbertlee-amd committed
2155
2156
          HIP_CALL(hipExtGetLinkTypeAndHopCount(RemappedIndex(exeList[i].second, false),
                                                RemappedIndex(srcList[j].second, false),
2157
2158
2159
2160
                                                &exeToSrcLinkType,
                                                &exeToSrcHopCount));
          isXgmiSrc = (exeToSrcLinkType == HSA_AMD_LINK_INFO_TYPE_XGMI);
          if (isXgmiSrc) numHopsSrc = exeToSrcHopCount;
2161
#endif
2162
2163
2164
2165
2166
2167
2168
2169
2170
2171
2172
2173
2174
2175
2176
        }
        else
        {
          isXgmiSrc = true;
          numHopsSrc = 0;
        }

        // Skip this SRC if it is not XGMI but only XGMI links may be used
        if (useXgmiOnly && !isXgmiSrc) continue;

        // Skip this SRC if XGMI distance is already past limit
        if (ev.sweepXgmiMax >= 0 && isXgmiSrc && numHopsSrc > ev.sweepXgmiMax) continue;
      }
      else if (useXgmiOnly) continue;

gilbertlee-amd's avatar
gilbertlee-amd committed
2177
2178
      tinfo.srcType  = srcList[j].first;
      tinfo.srcIndex = srcList[j].second;
2179
2180
2181
2182
2183
2184
2185
2186
2187

      bool isXgmiDst = false;
      int  numHopsDst = 0;
      for (int k = 0; k < numDsts; ++k)
      {
        if (IsGpuType(exeList[i].first) && IsGpuType(dstList[k].first))
        {
          if (exeList[i].second != dstList[k].second)
          {
2188
2189
2190
#if defined(__NVCC__)
            isXgmiSrc = false;
#else
2191
            uint32_t exeToDstLinkType, exeToDstHopCount;
gilbertlee-amd's avatar
gilbertlee-amd committed
2192
2193
            HIP_CALL(hipExtGetLinkTypeAndHopCount(RemappedIndex(exeList[i].second, false),
                                                  RemappedIndex(dstList[k].second, false),
2194
2195
2196
2197
                                                  &exeToDstLinkType,
                                                  &exeToDstHopCount));
            isXgmiDst = (exeToDstLinkType == HSA_AMD_LINK_INFO_TYPE_XGMI);
            if (isXgmiDst) numHopsDst = exeToDstHopCount;
2198
#endif
2199
2200
2201
2202
2203
2204
2205
2206
2207
2208
2209
2210
2211
2212
2213
2214
2215
          }
          else
          {
            isXgmiDst = true;
            numHopsDst = 0;
          }
        }

        // Skip this DST if it is not XGMI but only XGMI links may be used
        if (useXgmiOnly && !isXgmiDst) continue;

        // Skip this DST if total XGMI distance (SRC + DST) is less than min limit
        if (ev.sweepXgmiMin > 0 && (numHopsSrc + numHopsDst < ev.sweepXgmiMin)) continue;

        // Skip this DST if total XGMI distance (SRC + DST) is greater than max limit
        if (ev.sweepXgmiMax >= 0 && (numHopsSrc + numHopsDst) > ev.sweepXgmiMax) continue;

2216
2217
2218
2219
2220
2221
#if defined(__NVCC__)
        // Skip CPU executors on GPU memory on NVIDIA platform
        if (IsCpuType(exeList[i].first) && (IsGpuType(dstList[j].first) || IsGpuType(dstList[k].first)))
          continue;
#endif

gilbertlee-amd's avatar
gilbertlee-amd committed
2222
2223
        tinfo.dstType  = dstList[k].first;
        tinfo.dstIndex = dstList[k].second;
2224
2225
2226
2227

        possibleTransfers.push_back(tinfo);
      }
    }
Gilbert Lee's avatar
Gilbert Lee committed
2228
2229
  }

2230
2231
2232
  int const numPossible = (int)possibleTransfers.size();
  int maxParallelTransfers = (ev.sweepMax == 0 ? numPossible : ev.sweepMax);

Gilbert Lee's avatar
Gilbert Lee committed
2233
2234
2235
2236
2237
2238
  if (ev.sweepMin > numPossible)
  {
    printf("No valid test configurations exist\n");
    return;
  }

2239
2240
2241
2242
2243
2244
  if (ev.outputToCsv)
  {
    printf("\nTest#,Transfer#,NumBytes,Src,Exe,Dst,CUs,BW(GB/s),Time(ms),"
           "ExeToSrcLinkType,ExeToDstLinkType,SrcAddr,DstAddr\n");
  }

Gilbert Lee's avatar
Gilbert Lee committed
2245
2246
  int numTestsRun = 0;
  int M = ev.sweepMin;
gilbertlee-amd's avatar
gilbertlee-amd committed
2247
2248
2249
2250
2251
2252
2253
2254
2255
2256
2257
  std::uniform_int_distribution<int> randSize(1, numBytesPerTransfer / sizeof(float));
  std::uniform_int_distribution<int> distribution(ev.sweepMin, maxParallelTransfers);

  // Log sweep to configuration file
  FILE *fp = fopen("lastSweep.cfg", "w");
  if (!fp)
  {
    printf("[ERROR] Unable to open lastSweep.cfg.  Check permissions\n");
    exit(1);
  }

Gilbert Lee's avatar
Gilbert Lee committed
2258
2259
2260
2261
2262
2263
2264
2265
2266
  // Create bitmask of numPossible triplets, of which M will be chosen
  std::string bitmask(M, 1);  bitmask.resize(numPossible, 0);
  auto cpuStart = std::chrono::high_resolution_clock::now();
  while (1)
  {
    if (isRandom)
    {
      // Pick random number of simultaneous transfers to execute
      // NOTE: This currently skews distribution due to some #s having more possibilities than others
gilbertlee-amd's avatar
gilbertlee-amd committed
2267
      M = distribution(*ev.generator);
Gilbert Lee's avatar
Gilbert Lee committed
2268
2269
2270
2271

      // Generate a random bitmask
      for (int i = 0; i < numPossible; i++)
        bitmask[i] = (i < M) ? 1 : 0;
2272
      std::shuffle(bitmask.begin(), bitmask.end(), *ev.generator);
Gilbert Lee's avatar
Gilbert Lee committed
2273
2274
2275
2276
2277
2278
2279
2280
2281
2282
    }

    // Convert bitmask to list of Transfers
    std::vector<Transfer> transfers;
    for (int value = 0; value < numPossible; ++value)
    {
      if (bitmask[value])
      {
        // Convert integer value to (SRC->EXE->DST) triplet
        Transfer transfer;
gilbertlee-amd's avatar
gilbertlee-amd committed
2283
2284
2285
2286
2287
        transfer.numSrcs        = 1;
        transfer.numDsts        = 1;
        transfer.srcType        = {possibleTransfers[value].srcType};
        transfer.srcIndex       = {possibleTransfers[value].srcIndex};
        transfer.exeType        = possibleTransfers[value].exeType;
2288
        transfer.exeIndex       = possibleTransfers[value].exeIndex;
gilbertlee-amd's avatar
gilbertlee-amd committed
2289
2290
2291
        transfer.dstType        = {possibleTransfers[value].dstType};
        transfer.dstIndex       = {possibleTransfers[value].dstIndex};
        transfer.numSubExecs    = IsGpuType(transfer.exeType) ? numGpuSubExecs : numCpuSubExecs;
gilbertlee-amd's avatar
gilbertlee-amd committed
2292
        transfer.numBytes       = ev.sweepRandBytes ? randSize(*ev.generator) * sizeof(float) : 0;
Gilbert Lee's avatar
Gilbert Lee committed
2293
2294
2295
2296
        transfers.push_back(transfer);
      }
    }

gilbertlee-amd's avatar
gilbertlee-amd committed
2297
2298
    LogTransfers(fp, ++numTestsRun, transfers);
    ExecuteTransfers(ev, numTestsRun, numBytesPerTransfer / sizeof(float), transfers);
Gilbert Lee's avatar
Gilbert Lee committed
2299
2300
2301
2302
2303
2304
2305
2306
2307
2308
2309
2310
2311
2312
2313
2314
2315
2316
2317
2318
2319
2320
2321
2322
2323
2324
2325
2326
2327
2328
2329

    // Check for test limit
    if (numTestsRun == ev.sweepTestLimit)
    {
      printf("Test limit reached\n");
      break;
    }

    // Check for time limit
    auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
    double totalCpuTime = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count();
    if (ev.sweepTimeLimit && totalCpuTime > ev.sweepTimeLimit)
    {
      printf("Time limit exceeded\n");
      break;
    }

    // Increment bitmask if not random sweep
    if (!isRandom && !std::prev_permutation(bitmask.begin(), bitmask.end()))
    {
      M++;
      // Check for completion
      if (M > maxParallelTransfers)
      {
        printf("Sweep complete\n");
        break;
      }
      for (int i = 0; i < numPossible; i++)
        bitmask[i] = (i < M) ? 1 : 0;
    }
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
2330
2331
2332
2333
2334
2335
2336
2337
2338
2339
  fclose(fp);
}

void LogTransfers(FILE *fp, int const testNum, std::vector<Transfer> const& transfers)
{
  fprintf(fp, "# Test %d\n", testNum);
  fprintf(fp, "%d", -1 * (int)transfers.size());
  for (auto const& transfer : transfers)
  {
    fprintf(fp, " (%c%d->%c%d->%c%d %d %lu)",
gilbertlee-amd's avatar
gilbertlee-amd committed
2340
2341
2342
2343
            MemTypeStr[transfer.srcType[0]], transfer.srcIndex[0],
            ExeTypeStr[transfer.exeType],    transfer.exeIndex,
            MemTypeStr[transfer.dstType[0]], transfer.dstIndex[0],
            transfer.numSubExecs,
gilbertlee-amd's avatar
gilbertlee-amd committed
2344
2345
2346
2347
            transfer.numBytes);
  }
  fprintf(fp, "\n");
  fflush(fp);
Gilbert Lee's avatar
Gilbert Lee committed
2348
}
gilbertlee-amd's avatar
gilbertlee-amd committed
2349
2350
2351
2352
2353
2354
2355
2356
2357
2358
2359

std::string PtrVectorToStr(std::vector<float*> const& strVector, int const initOffset)
{
  std::stringstream ss;
  for (int i = 0; i < strVector.size(); ++i)
  {
    if (i) ss << " ";
    ss << (strVector[i] + initOffset);
  }
  return ss.str();
}