TransferBench.cpp 47.7 KB
Newer Older
Gilbert Lee's avatar
Gilbert Lee committed
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
/*
Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/

// This program measures simultaneous copy performance across multiple GPUs
// on the same node
#include <numa.h>
#include <numaif.h>
#include <stack>
#include <thread>

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

int main(int argc, char **argv)
{
  // 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
48
  // Determine number of bytes to run per Transfer
Gilbert Lee's avatar
Gilbert Lee committed
49
50
51
  // If a non-zero number of bytes is specified, use it
  // Otherwise generate array of bytes values to execute over
  std::vector<size_t> valuesOfN;
Gilbert Lee's avatar
Gilbert Lee committed
52
  size_t numBytesPerTransfer = argc > 2 ? atoll(argv[2]) : DEFAULT_BYTES_PER_TRANSFER;
Gilbert Lee's avatar
Gilbert Lee committed
53
54
55
56
57
58
  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
59
60
61
    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
62
63
    }
  }
Gilbert Lee's avatar
Gilbert Lee committed
64
  PopulateTestSizes(numBytesPerTransfer, ev.samplingFactor, valuesOfN);
Gilbert Lee's avatar
Gilbert Lee committed
65

Gilbert Lee's avatar
Gilbert Lee committed
66
  // Find the largest N to be used - memory will only be allocated once per set of simulatenous Transfers
Gilbert Lee's avatar
Gilbert Lee committed
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
  size_t maxN = valuesOfN[0];
  for (auto N : valuesOfN)
    maxN = std::max(maxN, N);

  // Execute only peer to peer benchmark mode, similar to rocm-bandwidth-test
  if (!strcmp(argv[1], "p2p") || !strcmp(argv[1], "p2p_rr") ||
      !strcmp(argv[1], "g2g") || !strcmp(argv[1], "g2g_rr"))
  {
    int numBlocksToUse = 0;
    if (argc > 3)
      numBlocksToUse = atoi(argv[3]);
    else
      HIP_CALL(hipDeviceGetAttribute(&numBlocksToUse, hipDeviceAttributeMultiprocessorCount, 0));

    // Perform either local read (+remote write) [EXE = SRC] or
    // remote read (+local write)                [EXE = DST]
    int readMode = (!strcmp(argv[1], "p2p_rr") || !strcmp(argv[1], "g2g_rr") ? 1 : 0);
    int skipCpu  = (!strcmp(argv[1], "g2g"   ) || !strcmp(argv[1], "g2g_rr") ? 1 : 0);

    // Execute peer to peer benchmark mode
Gilbert Lee's avatar
Gilbert Lee committed
87
    RunPeerToPeerBenchmarks(ev, numBytesPerTransfer / sizeof(float), numBlocksToUse, readMode, skipCpu);
Gilbert Lee's avatar
Gilbert Lee committed
88
89
90
    exit(0);
  }

Gilbert Lee's avatar
Gilbert Lee committed
91
  // Check that Transfer configuration file can be opened
Gilbert Lee's avatar
Gilbert Lee committed
92
93
94
  FILE* fp = fopen(argv[1], "r");
  if (!fp)
  {
Gilbert Lee's avatar
Gilbert Lee committed
95
    printf("[ERROR] Unable to open transfer configuration file: [%s]\n", argv[1]);
Gilbert Lee's avatar
Gilbert Lee committed
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
    exit(1);
  }

  // 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);
  }
  ev.DisplayEnvVars();

  int const initOffset = ev.byteOffset / sizeof(float);
  std::stack<std::thread> threads;

  // Collect the number of available CPUs/GPUs on this machine
  int numGpuDevices;
  HIP_CALL(hipGetDeviceCount(&numGpuDevices));
  int const numCpuDevices = numa_num_configured_nodes();

Gilbert Lee's avatar
Gilbert Lee committed
115
  // Track unique pair of transfers that get used
Gilbert Lee's avatar
Gilbert Lee committed
116
117
118
119
120
121
  std::set<std::pair<int, int>> peerAccessTracker;

  // Print CSV header
  if (ev.outputToCsv)
  {
    printf("Test,NumBytes,SrcMem,Executor,DstMem,CUs,BW(GB/s),Time(ms),"
Gilbert Lee's avatar
Gilbert Lee committed
122
           "TransferDesc,SrcAddr,DstAddr,ByteOffset,numWarmups,numIters\n");
Gilbert Lee's avatar
Gilbert Lee committed
123
124
  }

Gilbert Lee's avatar
Gilbert Lee committed
125
  // Loop over each line in the Transfer configuration file
Gilbert Lee's avatar
Gilbert Lee committed
126
127
128
129
130
131
132
  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
133
134
135
136
    // Parse transfers from configuration file
    TransferMap transferMap;
    ParseTransfers(line, numCpuDevices, numGpuDevices, transferMap);
    if (transferMap.size() == 0) continue;
Gilbert Lee's avatar
Gilbert Lee committed
137
138
139

    testNum++;

Gilbert Lee's avatar
Gilbert Lee committed
140
141
142
    // Prepare (maximum) memory for each transfer
    std::vector<Transfer*> transferList;
    for (auto& exeInfoPair : transferMap)
Gilbert Lee's avatar
Gilbert Lee committed
143
144
145
146
147
    {
      ExecutorInfo& exeInfo = exeInfoPair.second;
      exeInfo.totalTime = 0.0;
      exeInfo.totalBlocks = 0;

Gilbert Lee's avatar
Gilbert Lee committed
148
      for (Transfer& transfer : exeInfo.transfers)
Gilbert Lee's avatar
Gilbert Lee committed
149
      {
Gilbert Lee's avatar
Gilbert Lee committed
150
151
152
153
154
        // Get some aliases to transfer variables
        MemType const& exeMemType  = transfer.exeMemType;
        MemType const& srcMemType  = transfer.srcMemType;
        MemType const& dstMemType  = transfer.dstMemType;
        int     const& blocksToUse = transfer.numBlocksToUse;
Gilbert Lee's avatar
Gilbert Lee committed
155
156

        // Get potentially remapped device indices
Gilbert Lee's avatar
Gilbert Lee committed
157
158
159
        int const srcIndex = RemappedIndex(transfer.srcIndex, srcMemType);
        int const exeIndex = RemappedIndex(transfer.exeIndex, exeMemType);
        int const dstIndex = RemappedIndex(transfer.dstIndex, dstMemType);
Gilbert Lee's avatar
Gilbert Lee committed
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187

        // Enable peer-to-peer access if necessary (can only be called once per unique pair)
        if (exeMemType == MEM_GPU)
        {
          // Ensure executing GPU can access source memory
          if ((srcMemType == MEM_GPU || srcMemType == MEM_GPU_FINE) && srcIndex != exeIndex)
          {
            auto exeSrcPair = std::make_pair(exeIndex, srcIndex);
            if (!peerAccessTracker.count(exeSrcPair))
            {
              EnablePeerAccess(exeIndex, srcIndex);
              peerAccessTracker.insert(exeSrcPair);
            }
          }

          // Ensure executing GPU can access destination memory
          if ((dstMemType == MEM_GPU || dstMemType == MEM_GPU_FINE) && dstIndex != exeIndex)
          {
            auto exeDstPair = std::make_pair(exeIndex, dstIndex);
            if (!peerAccessTracker.count(exeDstPair))
            {
              EnablePeerAccess(exeIndex, dstIndex);
              peerAccessTracker.insert(exeDstPair);
            }
          }
        }

        // Allocate (maximum) source / destination memory based on type / device index
Gilbert Lee's avatar
Gilbert Lee committed
188
189
190
191
192
        AllocateMemory(srcMemType, srcIndex, maxN * sizeof(float) + ev.byteOffset, (void**)&transfer.srcMem);
        AllocateMemory(dstMemType, dstIndex, maxN * sizeof(float) + ev.byteOffset, (void**)&transfer.dstMem);
        transfer.blockParam.resize(exeMemType == MEM_CPU ? ev.numCpuPerTransfer : blocksToUse);
        exeInfo.totalBlocks += transfer.blockParam.size();
        transferList.push_back(&transfer);
Gilbert Lee's avatar
Gilbert Lee committed
193
194
195
196
197
198
199
200
201
202
      }

      // Prepare GPU resources for GPU executors
      MemType const exeMemType = exeInfoPair.first.first;
      int     const exeIndex   = RemappedIndex(exeInfoPair.first.second, exeMemType);
      if (exeMemType == MEM_GPU)
      {
        AllocateMemory(exeMemType, exeIndex, exeInfo.totalBlocks * sizeof(BlockParam),
                       (void**)&exeInfo.blockParamGpu);

Gilbert Lee's avatar
Gilbert Lee committed
203
204
205
206
207
        int const numTransfersToRun = ev.useSingleStream ? 1 : exeInfo.transfers.size();
        exeInfo.streams.resize(numTransfersToRun);
        exeInfo.startEvents.resize(numTransfersToRun);
        exeInfo.stopEvents.resize(numTransfersToRun);
        for (int i = 0; i < numTransfersToRun; ++i)
Gilbert Lee's avatar
Gilbert Lee committed
208
209
210
211
212
213
214
        {
          HIP_CALL(hipSetDevice(exeIndex));
          HIP_CALL(hipStreamCreate(&exeInfo.streams[i]));
          HIP_CALL(hipEventCreate(&exeInfo.startEvents[i]));
          HIP_CALL(hipEventCreate(&exeInfo.stopEvents[i]));
        }

Gilbert Lee's avatar
Gilbert Lee committed
215
216
        int transferOffset = 0;
        for (int i = 0; i < exeInfo.transfers.size(); i++)
Gilbert Lee's avatar
Gilbert Lee committed
217
        {
Gilbert Lee's avatar
Gilbert Lee committed
218
219
          exeInfo.transfers[i].blockParamGpuPtr = exeInfo.blockParamGpu + transferOffset;
          transferOffset += exeInfo.transfers[i].blockParam.size();
Gilbert Lee's avatar
Gilbert Lee committed
220
221
222
223
        }
      }
    }

Gilbert Lee's avatar
Gilbert Lee committed
224
    // Loop over all the different number of bytes to use per Transfer
Gilbert Lee's avatar
Gilbert Lee committed
225
226
227
228
229
    for (auto N : valuesOfN)
    {
      if (!ev.outputToCsv) printf("Test %d: [%lu bytes]\n", testNum, N * sizeof(float));

      // Prepare input memory and block parameters for current N
Gilbert Lee's avatar
Gilbert Lee committed
230
      for (auto& exeInfoPair : transferMap)
Gilbert Lee's avatar
Gilbert Lee committed
231
232
233
      {
        ExecutorInfo& exeInfo = exeInfoPair.second;

Gilbert Lee's avatar
Gilbert Lee committed
234
        int transferOffset = 0;
Gilbert Lee's avatar
Gilbert Lee committed
235

Gilbert Lee's avatar
Gilbert Lee committed
236
        for (int i = 0; i < exeInfo.transfers.size(); ++i)
Gilbert Lee's avatar
Gilbert Lee committed
237
        {
Gilbert Lee's avatar
Gilbert Lee committed
238
239
          Transfer& transfer = exeInfo.transfers[i];
          transfer.PrepareBlockParams(ev, N);
Gilbert Lee's avatar
Gilbert Lee committed
240
241

          // Copy block parameters to GPU for GPU executors
Gilbert Lee's avatar
Gilbert Lee committed
242
          if (transfer.exeMemType == MEM_GPU)
Gilbert Lee's avatar
Gilbert Lee committed
243
          {
Gilbert Lee's avatar
Gilbert Lee committed
244
245
246
            HIP_CALL(hipMemcpy(&exeInfo.blockParamGpu[transferOffset],
                               transfer.blockParam.data(),
                               transfer.blockParam.size() * sizeof(BlockParam),
Gilbert Lee's avatar
Gilbert Lee committed
247
                               hipMemcpyHostToDevice));
Gilbert Lee's avatar
Gilbert Lee committed
248
            transferOffset += transfer.blockParam.size();
Gilbert Lee's avatar
Gilbert Lee committed
249
250
251
252
253
254
          }
        }
      }

      // Launch kernels (warmup iterations are not counted)
      double totalCpuTime = 0;
Gilbert Lee's avatar
Gilbert Lee committed
255
256
      size_t numTimedIterations = 0;
      for (int iteration = -ev.numWarmups; ; iteration++)
Gilbert Lee's avatar
Gilbert Lee committed
257
      {
Gilbert Lee's avatar
Gilbert Lee committed
258
259
260
        if (ev.numIterations > 0 && iteration >= ev.numIterations) break;
        if (ev.numIterations < 0 && totalCpuTime > -ev.numIterations) break;

Gilbert Lee's avatar
Gilbert Lee committed
261
262
263
264
265
266
267
268
269
270
271
        // Pause before starting first timed iteration in interactive mode
        if (ev.useInteractive && iteration == 0)
        {
          printf("Hit <Enter> to continue: ");
          scanf("%*c");
          printf("\n");
        }

        // Start CPU timing for this iteration
        auto cpuStart = std::chrono::high_resolution_clock::now();

Gilbert Lee's avatar
Gilbert Lee committed
272
273
        // Execute all Transfers in parallel
        for (auto& exeInfoPair : transferMap)
Gilbert Lee's avatar
Gilbert Lee committed
274
275
        {
          ExecutorInfo& exeInfo = exeInfoPair.second;
Gilbert Lee's avatar
Gilbert Lee committed
276
277
278
          int const numTransfersToRun = ev.useSingleStream ? 1 : exeInfo.transfers.size();
          for (int i = 0; i < numTransfersToRun; ++i)
            threads.push(std::thread(RunTransfer, std::ref(ev), N, iteration, std::ref(exeInfo), i));
Gilbert Lee's avatar
Gilbert Lee committed
279
280
281
        }

        // Wait for all threads to finish
Gilbert Lee's avatar
Gilbert Lee committed
282
283
        int const numTransfers = threads.size();
        for (int i = 0; i < numTransfers; i++)
Gilbert Lee's avatar
Gilbert Lee committed
284
285
286
287
288
289
290
291
292
        {
          threads.top().join();
          threads.pop();
        }

        // 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();

Gilbert Lee's avatar
Gilbert Lee committed
293
294
295
296
297
        if (iteration >= 0)
        {
          ++numTimedIterations;
          totalCpuTime += deltaSec;
        }
Gilbert Lee's avatar
Gilbert Lee committed
298
299
300
301
302
303
304
305
306
307
      }

      // Pause for interactive mode
      if (ev.useInteractive)
      {
        printf("Transfers complete. Hit <Enter> to continue: ");
        scanf("%*c");
        printf("\n");
      }

Gilbert Lee's avatar
Gilbert Lee committed
308
309
310
311
      // Validate that each transfer has transferred correctly
      int const numTransfers = transferList.size();
      for (auto transfer : transferList)
        CheckOrFill(MODE_CHECK, N, ev.useMemset, ev.useHipCall, ev.fillPattern, transfer->dstMem + initOffset);
Gilbert Lee's avatar
Gilbert Lee committed
312
313

      // Report timings
Gilbert Lee's avatar
Gilbert Lee committed
314
315
      totalCpuTime = totalCpuTime / (1.0 * numTimedIterations) * 1000;
      double totalBandwidthGbs = (numTransfers * N * sizeof(float) / 1.0E6) / totalCpuTime;
Gilbert Lee's avatar
Gilbert Lee committed
316
317
318
319
      double maxGpuTime = 0;

      if (ev.useSingleStream)
      {
Gilbert Lee's avatar
Gilbert Lee committed
320
        for (auto& exeInfoPair : transferMap)
Gilbert Lee's avatar
Gilbert Lee committed
321
322
323
324
325
        {
          ExecutorInfo const& exeInfo = exeInfoPair.second;
          MemType const exeMemType    = exeInfoPair.first.first;
          int     const exeIndex      = exeInfoPair.first.second;

Gilbert Lee's avatar
Gilbert Lee committed
326
327
          double exeDurationMsec = exeInfo.totalTime / (1.0 * numTimedIterations);
          double exeBandwidthGbs = (exeInfo.transfers.size() * N * sizeof(float) / 1.0E9) / exeDurationMsec * 1000.0f;
Gilbert Lee's avatar
Gilbert Lee committed
328
329
330
331
          maxGpuTime = std::max(maxGpuTime, exeDurationMsec);

          if (!ev.outputToCsv)
          {
Gilbert Lee's avatar
Gilbert Lee committed
332
333
334
            printf(" Executor: %cPU %02d        (# Transfers %02lu)| %9.3f GB/s | %8.3f ms |\n",
                   MemTypeStr[exeMemType], exeIndex, exeInfo.transfers.size(), exeBandwidthGbs, exeDurationMsec);
            for (auto transfer : exeInfo.transfers)
Gilbert Lee's avatar
Gilbert Lee committed
335
            {
Gilbert Lee's avatar
Gilbert Lee committed
336
337
338
339
340
341
342
343
344
345
346
              double transferDurationMsec = transfer.transferTime / (1.0 * numTimedIterations);
              double transferBandwidthGbs = (N * sizeof(float) / 1.0E9) / transferDurationMsec * 1000.0f;

              printf("                            Transfer  %02d | %9.3f GB/s | %8.3f ms | %c%02d -> %c%02d:(%03d) -> %c%02d\n",
                     transfer.transferIndex,
                     transferBandwidthGbs,
                     transferDurationMsec,
                     MemTypeStr[transfer.srcMemType], transfer.srcIndex,
                     MemTypeStr[transfer.exeMemType], transfer.exeIndex,
                     transfer.exeMemType == MEM_CPU ? ev.numCpuPerTransfer : transfer.numBlocksToUse,
                     MemTypeStr[transfer.dstMemType], transfer.dstIndex);
Gilbert Lee's avatar
Gilbert Lee committed
347
348
349
350
            }
          }
          else
          {
Gilbert Lee's avatar
Gilbert Lee committed
351
            printf("%d,%lu,ALL,%c%02d,ALL,ALL,%.3f,%.3f,ALL,ALL,ALL,%d,%d,%lu\n",
Gilbert Lee's avatar
Gilbert Lee committed
352
353
354
355
                   testNum, N * sizeof(float),
                   MemTypeStr[exeMemType], exeIndex,
                   exeBandwidthGbs, exeDurationMsec,
                   ev.byteOffset,
Gilbert Lee's avatar
Gilbert Lee committed
356
                   ev.numWarmups, numTimedIterations);
Gilbert Lee's avatar
Gilbert Lee committed
357
358
359
360
361
          }
        }
      }
      else
      {
Gilbert Lee's avatar
Gilbert Lee committed
362
        for (auto transfer : transferList)
Gilbert Lee's avatar
Gilbert Lee committed
363
        {
Gilbert Lee's avatar
Gilbert Lee committed
364
365
366
          double transferDurationMsec = transfer->transferTime / (1.0 * numTimedIterations);
          double transferBandwidthGbs = (N * sizeof(float) / 1.0E9) / transferDurationMsec * 1000.0f;
          maxGpuTime = std::max(maxGpuTime, transferDurationMsec);
Gilbert Lee's avatar
Gilbert Lee committed
367
368
          if (!ev.outputToCsv)
          {
Gilbert Lee's avatar
Gilbert Lee committed
369
370
371
372
373
374
375
376
            printf(" Transfer %02d: %c%02d -> [%cPU %02d:%03d] -> %c%02d | %9.3f GB/s | %8.3f ms | %-16s\n",
                   transfer->transferIndex,
                   MemTypeStr[transfer->srcMemType], transfer->srcIndex,
                   MemTypeStr[transfer->exeMemType], transfer->exeIndex,
                   transfer->exeMemType == MEM_CPU ? ev.numCpuPerTransfer : transfer->numBlocksToUse,
                   MemTypeStr[transfer->dstMemType], transfer->dstIndex,
                   transferBandwidthGbs, transferDurationMsec,
                   GetTransferDesc(*transfer).c_str());
Gilbert Lee's avatar
Gilbert Lee committed
377
378
379
          }
          else
          {
Gilbert Lee's avatar
Gilbert Lee committed
380
            printf("%d,%lu,%c%02d,%c%02d,%c%02d,%d,%.3f,%.3f,%s,%p,%p,%d,%d,%lu\n",
Gilbert Lee's avatar
Gilbert Lee committed
381
                   testNum, N * sizeof(float),
Gilbert Lee's avatar
Gilbert Lee committed
382
383
384
385
386
387
388
                   MemTypeStr[transfer->srcMemType], transfer->srcIndex,
                   MemTypeStr[transfer->exeMemType], transfer->exeIndex,
                   MemTypeStr[transfer->dstMemType], transfer->dstIndex,
                   transfer->exeMemType == MEM_CPU ? ev.numCpuPerTransfer : transfer->numBlocksToUse,
                   transferBandwidthGbs, transferDurationMsec,
                   GetTransferDesc(*transfer).c_str(),
                   transfer->srcMem + initOffset, transfer->dstMem + initOffset,
Gilbert Lee's avatar
Gilbert Lee committed
389
                   ev.byteOffset,
Gilbert Lee's avatar
Gilbert Lee committed
390
                   ev.numWarmups, numTimedIterations);
Gilbert Lee's avatar
Gilbert Lee committed
391
392
393
394
395
396
397
          }
        }
      }

      // Display aggregate statistics
      if (!ev.outputToCsv)
      {
Gilbert Lee's avatar
Gilbert Lee committed
398
        printf(" Aggregate Bandwidth (CPU timed)         | %9.3f GB/s | %8.3f ms | Overhead: %.3f ms\n", totalBandwidthGbs, totalCpuTime,
Gilbert Lee's avatar
Gilbert Lee committed
399
400
401
402
               totalCpuTime - maxGpuTime);
      }
      else
      {
Gilbert Lee's avatar
Gilbert Lee committed
403
        printf("%d,%lu,ALL,ALL,ALL,ALL,%.3f,%.3f,ALL,ALL,ALL,%d,%d,%lu\n",
Gilbert Lee's avatar
Gilbert Lee committed
404
               testNum, N * sizeof(float), totalBandwidthGbs, totalCpuTime, ev.byteOffset,
Gilbert Lee's avatar
Gilbert Lee committed
405
               ev.numWarmups, numTimedIterations);
Gilbert Lee's avatar
Gilbert Lee committed
406
407
408
409
      }
    }

    // Release GPU memory
Gilbert Lee's avatar
Gilbert Lee committed
410
    for (auto exeInfoPair : transferMap)
Gilbert Lee's avatar
Gilbert Lee committed
411
412
    {
      ExecutorInfo& exeInfo = exeInfoPair.second;
Gilbert Lee's avatar
Gilbert Lee committed
413
      for (auto& transfer : exeInfo.transfers)
Gilbert Lee's avatar
Gilbert Lee committed
414
      {
Gilbert Lee's avatar
Gilbert Lee committed
415
416
417
418
        // Get some aliases to Transfer variables
        MemType const& exeMemType = transfer.exeMemType;
        MemType const& srcMemType = transfer.srcMemType;
        MemType const& dstMemType = transfer.dstMemType;
Gilbert Lee's avatar
Gilbert Lee committed
419
420

        // Allocate (maximum) source / destination memory based on type / device index
Gilbert Lee's avatar
Gilbert Lee committed
421
422
423
        DeallocateMemory(srcMemType, transfer.srcMem);
        DeallocateMemory(dstMemType, transfer.dstMem);
        transfer.blockParam.clear();
Gilbert Lee's avatar
Gilbert Lee committed
424
425
426
427
428
429
430
      }

      MemType const exeMemType = exeInfoPair.first.first;
      int     const exeIndex   = RemappedIndex(exeInfoPair.first.second, exeMemType);
      if (exeMemType == MEM_GPU)
      {
        DeallocateMemory(exeMemType, exeInfo.blockParamGpu);
Gilbert Lee's avatar
Gilbert Lee committed
431
432
        int const numTransfersToRun = ev.useSingleStream ? 1 : exeInfo.transfers.size();
        for (int i = 0; i < numTransfersToRun; ++i)
Gilbert Lee's avatar
Gilbert Lee committed
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
        {
          HIP_CALL(hipEventDestroy(exeInfo.startEvents[i]));
          HIP_CALL(hipEventDestroy(exeInfo.stopEvents[i]));
          HIP_CALL(hipStreamDestroy(exeInfo.streams[i]));
        }
      }
    }
  }
  fclose(fp);

  return 0;
}

void DisplayUsage(char const* cmdName)
{
Gilbert Lee's avatar
Gilbert Lee committed
448
  printf("TransferBench v%s\n", TB_VERSION);
Gilbert Lee's avatar
Gilbert Lee committed
449
450
451
452
453
454
455
456
457
458
459
460
461
  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
462
  printf("          - Filename of configFile containing Transfers to execute (see example.cfg for format)\n");
Gilbert Lee's avatar
Gilbert Lee committed
463
464
465
466
467
468
  printf("          - Name of preset benchmark:\n");
  printf("              p2p    - All CPU/GPU pairs benchmark\n");
  printf("              p2p_rr - All CPU/GPU pairs benchmark with remote reads\n");
  printf("              g2g    - All GPU/GPU pairs benchmark\n");
  printf("              g2g_rr - All GPU/GPU pairs benchmark with remote reads\n");
  printf("            - 3rd optional argument will be used as # of CUs to use (uses all by default)\n");
Gilbert Lee's avatar
Gilbert Lee committed
469
  printf("  N     : (Optional) Number of bytes to copy per Transfer.\n");
Gilbert Lee's avatar
Gilbert Lee committed
470
  printf("          If not specified, defaults to %lu bytes. Must be a multiple of 4 bytes\n",
Gilbert Lee's avatar
Gilbert Lee committed
471
         DEFAULT_BYTES_PER_TRANSFER);
Gilbert Lee's avatar
Gilbert Lee committed
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
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
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
  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();
}

int RemappedIndex(int const origIdx, MemType const memType)
{
  static std::vector<int> remapping;

  // No need to re-map CPU devices
  if (memType == MEM_CPU) return origIdx;

  // Build remapping on first use
  if (remapping.empty())
  {
    int numGpuDevices;
    HIP_CALL(hipGetDeviceCount(&numGpuDevices));
    remapping.resize(numGpuDevices);

    int const usePcieIndexing = getenv("USE_PCIE_INDEX") ? atoi(getenv("USE_PCIE_INDEX")) : 0;
    if (!usePcieIndexing)
    {
      // For HIP-based indexing no remapping is necessary
      for (int i = 0; i < numGpuDevices; ++i)
        remapping[i] = i;
    }
    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)
        remapping[i] = mapping[i].second;
    }
  }
  return remapping[origIdx];
}

void DisplayTopology(bool const outputToCsv)
{
  int numGpuDevices;
  HIP_CALL(hipGetDeviceCount(&numGpuDevices));

  if (outputToCsv)
  {
    printf("NumCpus,%d\n", numa_num_configured_nodes());
    printf("NumGpus,%d\n", numGpuDevices);
    printf("GPU");
    for (int j = 0; j < numGpuDevices; j++)
      printf(",GPU %02d", j);
    printf(",PCIe Bus ID,ClosestNUMA\n");
  }
  else
  {
    printf("\nDetected topology: %d CPU NUMA node(s)   %d GPU device(s)\n", numa_num_configured_nodes(), numGpuDevices);
    printf("        |");
    for (int j = 0; j < numGpuDevices; j++)
      printf(" GPU %02d |", j);
    printf(" PCIe Bus ID  | Closest NUMA\n");
    for (int j = 0; j <= numGpuDevices; j++)
      printf("--------+");
    printf("--------------+-------------\n");
  }

  char pciBusId[20];

  for (int i = 0; i < numGpuDevices; i++)
  {
    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;
        HIP_CALL(hipExtGetLinkTypeAndHopCount(RemappedIndex(i, MEM_GPU),
                                              RemappedIndex(j, MEM_GPU),
                                              &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 ? "," : " |");
      }
    }
    HIP_CALL(hipDeviceGetPCIBusId(pciBusId, 20, RemappedIndex(i, MEM_GPU)));
    if (outputToCsv)
      printf("%s,%d\n", pciBusId, GetClosestNumaNode(RemappedIndex(i, MEM_GPU)));
    else
      printf(" %11s |  %d  \n", pciBusId, GetClosestNumaNode(RemappedIndex(i, MEM_GPU)));
  }
}

Gilbert Lee's avatar
Gilbert Lee committed
583
void PopulateTestSizes(size_t const numBytesPerTransfer,
Gilbert Lee's avatar
Gilbert Lee committed
584
585
586
587
588
589
                       int const samplingFactor,
                       std::vector<size_t>& valuesOfN)
{
  valuesOfN.clear();

  // If the number of bytes is specified, use it
Gilbert Lee's avatar
Gilbert Lee committed
590
  if (numBytesPerTransfer != 0)
Gilbert Lee's avatar
Gilbert Lee committed
591
  {
Gilbert Lee's avatar
Gilbert Lee committed
592
    if (numBytesPerTransfer % 4)
Gilbert Lee's avatar
Gilbert Lee committed
593
    {
Gilbert Lee's avatar
Gilbert Lee committed
594
      printf("[ERROR] numBytesPerTransfer (%lu) must be a multiple of 4\n", numBytesPerTransfer);
Gilbert Lee's avatar
Gilbert Lee committed
595
596
      exit(1);
    }
Gilbert Lee's avatar
Gilbert Lee committed
597
    size_t N = numBytesPerTransfer / sizeof(float);
Gilbert Lee's avatar
Gilbert Lee committed
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
    valuesOfN.push_back(N);
  }
  else
  {
    // Otherwise generate a range of values
    // (Powers of 2, with samplingFactor samples between successive powers of 2)
    for (int N = 256; N <= (1<<27); N *= 2)
    {
      int delta = std::max(32, N / samplingFactor);
      int curr = N;
      while (curr < N * 2)
      {
        valuesOfN.push_back(curr);
        curr += delta;
      }
    }
  }
}

void ParseMemType(std::string const& token, int const numCpus, int const numGpus, MemType* memType, int* memIndex)
{
  char typeChar;
  if (sscanf(token.c_str(), " %c %d", &typeChar, memIndex) != 2)
  {
    printf("[ERROR] Unable to parse memory type token %s - expecting either 'B,C,G or F' followed by an index\n",
           token.c_str());
    exit(1);
  }

  switch (typeChar)
  {
  case 'C': case 'c': case 'B': case 'b':
    *memType = (typeChar == 'C' || typeChar == 'c') ? MEM_CPU : MEM_CPU_FINE;
    if (*memIndex < 0 || *memIndex >= numCpus)
    {
      printf("[ERROR] CPU index must be between 0 and %d (instead of %d)\n", numCpus-1, *memIndex);
      exit(1);
    }
    break;
  case 'G': case 'g': case 'F': case 'f':
    *memType = (typeChar == 'G' || typeChar == 'g') ? MEM_GPU : MEM_GPU_FINE;
    if (*memIndex < 0 || *memIndex >= numGpus)
    {
      printf("[ERROR] GPU index must be between 0 and %d (instead of %d)\n", numGpus-1, *memIndex);
      exit(1);
    }
    break;
  default:
    printf("[ERROR] Unrecognized memory type %s.  Expecting either 'B', 'C' or 'G' or 'F'\n", token.c_str());
    exit(1);
  }
}

Gilbert Lee's avatar
Gilbert Lee committed
651
652
// Helper function to parse a list of Transfer definitions
void ParseTransfers(char* line, int numCpus, int numGpus, TransferMap& transferMap)
Gilbert Lee's avatar
Gilbert Lee committed
653
654
655
656
657
{
  // 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
658
659
  transferMap.clear();
  int numTransfers = 0;
Gilbert Lee's avatar
Gilbert Lee committed
660
661

  std::istringstream iss(line);
Gilbert Lee's avatar
Gilbert Lee committed
662
  iss >> numTransfers;
Gilbert Lee's avatar
Gilbert Lee committed
663
664
665
666
667
  if (iss.fail()) return;

  std::string exeMem;
  std::string srcMem;
  std::string dstMem;
Gilbert Lee's avatar
Gilbert Lee committed
668
  if (numTransfers > 0)
Gilbert Lee's avatar
Gilbert Lee committed
669
670
671
672
673
674
675
676
677
  {
    // Method 1: Take in triples (srcMem, exeMem, dstMem)
    int numBlocksToUse;
    iss >> numBlocksToUse;
    if (numBlocksToUse <= 0 || iss.fail())
    {
      printf("Parsing error: Number of blocks to use (%d) must be greater than 0\n", numBlocksToUse);
      exit(1);
    }
Gilbert Lee's avatar
Gilbert Lee committed
678
    for (int i = 0; i < numTransfers; i++)
Gilbert Lee's avatar
Gilbert Lee committed
679
    {
Gilbert Lee's avatar
Gilbert Lee committed
680
681
      Transfer transfer;
      transfer.transferIndex = i;
Gilbert Lee's avatar
Gilbert Lee committed
682
683
684
      iss >> srcMem >> exeMem >> dstMem;
      if (iss.fail())
      {
Gilbert Lee's avatar
Gilbert Lee committed
685
        printf("Parsing error: Unable to read valid Transfer triplet (possibly missing a SRC or EXE or DST)\n");
Gilbert Lee's avatar
Gilbert Lee committed
686
687
        exit(1);
      }
Gilbert Lee's avatar
Gilbert Lee committed
688
689
690
691
      ParseMemType(srcMem, numCpus, numGpus, &transfer.srcMemType, &transfer.srcIndex);
      ParseMemType(exeMem, numCpus, numGpus, &transfer.exeMemType, &transfer.exeIndex);
      ParseMemType(dstMem, numCpus, numGpus, &transfer.dstMemType, &transfer.dstIndex);
      transfer.numBlocksToUse = numBlocksToUse;
Gilbert Lee's avatar
Gilbert Lee committed
692
693

      // Ensure executor is either CPU or GPU
Gilbert Lee's avatar
Gilbert Lee committed
694
      if (transfer.exeMemType != MEM_CPU && transfer.exeMemType != MEM_GPU)
Gilbert Lee's avatar
Gilbert Lee committed
695
696
      {
        printf("[ERROR] Executor must either be CPU ('C') or GPU ('G'), (from (%s->%s->%s %d))\n",
Gilbert Lee's avatar
Gilbert Lee committed
697
               srcMem.c_str(), exeMem.c_str(), dstMem.c_str(), transfer.numBlocksToUse);
Gilbert Lee's avatar
Gilbert Lee committed
698
699
700
        exit(1);
      }

Gilbert Lee's avatar
Gilbert Lee committed
701
702
703
704
      Executor executor(transfer.exeMemType, transfer.exeIndex);
      ExecutorInfo& executorInfo = transferMap[executor];
      executorInfo.totalBlocks += transfer.numBlocksToUse;
      executorInfo.transfers.push_back(transfer);
Gilbert Lee's avatar
Gilbert Lee committed
705
706
707
708
709
    }
  }
  else
  {
    // Method 2: Read in quads (srcMem, exeMem, dstMem,  Read common # blocks to use, then read (src, dst) doubles
Gilbert Lee's avatar
Gilbert Lee committed
710
    numTransfers *= -1;
Gilbert Lee's avatar
Gilbert Lee committed
711

Gilbert Lee's avatar
Gilbert Lee committed
712
    for (int i = 0; i < numTransfers; i++)
Gilbert Lee's avatar
Gilbert Lee committed
713
    {
Gilbert Lee's avatar
Gilbert Lee committed
714
715
716
      Transfer transfer;
      transfer.transferIndex = i;
      iss >> srcMem >> exeMem >> dstMem >> transfer.numBlocksToUse;
Gilbert Lee's avatar
Gilbert Lee committed
717
718
      if (iss.fail())
      {
Gilbert Lee's avatar
Gilbert Lee committed
719
        printf("Parsing error: Unable to read valid Transfer quadruple (possibly missing a SRC or EXE or DST or #CU)\n");
Gilbert Lee's avatar
Gilbert Lee committed
720
721
        exit(1);
      }
Gilbert Lee's avatar
Gilbert Lee committed
722
723
724
725
      ParseMemType(srcMem, numCpus, numGpus, &transfer.srcMemType, &transfer.srcIndex);
      ParseMemType(exeMem, numCpus, numGpus, &transfer.exeMemType, &transfer.exeIndex);
      ParseMemType(dstMem, numCpus, numGpus, &transfer.dstMemType, &transfer.dstIndex);
      if (transfer.exeMemType != MEM_CPU && transfer.exeMemType != MEM_GPU)
Gilbert Lee's avatar
Gilbert Lee committed
726
727
      {
        printf("[ERROR] Executor must either be CPU ('C') or GPU ('G'), (from (%s->%s->%s %d))\n"
Gilbert Lee's avatar
Gilbert Lee committed
728
,               srcMem.c_str(), exeMem.c_str(), dstMem.c_str(), transfer.numBlocksToUse);
Gilbert Lee's avatar
Gilbert Lee committed
729
730
731
        exit(1);
      }

Gilbert Lee's avatar
Gilbert Lee committed
732
733
734
735
      Executor executor(transfer.exeMemType, transfer.exeIndex);
      ExecutorInfo& executorInfo = transferMap[executor];
      executorInfo.totalBlocks += transfer.numBlocksToUse;
      executorInfo.transfers.push_back(transfer);
Gilbert Lee's avatar
Gilbert Lee committed
736
737
738
739
740
741
742
743
744
745
746
747
748
749
750
751
752
753
754
755
756
757
758
759
760
761
762
763
764
765
766
767
768
769
770
771
772
773
774
775
776
777
778
779
780
    }
  }
}

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));
  HIP_CALL(hipDeviceEnablePeerAccess(peerDeviceId, 0));
}

void AllocateMemory(MemType memType, int devIndex, size_t numBytes, void** memPtr)
{
  if (numBytes == 0)
  {
    printf("[ERROR] Unable to allocate 0 bytes\n");
    exit(1);
  }

  if (memType == MEM_CPU || memType == MEM_CPU_FINE)
  {
    // Set numa policy prior to call to hipHostMalloc
    // NOTE: It may be possible that the actual configured numa nodes do not start at 0
    //       so remapping may be necessary
    // Find the 'deviceId'-th available NUMA node
    int numaIdx = 0;
    for (int i = 0; i <= devIndex; i++)
      while (!numa_bitmask_isbitset(numa_get_mems_allowed(), numaIdx))
        ++numaIdx;

    unsigned long nodemask = (1ULL << numaIdx);
    long retCode = set_mempolicy(MPOL_BIND, &nodemask, sizeof(nodemask)*8);
    if (retCode)
    {
      printf("[ERROR] Unable to set NUMA memory policy to bind to NUMA node %d\n", numaIdx);
      exit(1);
    }

    // Allocate host-pinned memory (should respect NUMA mem policy)
Gilbert Lee's avatar
Gilbert Lee committed
781

Gilbert Lee's avatar
Gilbert Lee committed
782
783
784
785
786
787
788
789
790
791
792
793
794
795
796
797
798
799
800
801
802
803
804
805
806
807
808
809
810
811
812
813
814
815
816
817
818
819
820
821
822
823
824
825
826
827
828
829
830
831
832
833
834
835
836
837
838
839
840
841
842
843
844
845
846
847
848
849
850
851
852
853
854
855
856
857
858
859
860
861
862
863
864
865
866
867
868
869
870
871
872
873
874
875
876
877
878
879
880
881
882
883
884
885
886
887
888
889
890
891
892
893
894
895
896
897
898
899
900
901
902
903
904
905
906
907
908
909
910
911
912
913
914
915
916
917
918
919
920
921
922
923
924
925
926
927
928
929
930
931
932
933
934
935
936
937
938
939
940
941
942
943
944
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
    if (memType == MEM_CPU_FINE)
    {
      HIP_CALL(hipHostMalloc((void **)memPtr, numBytes, hipHostMallocNumaUser));
    }
    else
    {
      HIP_CALL(hipHostMalloc((void **)memPtr, numBytes, hipHostMallocNumaUser | hipHostMallocNonCoherent));
    }

    // Check that the allocated pages are actually on the correct NUMA node
    CheckPages((char*)*memPtr, numBytes, numaIdx);

    // Reset to default numa mem policy
    retCode = set_mempolicy(MPOL_DEFAULT, NULL, 8);
    if (retCode)
    {
      printf("[ERROR] Unable reset to default NUMA memory policy\n");
      exit(1);
    }
  }
  else 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)
  {
    HIP_CALL(hipSetDevice(devIndex));
    HIP_CALL(hipExtMallocWithFlags((void**)memPtr, numBytes, hipDeviceMallocFinegrained));
  }
  else
  {
    printf("[ERROR] Unsupported memory type %d\n", memType);
    exit(1);
  }
}

void DeallocateMemory(MemType memType, void* memPtr)
{
  if (memType == MEM_CPU || memType == MEM_CPU_FINE)
  {
    HIP_CALL(hipHostFree(memPtr));
  }
  else if (memType == MEM_GPU || memType == MEM_GPU_FINE)
  {
    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);
    printf("[ERROR] Ensure up-to-date ROCm is installed\n");
    exit(1);
  }
}

// Helper function to either fill a device pointer with pseudo-random data, or to check to see if it matches
void CheckOrFill(ModeType mode, int N, bool isMemset, bool isHipCall, std::vector<float>const& fillPattern, float* ptr)
{
  // Prepare reference resultx
  float* refBuffer = (float*)malloc(N * sizeof(float));
  if (isMemset)
  {
    if (isHipCall)
    {
      memset(refBuffer, 42, N * sizeof(float));
    }
    else
    {
      for (int i = 0; i < N; i++)
        refBuffer[i] = 1234.0f;
    }
  }
  else
  {
    // Fill with repeated pattern if specified
    size_t patternLen = fillPattern.size();
    if (patternLen > 0)
    {
      for (int i = 0; i < N; i++)
        refBuffer[i] = fillPattern[i % patternLen];
    }
    else // Otherwise fill with pseudo-random values
    {
      for (int i = 0; i < N; i++)
        refBuffer[i] = (i % 383 + 31);
    }
  }

  // Either fill the memory with the reference buffer, or compare against it
  if (mode == MODE_FILL)
  {
    HIP_CALL(hipMemcpy(ptr, refBuffer, N * sizeof(float), hipMemcpyDefault));
  }
  else if (mode == MODE_CHECK)
  {
    float* hostBuffer = (float*) malloc(N * sizeof(float));
    HIP_CALL(hipMemcpy(hostBuffer, ptr, N * sizeof(float), hipMemcpyDefault));
    for (int i = 0; i < N; i++)
    {
      if (refBuffer[i] != hostBuffer[i])
      {
        printf("[ERROR] Mismatch at element %d Ref: %f Actual: %f\n", i, refBuffer[i], hostBuffer[i]);
        exit(1);
      }
    }
    free(hostBuffer);
  }

  free(refBuffer);
}

std::string GetLinkTypeDesc(uint32_t linkType, uint32_t hopCount)
{
  char result[10];

  switch (linkType)
  {
  case HSA_AMD_LINK_INFO_TYPE_HYPERTRANSPORT: sprintf(result, "  HT-%d", hopCount); break;
  case HSA_AMD_LINK_INFO_TYPE_QPI           : sprintf(result, " QPI-%d", hopCount); break;
  case HSA_AMD_LINK_INFO_TYPE_PCIE          : sprintf(result, "PCIE-%d", hopCount); break;
  case HSA_AMD_LINK_INFO_TYPE_INFINBAND     : sprintf(result, "INFB-%d", hopCount); break;
  case HSA_AMD_LINK_INFO_TYPE_XGMI          : sprintf(result, "XGMI-%d", hopCount); break;
  default: sprintf(result, "??????");
  }
  return result;
}

std::string GetDesc(MemType srcMemType, int srcIndex,
                    MemType dstMemType, int dstIndex)
{
  if (srcMemType == MEM_CPU || srcMemType == MEM_CPU_FINE)
  {
    if (dstMemType == MEM_CPU || dstMemType == MEM_CPU_FINE)
      return (srcIndex == dstIndex) ? "LOCAL" : "NUMA";
    else if (dstMemType == MEM_GPU || dstMemType == MEM_GPU_FINE)
      return "PCIE";
    else
      goto error;
  }
  else if (srcMemType == MEM_GPU || srcMemType == MEM_GPU_FINE)
  {
    if (dstMemType == MEM_CPU || dstMemType == MEM_CPU_FINE)
      return "PCIE";
    else if (dstMemType == MEM_GPU || dstMemType == MEM_GPU_FINE)
    {
      if (srcIndex == dstIndex) return "LOCAL";
      else
      {
        uint32_t linkType, hopCount;
        HIP_CALL(hipExtGetLinkTypeAndHopCount(RemappedIndex(srcIndex, MEM_GPU),
                                              RemappedIndex(dstIndex, MEM_GPU),
                                              &linkType, &hopCount));
        return GetLinkTypeDesc(linkType, hopCount);
      }
    }
    else
      goto error;
  }
error:
  printf("[ERROR] Unrecognized memory type\n");
  exit(1);
}

Gilbert Lee's avatar
Gilbert Lee committed
979
std::string GetTransferDesc(Transfer const& transfer)
Gilbert Lee's avatar
Gilbert Lee committed
980
{
Gilbert Lee's avatar
Gilbert Lee committed
981
982
  return GetDesc(transfer.srcMemType, transfer.srcIndex, transfer.exeMemType, transfer.exeIndex) + "-"
    + GetDesc(transfer.exeMemType, transfer.exeIndex, transfer.dstMemType, transfer.dstIndex);
Gilbert Lee's avatar
Gilbert Lee committed
983
984
}

Gilbert Lee's avatar
Gilbert Lee committed
985
void RunTransfer(EnvVars const& ev, size_t const N, int const iteration, ExecutorInfo& exeInfo, int const transferIdx)
Gilbert Lee's avatar
Gilbert Lee committed
986
{
Gilbert Lee's avatar
Gilbert Lee committed
987
  Transfer& transfer = exeInfo.transfers[transferIdx];
Gilbert Lee's avatar
Gilbert Lee committed
988
989

  // GPU execution agent
Gilbert Lee's avatar
Gilbert Lee committed
990
  if (transfer.exeMemType == MEM_GPU)
Gilbert Lee's avatar
Gilbert Lee committed
991
992
  {
    // Switch to executing GPU
Gilbert Lee's avatar
Gilbert Lee committed
993
    int const exeIndex = RemappedIndex(transfer.exeIndex, MEM_GPU);
Gilbert Lee's avatar
Gilbert Lee committed
994
995
    HIP_CALL(hipSetDevice(exeIndex));

Gilbert Lee's avatar
Gilbert Lee committed
996
997
998
    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
999
1000
1001
1002
1003
1004

    int const initOffset = ev.byteOffset / sizeof(float);

    if (ev.useHipCall)
    {
      // Record start event
Gilbert Lee's avatar
Gilbert Lee committed
1005
      HIP_CALL(hipEventRecord(startEvent, stream));
Gilbert Lee's avatar
Gilbert Lee committed
1006
1007
1008

      // Execute hipMemset / hipMemcpy
      if (ev.useMemset)
Gilbert Lee's avatar
Gilbert Lee committed
1009
        HIP_CALL(hipMemsetAsync(transfer.dstMem + initOffset, 42, N * sizeof(float), stream));
Gilbert Lee's avatar
Gilbert Lee committed
1010
      else
Gilbert Lee's avatar
Gilbert Lee committed
1011
1012
        HIP_CALL(hipMemcpyAsync(transfer.dstMem + initOffset,
                                transfer.srcMem + initOffset,
Gilbert Lee's avatar
Gilbert Lee committed
1013
1014
1015
                                N * sizeof(float), hipMemcpyDefault,
                                stream));
      // Record stop event
Gilbert Lee's avatar
Gilbert Lee committed
1016
      HIP_CALL(hipEventRecord(stopEvent, stream));
Gilbert Lee's avatar
Gilbert Lee committed
1017
1018
1019
    }
    else
    {
Gilbert Lee's avatar
Gilbert Lee committed
1020
      int const numBlocksToRun = ev.useSingleStream ? exeInfo.totalBlocks : transfer.numBlocksToUse;
Gilbert Lee's avatar
Gilbert Lee committed
1021
1022
1023
1024
      hipExtLaunchKernelGGL(ev.useMemset ? GpuMemsetKernel : GpuCopyKernel,
                            dim3(numBlocksToRun, 1, 1),
                            dim3(BLOCKSIZE, 1, 1),
                            ev.sharedMemBytes, stream,
Gilbert Lee's avatar
Gilbert Lee committed
1025
1026
                            startEvent, stopEvent,
                            0, transfer.blockParamGpuPtr);
Gilbert Lee's avatar
Gilbert Lee committed
1027
1028
1029
1030
    }

    // 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
1031
    HIP_CALL(hipStreamSynchronize(stream));
Gilbert Lee's avatar
Gilbert Lee committed
1032
1033
1034
1035

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

Gilbert Lee's avatar
Gilbert Lee committed
1039
1040
1041
      if (ev.useSingleStream)
      {
        for (Transfer& currTransfer : exeInfo.transfers)
Gilbert Lee's avatar
Gilbert Lee committed
1042
        {
Gilbert Lee's avatar
Gilbert Lee committed
1043
1044
1045
          long long minStartCycle = currTransfer.blockParamGpuPtr[0].startCycle;
          long long maxStopCycle  = currTransfer.blockParamGpuPtr[0].stopCycle;
          for (int i = 1; i < currTransfer.numBlocksToUse; i++)
Gilbert Lee's avatar
Gilbert Lee committed
1046
          {
Gilbert Lee's avatar
Gilbert Lee committed
1047
1048
            minStartCycle = std::min(minStartCycle, currTransfer.blockParamGpuPtr[i].startCycle);
            maxStopCycle  = std::max(maxStopCycle,  currTransfer.blockParamGpuPtr[i].stopCycle);
Gilbert Lee's avatar
Gilbert Lee committed
1049
          }
Gilbert Lee's avatar
Gilbert Lee committed
1050
1051
1052
          int const wallClockRate = GetWallClockRate(exeIndex);
          double iterationTimeMs = (maxStopCycle - minStartCycle) / (double)(wallClockRate);
          currTransfer.transferTime += iterationTimeMs;
Gilbert Lee's avatar
Gilbert Lee committed
1053
        }
Gilbert Lee's avatar
Gilbert Lee committed
1054
1055
1056
1057
1058
        exeInfo.totalTime += gpuDeltaMsec;
      }
      else
      {
        transfer.transferTime += gpuDeltaMsec;
Gilbert Lee's avatar
Gilbert Lee committed
1059
1060
1061
      }
    }
  }
Gilbert Lee's avatar
Gilbert Lee committed
1062
  else if (transfer.exeMemType == MEM_CPU) // CPU execution agent
Gilbert Lee's avatar
Gilbert Lee committed
1063
1064
  {
    // Force this thread and all child threads onto correct NUMA node
Gilbert Lee's avatar
Gilbert Lee committed
1065
    if (numa_run_on_node(transfer.exeIndex))
Gilbert Lee's avatar
Gilbert Lee committed
1066
    {
Gilbert Lee's avatar
Gilbert Lee committed
1067
      printf("[ERROR] Unable to set CPU to NUMA node %d\n", transfer.exeIndex);
Gilbert Lee's avatar
Gilbert Lee committed
1068
1069
1070
1071
1072
1073
1074
1075
      exit(1);
    }

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

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

    // Launch child-threads to perform memcopies
Gilbert Lee's avatar
Gilbert Lee committed
1076
1077
    for (int i = 0; i < ev.numCpuPerTransfer; i++)
      childThreads.push_back(std::thread(ev.useMemset ? CpuMemsetKernel : CpuCopyKernel, std::ref(transfer.blockParam[i])));
Gilbert Lee's avatar
Gilbert Lee committed
1078
1079

    // Wait for child-threads to finish
Gilbert Lee's avatar
Gilbert Lee committed
1080
    for (int i = 0; i < ev.numCpuPerTransfer; i++)
Gilbert Lee's avatar
Gilbert Lee committed
1081
1082
1083
1084
1085
1086
      childThreads[i].join();

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

    // Record time if not a warmup iteration
    if (iteration >= 0)
Gilbert Lee's avatar
Gilbert Lee committed
1087
      transfer.transferTime += (std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() * 1000.0);
Gilbert Lee's avatar
Gilbert Lee committed
1088
1089
1090
1091
1092
1093
1094
1095
1096
1097
1098
1099
1100
1101
1102
1103
1104
1105
1106
  }
}

void RunPeerToPeerBenchmarks(EnvVars const& ev, size_t N, int numBlocksToUse, int readMode, int skipCpu)
{
  // Collect the number of available CPUs/GPUs on this machine
  int numGpus;
  HIP_CALL(hipGetDeviceCount(&numGpus));
  int const numCpus = numa_num_configured_nodes();
  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);

  if (!ev.outputToCsv)
  {
    printf("Performing copies in each direction of %lu bytes\n", N * sizeof(float));
Gilbert Lee's avatar
Gilbert Lee committed
1107
    printf("Using %d threads per NUMA node for CPU copies\n", ev.numCpuPerTransfer);
Gilbert Lee's avatar
Gilbert Lee committed
1108
1109
1110
1111
1112
1113
1114
1115
1116
1117
1118
1119
1120
1121
1122
1123
1124
1125
1126
1127
1128
1129
1130
1131
1132
1133
1134
1135
1136
1137
1138
1139
1140
1141
1142
1143
1144
1145
1146
1147
1148
1149
1150
1151
1152
1153
1154
1155
1156
1157
1158
1159
1160
1161
1162
1163
1164
1165
1166
1167
1168
1169
1170
1171
1172
1173
1174
1175
1176
1177
1178
1179
1180
1181
1182
1183
1184
1185
1186
1187
1188
1189
1190
1191
    printf("Using %d CUs per transfer\n", numBlocksToUse);
  }
  else
  {
    printf("SRC,DST,Direction,ReadMode,BW(GB/s),Bytes\n");
  }

  // Perform unidirectional / bidirectional
  for (int isBidirectional = 0; isBidirectional <= 1; isBidirectional++)
  {
    // Print header
    if (!ev.outputToCsv)
    {
      printf("%sdirectional copy peak bandwidth GB/s [%s read / %s write]\n", isBidirectional ? "Bi" : "Uni",
             readMode == 0 ? "Local" : "Remote",
             readMode == 0 ? "Remote" : "Local");
      printf("%10s", "D/D");
      if (!skipCpu)
      {
        for (int i = 0; i < numCpus; i++)
          printf("%7s %02d", "CPU", i);
      }
      for (int i = 0; i < numGpus; i++)
        printf("%7s %02d", "GPU", i);
      printf("\n");
    }

    // Loop over all possible src/dst pairs
    for (int src = 0; src < numDevices; src++)
    {
      MemType const& srcMemType = (src < numCpus ? MEM_CPU : MEM_GPU);
      if (skipCpu && srcMemType == MEM_CPU) continue;
      int srcIndex = (srcMemType == MEM_CPU ? src : src - numCpus);
      if (!ev.outputToCsv)
        printf("%7s %02d", (srcMemType == MEM_CPU) ? "CPU" : "GPU", srcIndex);
      for (int dst = 0; dst < numDevices; dst++)
      {
        MemType const& dstMemType = (dst < numCpus ? MEM_CPU : MEM_GPU);
        if (skipCpu && dstMemType == MEM_CPU) continue;
        int dstIndex = (dstMemType == MEM_CPU ? dst : dst - numCpus);
        double bandwidth = GetPeakBandwidth(ev, N, isBidirectional, readMode, numBlocksToUse,
                                            srcMemType, srcIndex, dstMemType, dstIndex);
        if (!ev.outputToCsv)
        {
          if (bandwidth == 0)
            printf("%10s", "N/A");
          else
            printf("%10.2f", bandwidth);
        }
        else
        {
          printf("%s %02d,%s %02d,%s,%s,%.2f,%lu\n",
                 srcMemType == MEM_CPU ? "CPU" : "GPU",
                 srcIndex,
                 dstMemType == MEM_CPU ? "CPU" : "GPU",
                 dstIndex,
                 isBidirectional ? "bidirectional" : "unidirectional",
                 readMode == 0 ? "Local" : "Remote",
                 bandwidth,
                 N * sizeof(float));
        }
        fflush(stdout);
      }
      if (!ev.outputToCsv) printf("\n");
    }
    if (!ev.outputToCsv) printf("\n");
  }
}

double GetPeakBandwidth(EnvVars const& ev,
                        size_t  const  N,
                        int     const  isBidirectional,
                        int     const  readMode,
                        int     const  numBlocksToUse,
                        MemType const  srcMemType,
                        int     const  srcIndex,
                        MemType const  dstMemType,
                        int     const  dstIndex)
{
  // Skip bidirectional on same device
  if (isBidirectional && srcMemType == dstMemType && srcIndex == dstIndex) return 0.0f;

  int const initOffset = ev.byteOffset / sizeof(float);

Gilbert Lee's avatar
Gilbert Lee committed
1192
1193
  // Prepare Transfers
  std::vector<Transfer*> transfers;
Gilbert Lee's avatar
Gilbert Lee committed
1194
1195
1196
  ExecutorInfo exeInfo[2];
  for (int i = 0; i < 2; i++)
  {
Gilbert Lee's avatar
Gilbert Lee committed
1197
    exeInfo[i].transfers.resize(1);
Gilbert Lee's avatar
Gilbert Lee committed
1198
1199
1200
    exeInfo[i].streams.resize(1);
    exeInfo[i].startEvents.resize(1);
    exeInfo[i].stopEvents.resize(1);
Gilbert Lee's avatar
Gilbert Lee committed
1201
    transfers.push_back(&exeInfo[i].transfers[0]);
Gilbert Lee's avatar
Gilbert Lee committed
1202
1203
  }

Gilbert Lee's avatar
Gilbert Lee committed
1204
1205
1206
1207
  transfers[0]->srcMemType = transfers[1]->dstMemType = srcMemType;
  transfers[0]->dstMemType = transfers[1]->srcMemType = dstMemType;
  transfers[0]->srcIndex   = transfers[1]->dstIndex   = RemappedIndex(srcIndex, srcMemType);
  transfers[0]->dstIndex   = transfers[1]->srcIndex   = RemappedIndex(dstIndex, dstMemType);
Gilbert Lee's avatar
Gilbert Lee committed
1208
1209

  // Either perform (local read + remote write), or (remote read + local write)
Gilbert Lee's avatar
Gilbert Lee committed
1210
1211
1212
1213
  transfers[0]->exeMemType = (readMode == 0 ? srcMemType : dstMemType);
  transfers[1]->exeMemType = (readMode == 0 ? dstMemType : srcMemType);
  transfers[0]->exeIndex   = RemappedIndex((readMode == 0 ? srcIndex : dstIndex), transfers[0]->exeMemType);
  transfers[1]->exeIndex   = RemappedIndex((readMode == 0 ? dstIndex : srcIndex), transfers[1]->exeMemType);
Gilbert Lee's avatar
Gilbert Lee committed
1214
1215
1216

  for (int i = 0; i <= isBidirectional; i++)
  {
Gilbert Lee's avatar
Gilbert Lee committed
1217
1218
1219
1220
    AllocateMemory(transfers[i]->srcMemType, transfers[i]->srcIndex,
                   N * sizeof(float) + ev.byteOffset, (void**)&transfers[i]->srcMem);
    AllocateMemory(transfers[i]->dstMemType, transfers[i]->dstIndex,
                   N * sizeof(float) + ev.byteOffset, (void**)&transfers[i]->dstMem);
Gilbert Lee's avatar
Gilbert Lee committed
1221
1222

    // Prepare block parameters on CPU
Gilbert Lee's avatar
Gilbert Lee committed
1223
1224
1225
    transfers[i]->numBlocksToUse = (transfers[i]->exeMemType == MEM_GPU) ? numBlocksToUse : ev.numCpuPerTransfer;
    transfers[i]->blockParam.resize(transfers[i]->numBlocksToUse);
    transfers[i]->PrepareBlockParams(ev, N);
Gilbert Lee's avatar
Gilbert Lee committed
1226

Gilbert Lee's avatar
Gilbert Lee committed
1227
    if (transfers[i]->exeMemType == MEM_GPU)
Gilbert Lee's avatar
Gilbert Lee committed
1228
1229
    {
      // Copy block parameters onto GPU
Gilbert Lee's avatar
Gilbert Lee committed
1230
1231
1232
1233
      AllocateMemory(MEM_GPU, transfers[i]->exeIndex, numBlocksToUse * sizeof(BlockParam),
                     (void **)&transfers[i]->blockParamGpuPtr);
      HIP_CALL(hipMemcpy(transfers[i]->blockParamGpuPtr,
                         transfers[i]->blockParam.data(),
Gilbert Lee's avatar
Gilbert Lee committed
1234
1235
1236
1237
                         numBlocksToUse * sizeof(BlockParam),
                         hipMemcpyHostToDevice));

      // Prepare GPU resources
Gilbert Lee's avatar
Gilbert Lee committed
1238
      HIP_CALL(hipSetDevice(transfers[i]->exeIndex));
Gilbert Lee's avatar
Gilbert Lee committed
1239
1240
1241
1242
1243
1244
1245
1246
1247
1248
1249
1250
1251
      HIP_CALL(hipStreamCreate(&exeInfo[i].streams[0]));
      HIP_CALL(hipEventCreate(&exeInfo[i].startEvents[0]));
      HIP_CALL(hipEventCreate(&exeInfo[i].stopEvents[0]));
    }
  }

  std::stack<std::thread> threads;

  // Perform iteration
  for (int iteration = -ev.numWarmups; iteration < ev.numIterations; iteration++)
  {
    // Perform timed iterations
    for (int i = 0; i <= isBidirectional; i++)
Gilbert Lee's avatar
Gilbert Lee committed
1252
      threads.push(std::thread(RunTransfer, std::ref(ev), N, iteration, std::ref(exeInfo[i]), 0));
Gilbert Lee's avatar
Gilbert Lee committed
1253
1254
1255
1256
1257
1258
1259
1260
1261

    // Wait for all threads to finish
    for (int i = 0; i <= isBidirectional; i++)
    {
      threads.top().join();
      threads.pop();
    }
  }

Gilbert Lee's avatar
Gilbert Lee committed
1262
  // Validate that each Transfer has transferred correctly
Gilbert Lee's avatar
Gilbert Lee committed
1263
  for (int i = 0; i <= isBidirectional; i++)
Gilbert Lee's avatar
Gilbert Lee committed
1264
    CheckOrFill(MODE_CHECK, N, ev.useMemset, ev.useHipCall, ev.fillPattern, transfers[i]->dstMem + initOffset);
Gilbert Lee's avatar
Gilbert Lee committed
1265
1266
1267
1268
1269

  // Collect aggregate bandwidth
  double totalBandwidth = 0;
  for (int i = 0; i <= isBidirectional; i++)
  {
Gilbert Lee's avatar
Gilbert Lee committed
1270
1271
1272
    double transferDurationMsec = transfers[i]->transferTime / (1.0 * ev.numIterations);
    double transferBandwidthGbs = (N * sizeof(float) / 1.0E9) / transferDurationMsec * 1000.0f;
    totalBandwidth += transferBandwidthGbs;
Gilbert Lee's avatar
Gilbert Lee committed
1273
1274
1275
1276
1277
  }

  // Release GPU memory
  for (int i = 0; i <= isBidirectional; i++)
  {
Gilbert Lee's avatar
Gilbert Lee committed
1278
1279
    DeallocateMemory(transfers[i]->srcMemType, transfers[i]->srcMem);
    DeallocateMemory(transfers[i]->dstMemType, transfers[i]->dstMem);
Gilbert Lee's avatar
Gilbert Lee committed
1280

Gilbert Lee's avatar
Gilbert Lee committed
1281
    if (transfers[i]->exeMemType == MEM_GPU)
Gilbert Lee's avatar
Gilbert Lee committed
1282
    {
Gilbert Lee's avatar
Gilbert Lee committed
1283
      DeallocateMemory(MEM_GPU, transfers[i]->blockParamGpuPtr);
Gilbert Lee's avatar
Gilbert Lee committed
1284
1285
1286
1287
1288
1289
1290
1291
      HIP_CALL(hipStreamDestroy(exeInfo[i].streams[0]));
      HIP_CALL(hipEventDestroy(exeInfo[i].startEvents[0]));
      HIP_CALL(hipEventDestroy(exeInfo[i].stopEvents[0]));
    }
  }
  return totalBandwidth;
}

Gilbert Lee's avatar
Gilbert Lee committed
1292
void Transfer::PrepareBlockParams(EnvVars const& ev, size_t const N)
Gilbert Lee's avatar
Gilbert Lee committed
1293
1294
1295
1296
1297
1298
1299
{
  int const initOffset = ev.byteOffset / sizeof(float);

  // Initialize source memory with patterned data
  CheckOrFill(MODE_FILL, N, ev.useMemset, ev.useHipCall, ev.fillPattern, this->srcMem + initOffset);

  // Each block needs to know src/dst pointers and how many elements to transfer
Gilbert Lee's avatar
Gilbert Lee committed
1300
  // Figure out the sub-array each block does for this Transfer
Gilbert Lee's avatar
Gilbert Lee committed
1301
1302
1303
1304
1305
1306
1307
1308
1309
1310
1311
1312
1313
1314
1315
1316
1317
1318
1319
1320
  // - Partition N as evenly as possible, but try to keep blocks as multiples of BLOCK_BYTES bytes,
  //   except the very last one, for alignment reasons
  int const targetMultiple = ev.blockBytes / sizeof(float);
  int const maxNumBlocksToUse = std::min((N + targetMultiple - 1) / targetMultiple, this->blockParam.size());
  size_t assigned = 0;
  for (int j = 0; j < this->blockParam.size(); j++)
  {
    int    const blocksLeft = std::max(0, maxNumBlocksToUse - j);
    size_t const leftover   = N - assigned;
    size_t const roundedN   = (leftover + targetMultiple - 1) / targetMultiple;

    BlockParam& param = this->blockParam[j];
    param.N          = blocksLeft ? std::min(leftover, ((roundedN / blocksLeft) * targetMultiple)) : 0;
    param.src        = this->srcMem + assigned + initOffset;
    param.dst        = this->dstMem + assigned + initOffset;
    param.startCycle = 0;
    param.stopCycle  = 0;
    assigned += param.N;
  }

Gilbert Lee's avatar
Gilbert Lee committed
1321
  this->transferTime = 0.0;
Gilbert Lee's avatar
Gilbert Lee committed
1322
1323
1324
1325
1326
1327
1328
1329
1330
1331
1332
1333
1334
1335
1336
1337
1338
1339
1340
1341
1342
1343
1344
1345
1346
1347
1348
1349
1350
}

// 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);

    hipDeviceProp_t prop;
    for (int i = 0; i < numGpuDevices; i++)
    {
      HIP_CALL(hipGetDeviceProperties(&prop, i));
      int value = 25000;
      switch (prop.gcnArch)
      {
      case 906: case 910: value = 25000; break;
      default:
        printf("Unrecognized GCN arch %d\n", prop.gcnArch);
      }
      wallClockPerDeviceMhz[i] = value;
    }
  }
  return wallClockPerDeviceMhz[deviceId];
}