TransferBench.cpp 69.3 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
25
26

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>
Gilbert Lee's avatar
Gilbert Lee committed
27
#include <random>
Gilbert Lee's avatar
Gilbert Lee committed
28
29
30
31
32
33
34
35
#include <stack>
#include <thread>

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

int main(int argc, char **argv)
{
Gilbert Lee's avatar
Gilbert Lee committed
36
37
38
39
40
41
42
  // 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
43
44
45
46
47
48
49
50
51
52
53
54
  // 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
55
56
  // 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
57
58
59
60
61
62
  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
63
64
65
    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
66
67
    }
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
68
69
70
71
72
  if (numBytesPerTransfer % 4)
  {
    printf("[ERROR] numBytesPerTransfer (%lu) must be a multiple of 4\n", numBytesPerTransfer);
    exit(1);
  }
Gilbert Lee's avatar
Gilbert Lee committed
73

Gilbert Lee's avatar
Gilbert Lee committed
74
75
76
77
  // 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
78
79
    int numGpuSubExecs = (argc > 3 ? atoi(argv[3]) : 4);
    int numCpuSubExecs = (argc > 4 ? atoi(argv[4]) : 4);
gilbertlee-amd's avatar
gilbertlee-amd committed
80

81
    ev.configMode = CFG_SWEEP;
gilbertlee-amd's avatar
gilbertlee-amd committed
82
    RunSweepPreset(ev, numBytesPerTransfer, numGpuSubExecs, numCpuSubExecs, !strcmp(argv[1], "rsweep"));
Gilbert Lee's avatar
Gilbert Lee committed
83
84
85
    exit(0);
  }
  // - Tests that benchmark peer-to-peer performance
gilbertlee-amd's avatar
gilbertlee-amd committed
86
  else if (!strcmp(argv[1], "p2p"))
Gilbert Lee's avatar
Gilbert Lee committed
87
  {
88
    ev.configMode = CFG_P2P;
gilbertlee-amd's avatar
gilbertlee-amd committed
89
    RunPeerToPeerBenchmarks(ev, numBytesPerTransfer / sizeof(float));
Gilbert Lee's avatar
Gilbert Lee committed
90
91
    exit(0);
  }
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
  // - 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
107
108
109
110
111
112
113
114
115
116
117
  // - 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);
  }
Gilbert Lee's avatar
Gilbert Lee committed
118

Gilbert Lee's avatar
Gilbert Lee committed
119
  // Check that Transfer configuration file can be opened
120
  ev.configMode = CFG_FILE;
Gilbert Lee's avatar
Gilbert Lee committed
121
122
123
  FILE* fp = fopen(argv[1], "r");
  if (!fp)
  {
Gilbert Lee's avatar
Gilbert Lee committed
124
    printf("[ERROR] Unable to open transfer configuration file: [%s]\n", argv[1]);
Gilbert Lee's avatar
Gilbert Lee committed
125
126
127
    exit(1);
  }

Gilbert Lee's avatar
Gilbert Lee committed
128
  // Print environment variables and CSV header
Gilbert Lee's avatar
Gilbert Lee committed
129
130
131
  ev.DisplayEnvVars();
  if (ev.outputToCsv)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
132
    printf("Test#,Transfer#,NumBytes,Src,Exe,Dst,CUs,BW(GB/s),Time(ms),SrcAddr,DstAddr\n");
Gilbert Lee's avatar
Gilbert Lee committed
133
134
135
136
137
138
139
140
141
  }

  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
142
143
144
145
    // 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
146

gilbertlee-amd's avatar
gilbertlee-amd committed
147
148
149
150
151
152
153
154
155
156
157
    // 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
158
        int delta = std::max(1, N / ev.samplingFactor);
gilbertlee-amd's avatar
gilbertlee-amd committed
159
160
161
        int curr = N;
        while (curr < N * 2)
        {
gilbertlee-amd's avatar
gilbertlee-amd committed
162
          ExecuteTransfers(ev, ++testNum, curr, transfers);
gilbertlee-amd's avatar
gilbertlee-amd committed
163
164
165
166
          curr += delta;
        }
      }
    }
Gilbert Lee's avatar
Gilbert Lee committed
167
168
  }
  fclose(fp);
Gilbert Lee's avatar
Gilbert Lee committed
169

Gilbert Lee's avatar
Gilbert Lee committed
170
171
  return 0;
}
Gilbert Lee's avatar
Gilbert Lee committed
172

Gilbert Lee's avatar
Gilbert Lee committed
173
void ExecuteTransfers(EnvVars const& ev,
gilbertlee-amd's avatar
gilbertlee-amd committed
174
175
176
                      int const testNum,
                      size_t const N,
                      std::vector<Transfer>& transfers,
gilbertlee-amd's avatar
gilbertlee-amd committed
177
178
                      bool verbose,
                      double* totalBandwidthCpu)
Gilbert Lee's avatar
Gilbert Lee committed
179
180
{
  int const initOffset = ev.byteOffset / sizeof(float);
Gilbert Lee's avatar
Gilbert Lee committed
181

Gilbert Lee's avatar
Gilbert Lee committed
182
183
  // Map transfers by executor
  TransferMap transferMap;
gilbertlee-amd's avatar
gilbertlee-amd committed
184
  for (int i = 0; i < transfers.size(); i++)
Gilbert Lee's avatar
Gilbert Lee committed
185
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
186
187
    Transfer& transfer = transfers[i];
    transfer.transferIndex = i;
gilbertlee-amd's avatar
gilbertlee-amd committed
188
    Executor executor(transfer.exeType, transfer.exeIndex);
Gilbert Lee's avatar
Gilbert Lee committed
189
    ExecutorInfo& executorInfo = transferMap[executor];
gilbertlee-amd's avatar
gilbertlee-amd committed
190
    executorInfo.transfers.push_back(&transfer);
Gilbert Lee's avatar
Gilbert Lee committed
191
  }
Gilbert Lee's avatar
Gilbert Lee committed
192

gilbertlee-amd's avatar
gilbertlee-amd committed
193
  // Loop over each executor and prepare sub-executors
gilbertlee-amd's avatar
gilbertlee-amd committed
194
  std::map<int, Transfer*> transferList;
Gilbert Lee's avatar
Gilbert Lee committed
195
196
197
  for (auto& exeInfoPair : transferMap)
  {
    Executor const& executor = exeInfoPair.first;
gilbertlee-amd's avatar
gilbertlee-amd committed
198
199
200
201
    ExecutorInfo& exeInfo    = exeInfoPair.second;
    ExeType const exeType    = executor.first;
    int     const exeIndex   = RemappedIndex(executor.second, IsCpuType(exeType));

Gilbert Lee's avatar
Gilbert Lee committed
202
    exeInfo.totalTime = 0.0;
gilbertlee-amd's avatar
gilbertlee-amd committed
203
    exeInfo.totalSubExecs = 0;
Gilbert Lee's avatar
Gilbert Lee committed
204
205

    // Loop over each transfer this executor is involved in
gilbertlee-amd's avatar
gilbertlee-amd committed
206
    for (Transfer* transfer : exeInfo.transfers)
Gilbert Lee's avatar
Gilbert Lee committed
207
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
208
209
210
211
212
213
      // 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
214
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
215
216
217
        MemType const& srcType  = transfer->srcType[iSrc];
        int     const  srcIndex    = RemappedIndex(transfer->srcIndex[iSrc], IsCpuType(srcType));

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

gilbertlee-amd's avatar
gilbertlee-amd committed
222
223
224
225
226
227
228
229
230
231
        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
232
        // Ensure executing GPU can access destination memory
gilbertlee-amd's avatar
gilbertlee-amd committed
233
        if (IsGpuType(exeType) == MEM_GPU && IsGpuType(dstType) && dstIndex != exeIndex)
Gilbert Lee's avatar
Gilbert Lee committed
234
235
          EnablePeerAccess(exeIndex, dstIndex);

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

gilbertlee-amd's avatar
gilbertlee-amd committed
239
      exeInfo.totalSubExecs += transfer->numSubExecs;
gilbertlee-amd's avatar
gilbertlee-amd committed
240
      transferList[transfer->transferIndex] = transfer;
Gilbert Lee's avatar
Gilbert Lee committed
241
242
    }

gilbertlee-amd's avatar
gilbertlee-amd committed
243
244
    // Prepare additional requirement for GPU-based executors
    if (IsGpuType(exeType))
Gilbert Lee's avatar
Gilbert Lee committed
245
    {
246
247
      HIP_CALL(hipSetDevice(exeIndex));

gilbertlee-amd's avatar
gilbertlee-amd committed
248
249
250
251
252
253
      // 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
254
      {
Gilbert Lee's avatar
Gilbert Lee committed
255
256
257
258
        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
259

gilbertlee-amd's avatar
gilbertlee-amd committed
260
      if (exeType == EXE_GPU_GFX)
Gilbert Lee's avatar
Gilbert Lee committed
261
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
262
263
264
265
        // 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
        AllocateMemory(MEM_GPU, exeIndex, exeInfo.totalSubExecs * sizeof(SubExecParam),
                       (void**)&exeInfo.subExecParamGpu);
Gilbert Lee's avatar
Gilbert Lee committed
266
267
268
      }
    }
  }
Gilbert Lee's avatar
Gilbert Lee committed
269

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

  // Prepare input memory and block parameters for current N
273
  bool isSrcCorrect = true;
gilbertlee-amd's avatar
gilbertlee-amd committed
274
  for (auto& exeInfoPair : transferMap)
Gilbert Lee's avatar
Gilbert Lee committed
275
  {
276
277
278
279
280
    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
281
    exeInfo.totalBytes = 0;
Gilbert Lee's avatar
Gilbert Lee committed
282

gilbertlee-amd's avatar
gilbertlee-amd committed
283
284
    int transferOffset = 0;
    for (int i = 0; i < exeInfo.transfers.size(); ++i)
Gilbert Lee's avatar
Gilbert Lee committed
285
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
286
287
      // 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
288
      transfer->PrepareSubExecParams(ev);
289
      isSrcCorrect &= transfer->PrepareSrc(ev);
gilbertlee-amd's avatar
gilbertlee-amd committed
290
      exeInfo.totalBytes += transfer->numBytesActual;
Gilbert Lee's avatar
Gilbert Lee committed
291

gilbertlee-amd's avatar
gilbertlee-amd committed
292
      // Copy block parameters to GPU for GPU executors
gilbertlee-amd's avatar
gilbertlee-amd committed
293
      if (transfer->exeType == EXE_GPU_GFX)
Gilbert Lee's avatar
Gilbert Lee committed
294
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
295
        exeInfo.transfers[i]->subExecParamGpuPtr = exeInfo.subExecParamGpu + transferOffset;
296
        HIP_CALL(hipSetDevice(exeIndex));
gilbertlee-amd's avatar
gilbertlee-amd committed
297
298
299
        HIP_CALL(hipMemcpy(&exeInfo.subExecParamGpu[transferOffset],
                           transfer->subExecParam.data(),
                           transfer->subExecParam.size() * sizeof(SubExecParam),
gilbertlee-amd's avatar
gilbertlee-amd committed
300
                           hipMemcpyHostToDevice));
301
302
        HIP_CALL(hipDeviceSynchronize());

gilbertlee-amd's avatar
gilbertlee-amd committed
303
        transferOffset += transfer->subExecParam.size();
Gilbert Lee's avatar
Gilbert Lee committed
304
305
      }
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
306
  }
Gilbert Lee's avatar
Gilbert Lee committed
307

gilbertlee-amd's avatar
gilbertlee-amd committed
308
309
310
311
  // Launch kernels (warmup iterations are not counted)
  double totalCpuTime = 0;
  size_t numTimedIterations = 0;
  std::stack<std::thread> threads;
312
  for (int iteration = -ev.numWarmups; isSrcCorrect; iteration++)
gilbertlee-amd's avatar
gilbertlee-amd committed
313
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
314
    if (ev.numIterations > 0 && iteration    >= ev.numIterations) break;
gilbertlee-amd's avatar
gilbertlee-amd committed
315
    if (ev.numIterations < 0 && totalCpuTime > -ev.numIterations) break;
Gilbert Lee's avatar
Gilbert Lee committed
316

gilbertlee-amd's avatar
gilbertlee-amd committed
317
318
    // Pause before starting first timed iteration in interactive mode
    if (verbose && ev.useInteractive && iteration == 0)
Gilbert Lee's avatar
Gilbert Lee committed
319
    {
320
321
322
323
      printf("Memory prepared:\n");

      for (Transfer& transfer : transfers)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
324
325
326
327
328
        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]);
329
      }
gilbertlee-amd's avatar
gilbertlee-amd committed
330
      printf("Hit <Enter> to continue: ");
331
332
333
334
335
      if (scanf("%*c") != 0)
      {
        printf("[ERROR] Unexpected input\n");
        exit(1);
      }
Gilbert Lee's avatar
Gilbert Lee committed
336
337
      printf("\n");
    }
Gilbert Lee's avatar
Gilbert Lee committed
338

gilbertlee-amd's avatar
gilbertlee-amd committed
339
340
341
342
343
    // Start CPU timing for this iteration
    auto cpuStart = std::chrono::high_resolution_clock::now();

    // Execute all Transfers in parallel
    for (auto& exeInfoPair : transferMap)
344
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
345
      ExecutorInfo& exeInfo = exeInfoPair.second;
gilbertlee-amd's avatar
gilbertlee-amd committed
346
347
348
      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
349
350
      for (int i = 0; i < numTransfersToRun; ++i)
        threads.push(std::thread(RunTransfer, std::ref(ev), iteration, std::ref(exeInfo), i));
351
    }
Gilbert Lee's avatar
Gilbert Lee committed
352

gilbertlee-amd's avatar
gilbertlee-amd committed
353
354
355
356
357
358
359
    // 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
360

gilbertlee-amd's avatar
gilbertlee-amd committed
361
362
363
364
365
    // 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
366
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
367
368
369
370
      ++numTimedIterations;
      totalCpuTime += deltaSec;
    }
  }
Gilbert Lee's avatar
Gilbert Lee committed
371

gilbertlee-amd's avatar
gilbertlee-amd committed
372
  // Pause for interactive mode
373
  if (verbose && isSrcCorrect && ev.useInteractive)
gilbertlee-amd's avatar
gilbertlee-amd committed
374
375
  {
    printf("Transfers complete. Hit <Enter> to continue: ");
376
377
378
379
380
    if (scanf("%*c") != 0)
    {
      printf("[ERROR] Unexpected input\n");
      exit(1);
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
381
382
    printf("\n");
  }
Gilbert Lee's avatar
Gilbert Lee committed
383

gilbertlee-amd's avatar
gilbertlee-amd committed
384
385
386
  // Validate that each transfer has transferred correctly
  size_t totalBytesTransferred = 0;
  int const numTransfers = transferList.size();
gilbertlee-amd's avatar
gilbertlee-amd committed
387

gilbertlee-amd's avatar
gilbertlee-amd committed
388
389
390
  for (auto transferPair : transferList)
  {
    Transfer* transfer = transferPair.second;
gilbertlee-amd's avatar
gilbertlee-amd committed
391
392
    transfer->ValidateDst(ev);
    totalBytesTransferred += transfer->numBytesActual;
gilbertlee-amd's avatar
gilbertlee-amd committed
393
  }
Gilbert Lee's avatar
Gilbert Lee committed
394

gilbertlee-amd's avatar
gilbertlee-amd committed
395
396
397
  // Report timings
  totalCpuTime = totalCpuTime / (1.0 * numTimedIterations) * 1000;
  double totalBandwidthGbs = (totalBytesTransferred / 1.0E6) / totalCpuTime;
gilbertlee-amd's avatar
gilbertlee-amd committed
398
399
  if (totalBandwidthCpu) *totalBandwidthCpu = totalBandwidthGbs;

gilbertlee-amd's avatar
gilbertlee-amd committed
400
  double maxGpuTime = 0;
Gilbert Lee's avatar
Gilbert Lee committed
401

402
  if (!isSrcCorrect) goto cleanup;
gilbertlee-amd's avatar
gilbertlee-amd committed
403
404
405
406
  if (ev.useSingleStream)
  {
    for (auto& exeInfoPair : transferMap)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
407
408
409
      ExecutorInfo  exeInfo  = exeInfoPair.second;
      ExeType const exeType  = exeInfoPair.first.first;
      int     const exeIndex = exeInfoPair.first.second;
Gilbert Lee's avatar
Gilbert Lee committed
410

gilbertlee-amd's avatar
gilbertlee-amd committed
411
412
      // Compute total time for non GPU executors
      if (exeType != EXE_GPU_GFX)
gilbertlee-amd's avatar
gilbertlee-amd committed
413
414
415
416
417
      {
        exeInfo.totalTime = 0;
        for (auto const& transfer : exeInfo.transfers)
          exeInfo.totalTime = std::max(exeInfo.totalTime, transfer->transferTime);
      }
418

gilbertlee-amd's avatar
gilbertlee-amd committed
419
420
421
      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
422

gilbertlee-amd's avatar
gilbertlee-amd committed
423
424
      if (verbose && !ev.outputToCsv)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
425
426
        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
427
      }
gilbertlee-amd's avatar
gilbertlee-amd committed
428
429
430

      int totalCUs = 0;
      for (auto const& transfer : exeInfo.transfers)
Gilbert Lee's avatar
Gilbert Lee committed
431
      {
Gilbert Lee's avatar
Gilbert Lee committed
432
        double transferDurationMsec = transfer->transferTime / (1.0 * numTimedIterations);
gilbertlee-amd's avatar
gilbertlee-amd committed
433
434
        double transferBandwidthGbs = (transfer->numBytesActual / 1.0E9) / transferDurationMsec * 1000.0f;
        totalCUs += transfer->numSubExecs;
gilbertlee-amd's avatar
gilbertlee-amd committed
435
436

        if (!verbose) continue;
Gilbert Lee's avatar
Gilbert Lee committed
437
438
        if (!ev.outputToCsv)
        {
gilbertlee-amd's avatar
gilbertlee-amd committed
439
          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
440
                 transfer->transferIndex,
gilbertlee-amd's avatar
gilbertlee-amd committed
441
442
                 transferBandwidthGbs,
                 transferDurationMsec,
gilbertlee-amd's avatar
gilbertlee-amd committed
443
444
445
446
447
                 transfer->numBytesActual,
                 transfer->SrcToStr().c_str(),
                 ExeTypeName[transfer->exeType], transfer->exeIndex,
                 transfer->numSubExecs,
                 transfer->DstToStr().c_str());
Gilbert Lee's avatar
Gilbert Lee committed
448
449
450
        }
        else
        {
gilbertlee-amd's avatar
gilbertlee-amd committed
451
452
453
454
455
456
          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
457
                 transferBandwidthGbs, transferDurationMsec,
gilbertlee-amd's avatar
gilbertlee-amd committed
458
459
                 PtrVectorToStr(transfer->srcMem, initOffset).c_str(),
                 PtrVectorToStr(transfer->dstMem, initOffset).c_str());
Gilbert Lee's avatar
Gilbert Lee committed
460
        }
Gilbert Lee's avatar
Gilbert Lee committed
461
      }
gilbertlee-amd's avatar
gilbertlee-amd committed
462
463
464

      if (verbose && ev.outputToCsv)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
465
        printf("%d,ALL,%lu,ALL,%c%02d,ALL,%d,%.3f,%.3f,ALL,ALL\n",
gilbertlee-amd's avatar
gilbertlee-amd committed
466
               testNum, totalBytesTransferred,
gilbertlee-amd's avatar
gilbertlee-amd committed
467
               MemTypeStr[exeType], exeIndex, totalCUs,
gilbertlee-amd's avatar
gilbertlee-amd committed
468
469
               exeBandwidthGbs, exeDurationMsec);
      }
Gilbert Lee's avatar
Gilbert Lee committed
470
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
471
472
473
474
475
476
477
  }
  else
  {
    for (auto const& transferPair : transferList)
    {
      Transfer* transfer = transferPair.second;
      double transferDurationMsec = transfer->transferTime / (1.0 * numTimedIterations);
gilbertlee-amd's avatar
gilbertlee-amd committed
478
      double transferBandwidthGbs = (transfer->numBytesActual / 1.0E9) / transferDurationMsec * 1000.0f;
gilbertlee-amd's avatar
gilbertlee-amd committed
479
480
481
482
      maxGpuTime = std::max(maxGpuTime, transferDurationMsec);
      if (!verbose) continue;
      if (!ev.outputToCsv)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
483
        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
484
485
               transfer->transferIndex,
               transferBandwidthGbs, transferDurationMsec,
gilbertlee-amd's avatar
gilbertlee-amd committed
486
487
488
489
490
               transfer->numBytesActual,
               transfer->SrcToStr().c_str(),
               ExeTypeName[transfer->exeType], transfer->exeIndex,
               transfer->numSubExecs,
               transfer->DstToStr().c_str());
gilbertlee-amd's avatar
gilbertlee-amd committed
491
492
493
      }
      else
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
494
495
496
497
498
499
        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
500
               transferBandwidthGbs, transferDurationMsec,
gilbertlee-amd's avatar
gilbertlee-amd committed
501
502
               PtrVectorToStr(transfer->srcMem, initOffset).c_str(),
               PtrVectorToStr(transfer->dstMem, initOffset).c_str());
gilbertlee-amd's avatar
gilbertlee-amd committed
503
504
505
      }
    }
  }
Gilbert Lee's avatar
Gilbert Lee committed
506

gilbertlee-amd's avatar
gilbertlee-amd committed
507
508
509
  // Display aggregate statistics
  if (verbose)
  {
Gilbert Lee's avatar
Gilbert Lee committed
510
    if (!ev.outputToCsv)
Gilbert Lee's avatar
Gilbert Lee committed
511
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
512
      printf(" Aggregate (CPU)  | %7.3f GB/s | %8.3f ms | %12lu bytes | Overhead: %.3f ms\n",
513
             totalBandwidthGbs, totalCpuTime, totalBytesTransferred, totalCpuTime - maxGpuTime);
Gilbert Lee's avatar
Gilbert Lee committed
514
515
516
    }
    else
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
517
      printf("%d,ALL,%lu,ALL,ALL,ALL,ALL,%.3f,%.3f,ALL,ALL\n",
518
             testNum, totalBytesTransferred, totalBandwidthGbs, totalCpuTime);
Gilbert Lee's avatar
Gilbert Lee committed
519
520
    }
  }
Gilbert Lee's avatar
Gilbert Lee committed
521

Gilbert Lee's avatar
Gilbert Lee committed
522
  // Release GPU memory
523
cleanup:
Gilbert Lee's avatar
Gilbert Lee committed
524
525
  for (auto exeInfoPair : transferMap)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
526
527
528
529
    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
530
531
    for (auto& transfer : exeInfo.transfers)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
532
533
534
535
536
537
538
539
540
541
542
      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
543
544
    }

gilbertlee-amd's avatar
gilbertlee-amd committed
545
    if (IsGpuType(exeType))
Gilbert Lee's avatar
Gilbert Lee committed
546
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
547
548
      int const numStreams = (int)exeInfo.streams.size();
      for (int i = 0; i < numStreams; ++i)
Gilbert Lee's avatar
Gilbert Lee committed
549
      {
Gilbert Lee's avatar
Gilbert Lee committed
550
551
552
        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
553
      }
gilbertlee-amd's avatar
gilbertlee-amd committed
554
555
556
557
558

      if (exeType == EXE_GPU_GFX)
      {
        DeallocateMemory(MEM_GPU, exeInfo.subExecParamGpu);
      }
Gilbert Lee's avatar
Gilbert Lee committed
559
560
561
562
563
564
    }
  }
}

void DisplayUsage(char const* cmdName)
{
Gilbert Lee's avatar
Gilbert Lee committed
565
  printf("TransferBench v%s\n", TB_VERSION);
Gilbert Lee's avatar
Gilbert Lee committed
566
567
568
569
570
571
572
573
574
575
576
577
578
  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
579
  printf("          - Filename of configFile containing Transfers to execute (see example.cfg for format)\n");
gilbertlee-amd's avatar
gilbertlee-amd committed
580
581
582
  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");
583
584
585
586
587
  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
588
589
  printf("              a2a          - GPU All-To-All benchmark\n");
  printf("                             - 3rd optional arg: # of SubExecs to use\n");
Gilbert Lee's avatar
Gilbert Lee committed
590
  printf("  N     : (Optional) Number of bytes to copy per Transfer.\n");
Gilbert Lee's avatar
Gilbert Lee committed
591
  printf("          If not specified, defaults to %lu bytes. Must be a multiple of 4 bytes\n",
Gilbert Lee's avatar
Gilbert Lee committed
592
         DEFAULT_BYTES_PER_TRANSFER);
Gilbert Lee's avatar
Gilbert Lee committed
593
594
595
596
597
598
599
  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
600
int RemappedIndex(int const origIdx, bool const isCpuType)
Gilbert Lee's avatar
Gilbert Lee committed
601
{
602
603
  static std::vector<int> remappingCpu;
  static std::vector<int> remappingGpu;
Gilbert Lee's avatar
Gilbert Lee committed
604

605
606
607
608
609
610
611
612
  // 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
613

614
615
  // Build remappingGpu on first use
  if (remappingGpu.empty())
Gilbert Lee's avatar
Gilbert Lee committed
616
617
618
  {
    int numGpuDevices;
    HIP_CALL(hipGetDeviceCount(&numGpuDevices));
619
    remappingGpu.resize(numGpuDevices);
Gilbert Lee's avatar
Gilbert Lee committed
620
621
622
623

    int const usePcieIndexing = getenv("USE_PCIE_INDEX") ? atoi(getenv("USE_PCIE_INDEX")) : 0;
    if (!usePcieIndexing)
    {
624
      // For HIP-based indexing no remappingGpu is necessary
Gilbert Lee's avatar
Gilbert Lee committed
625
      for (int i = 0; i < numGpuDevices; ++i)
626
        remappingGpu[i] = i;
Gilbert Lee's avatar
Gilbert Lee committed
627
628
629
630
631
632
633
634
635
636
637
638
639
640
    }
    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)
641
        remappingGpu[i] = mapping[i].second;
Gilbert Lee's avatar
Gilbert Lee committed
642
643
    }
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
644
  return isCpuType ? remappingCpu[origIdx] : remappingGpu[origIdx];
Gilbert Lee's avatar
Gilbert Lee committed
645
646
647
648
}

void DisplayTopology(bool const outputToCsv)
{
649

650
  int numCpuDevices = numa_num_configured_nodes();
Gilbert Lee's avatar
Gilbert Lee committed
651
652
653
654
655
  int numGpuDevices;
  HIP_CALL(hipGetDeviceCount(&numGpuDevices));

  if (outputToCsv)
  {
656
    printf("NumCpus,%d\n", numCpuDevices);
Gilbert Lee's avatar
Gilbert Lee committed
657
    printf("NumGpus,%d\n", numGpuDevices);
658
659
660
  }
  else
  {
661
662
    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);
663
664
665
666
667
668
669
670
  }

  // Print out detected CPU topology
  if (outputToCsv)
  {
    printf("NUMA");
    for (int j = 0; j < numCpuDevices; j++)
      printf(",NUMA%02d", j);
671
    printf(",# CPUs,ClosestGPUs,ActualNode\n");
672
673
674
  }
  else
  {
675
    printf("            |");
676
    for (int j = 0; j < numCpuDevices; j++)
677
678
679
680
      printf("NUMA %02d|", j);
    printf(" #Cpus | Closest GPU(s)\n");

    printf("------------+");
681
    for (int j = 0; j <= numCpuDevices; j++)
682
683
      printf("-------+");
    printf("---------------\n");
684
685
686
687
  }

  for (int i = 0; i < numCpuDevices; i++)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
688
    int nodeI = RemappedIndex(i, true);
689
    printf("NUMA %02d (%02d)%s", i, nodeI, outputToCsv ? "," : "|");
690
691
    for (int j = 0; j < numCpuDevices; j++)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
692
      int nodeJ = RemappedIndex(j, true);
693
      int numaDist = numa_distance(nodeI, nodeJ);
694
      if (outputToCsv)
gilbertlee-amd's avatar
gilbertlee-amd committed
695
        printf("%d,", numaDist);
696
      else
697
        printf(" %5d |", numaDist);
698
699
700
701
    }

    int numCpus = 0;
    for (int j = 0; j < numa_num_configured_cpus(); j++)
702
      if (numa_node_of_cpu(j) == nodeI) numCpus++;
703
704
705
    if (outputToCsv)
      printf("%d,", numCpus);
    else
706
      printf(" %5d | ", numCpus);
707

708
#if !defined(__NVCC__)
709
710
711
    bool isFirst = true;
    for (int j = 0; j < numGpuDevices; j++)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
712
      if (GetClosestNumaNode(RemappedIndex(j, false)) == i)
713
714
      {
        if (isFirst) isFirst = false;
gilbertlee-amd's avatar
gilbertlee-amd committed
715
716
        else printf(",");
        printf("%d", j);
717
718
      }
    }
719
#endif
720
721
722
723
    printf("\n");
  }
  printf("\n");

724
725
726
727
728
#if defined(__NVCC__)
  // No further topology detection done for NVIDIA platforms
  return;
#endif

729
730
731
  // Print out detected GPU topology
  if (outputToCsv)
  {
Gilbert Lee's avatar
Gilbert Lee committed
732
733
734
735
736
737
738
    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
739
740
741
742
743
744
745
746
747
748
    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
749
750
751
    printf("        |");
    for (int j = 0; j < numGpuDevices; j++)
      printf(" GPU %02d |", j);
gilbertlee-amd's avatar
gilbertlee-amd committed
752
    printf(" PCIe Bus ID  | #CUs | Closest NUMA\n");
Gilbert Lee's avatar
Gilbert Lee committed
753
754
    for (int j = 0; j <= numGpuDevices; j++)
      printf("--------+");
gilbertlee-amd's avatar
gilbertlee-amd committed
755
    printf("--------------+------+-------------\n");
Gilbert Lee's avatar
Gilbert Lee committed
756
757
  }

758
#if !defined(__NVCC__)
Gilbert Lee's avatar
Gilbert Lee committed
759
760
761
  char pciBusId[20];
  for (int i = 0; i < numGpuDevices; i++)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
762
    int const deviceIdx = RemappedIndex(i, false);
Gilbert Lee's avatar
Gilbert Lee committed
763
764
765
766
767
768
769
770
771
772
773
774
775
    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
776
777
        HIP_CALL(hipExtGetLinkTypeAndHopCount(deviceIdx,
                                              RemappedIndex(j, false),
Gilbert Lee's avatar
Gilbert Lee committed
778
779
780
781
782
783
784
785
786
787
788
                                              &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
789
790
791
792
793
    HIP_CALL(hipDeviceGetPCIBusId(pciBusId, 20, deviceIdx));

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

Gilbert Lee's avatar
Gilbert Lee committed
794
    if (outputToCsv)
gilbertlee-amd's avatar
gilbertlee-amd committed
795
      printf("%s,%d,%d\n", pciBusId, numDeviceCUs, GetClosestNumaNode(deviceIdx));
Gilbert Lee's avatar
Gilbert Lee committed
796
    else
gilbertlee-amd's avatar
gilbertlee-amd committed
797
      printf(" %11s | %4d | %d\n", pciBusId, numDeviceCUs, GetClosestNumaNode(deviceIdx));
Gilbert Lee's avatar
Gilbert Lee committed
798
  }
799
#endif
Gilbert Lee's avatar
Gilbert Lee committed
800
801
}

gilbertlee-amd's avatar
gilbertlee-amd committed
802
803
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
804
805
{
  char typeChar;
gilbertlee-amd's avatar
gilbertlee-amd committed
806
807
  int offset = 0, devIndex, inc;
  bool found = false;
Gilbert Lee's avatar
Gilbert Lee committed
808

gilbertlee-amd's avatar
gilbertlee-amd committed
809
810
811
  memTypes.clear();
  memIndices.clear();
  while (sscanf(token.c_str() + offset, " %c %d%n", &typeChar, &devIndex, &inc) == 2)
Gilbert Lee's avatar
Gilbert Lee committed
812
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
813
814
815
816
    offset += inc;
    MemType memType = CharToMemType(typeChar);

    if (IsCpuType(memType) && (devIndex < 0 || devIndex >= numCpus))
Gilbert Lee's avatar
Gilbert Lee committed
817
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
818
      printf("[ERROR] CPU index must be between 0 and %d (instead of %d)\n", numCpus-1, devIndex);
Gilbert Lee's avatar
Gilbert Lee committed
819
820
      exit(1);
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
821
    if (IsGpuType(memType) && (devIndex < 0 || devIndex >= numGpus))
Gilbert Lee's avatar
Gilbert Lee committed
822
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
823
      printf("[ERROR] GPU index must be between 0 and %d (instead of %d)\n", numGpus-1, devIndex);
Gilbert Lee's avatar
Gilbert Lee committed
824
825
      exit(1);
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
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

    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
862
863
864
865
    exit(1);
  }
}

Gilbert Lee's avatar
Gilbert Lee committed
866
// Helper function to parse a list of Transfer definitions
Gilbert Lee's avatar
Gilbert Lee committed
867
void ParseTransfers(char* line, int numCpus, int numGpus, std::vector<Transfer>& transfers)
Gilbert Lee's avatar
Gilbert Lee committed
868
869
870
871
872
{
  // 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
873
  transfers.clear();
Gilbert Lee's avatar
Gilbert Lee committed
874

Gilbert Lee's avatar
Gilbert Lee committed
875
  int numTransfers = 0;
Gilbert Lee's avatar
Gilbert Lee committed
876
  std::istringstream iss(line);
Gilbert Lee's avatar
Gilbert Lee committed
877
  iss >> numTransfers;
Gilbert Lee's avatar
Gilbert Lee committed
878
879
880
881
882
  if (iss.fail()) return;

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

gilbertlee-amd's avatar
gilbertlee-amd committed
884
  // If numTransfers < 0, read 5-tuple (srcMem, exeMem, dstMem, #CUs, #Bytes)
Gilbert Lee's avatar
Gilbert Lee committed
885
  // otherwise read triples (srcMem, exeMem, dstMem)
gilbertlee-amd's avatar
gilbertlee-amd committed
886
  bool const advancedMode = (numTransfers < 0);
Gilbert Lee's avatar
Gilbert Lee committed
887
888
  numTransfers = abs(numTransfers);

gilbertlee-amd's avatar
gilbertlee-amd committed
889
  int numSubExecs;
gilbertlee-amd's avatar
gilbertlee-amd committed
890
  if (!advancedMode)
Gilbert Lee's avatar
Gilbert Lee committed
891
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
892
893
    iss >> numSubExecs;
    if (numSubExecs <= 0 || iss.fail())
Gilbert Lee's avatar
Gilbert Lee committed
894
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
895
      printf("Parsing error: Number of blocks to use (%d) must be greater than 0\n", numSubExecs);
Gilbert Lee's avatar
Gilbert Lee committed
896
897
898
899
      exit(1);
    }
  }

gilbertlee-amd's avatar
gilbertlee-amd committed
900
  size_t numBytes = 0;
Gilbert Lee's avatar
Gilbert Lee committed
901
902
903
  for (int i = 0; i < numTransfers; i++)
  {
    Transfer transfer;
gilbertlee-amd's avatar
gilbertlee-amd committed
904
    transfer.numBytes = 0;
gilbertlee-amd's avatar
gilbertlee-amd committed
905
    transfer.numBytesActual = 0;
gilbertlee-amd's avatar
gilbertlee-amd committed
906
    if (!advancedMode)
Gilbert Lee's avatar
Gilbert Lee committed
907
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
908
909
910
911
912
913
914
915
916
917
      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
918
      iss >> srcMem >> exeMem >> dstMem >> numSubExecs >> numBytesToken;
gilbertlee-amd's avatar
gilbertlee-amd committed
919
920
921
922
923
924
925
926
927
928
929
      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
930
      switch (toupper(units))
gilbertlee-amd's avatar
gilbertlee-amd committed
931
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
932
933
934
      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
935
      }
Gilbert Lee's avatar
Gilbert Lee committed
936
    }
Gilbert Lee's avatar
Gilbert Lee committed
937

gilbertlee-amd's avatar
gilbertlee-amd committed
938
939
940
941
942
943
944
945
946
947
948
949
950
951
952
953
954
955
956
    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
957
    transfer.numBytes = numBytes;
Gilbert Lee's avatar
Gilbert Lee committed
958
    transfers.push_back(transfer);
Gilbert Lee's avatar
Gilbert Lee committed
959
960
961
962
963
964
965
966
967
968
969
970
971
  }
}

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
972
973
974
975
976
977
978
  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
979
980
981
982
983
984
985
986
987
}

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
988
  *memPtr = nullptr;
gilbertlee-amd's avatar
gilbertlee-amd committed
989
  if (IsCpuType(memType))
Gilbert Lee's avatar
Gilbert Lee committed
990
991
  {
    // Set numa policy prior to call to hipHostMalloc
992
    numa_set_preferred(devIndex);
Gilbert Lee's avatar
Gilbert Lee committed
993
994
995
996

    // Allocate host-pinned memory (should respect NUMA mem policy)
    if (memType == MEM_CPU_FINE)
    {
997
998
999
1000
#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
1001
      HIP_CALL(hipHostMalloc((void **)memPtr, numBytes, hipHostMallocNumaUser));
1002
#endif
Gilbert Lee's avatar
Gilbert Lee committed
1003
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
1004
    else if (memType == MEM_CPU)
Gilbert Lee's avatar
Gilbert Lee committed
1005
    {
1006
1007
1008
#if defined (__NVCC__)
      if (hipHostMalloc((void **)memPtr, numBytes, 0) != hipSuccess)
#else
1009
      if (hipHostMalloc((void **)memPtr, numBytes, hipHostMallocNumaUser | hipHostMallocNonCoherent) != hipSuccess)
1010
#endif
1011
1012
1013
1014
      {
        printf("[ERROR] Unable to allocate non-coherent host memory on NUMA node %d\n", devIndex);
        exit(1);
      }
Gilbert Lee's avatar
Gilbert Lee committed
1015
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
1016
1017
1018
1019
    else if (memType == MEM_CPU_UNPINNED)
    {
      *memPtr = numa_alloc_onnode(numBytes, devIndex);
    }
Gilbert Lee's avatar
Gilbert Lee committed
1020
1021

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

    // Reset to default numa mem policy
1026
    numa_set_preferred(-1);
Gilbert Lee's avatar
Gilbert Lee committed
1027
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
1028
  else if (IsGpuType(memType))
Gilbert Lee's avatar
Gilbert Lee committed
1029
  {
1030
1031
1032
1033
1034
1035
1036
1037
    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)
    {
1038
#if defined (__NVCC__)
1039
1040
      printf("[ERROR] Fine-grained GPU memory not supported on NVIDIA platform\n");
      exit(1);
1041
#else
1042
1043
      HIP_CALL(hipSetDevice(devIndex));

gilbertlee-amd's avatar
gilbertlee-amd committed
1044
1045
1046
1047
1048
1049
      // 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));
1050
#endif
1051
1052
    }
    HIP_CALL(hipMemset(*memPtr, 0, numBytes));
gilbertlee-amd's avatar
gilbertlee-amd committed
1053
    HIP_CALL(hipDeviceSynchronize());
Gilbert Lee's avatar
Gilbert Lee committed
1054
1055
1056
1057
1058
1059
1060
1061
  }
  else
  {
    printf("[ERROR] Unsupported memory type %d\n", memType);
    exit(1);
  }
}

gilbertlee-amd's avatar
gilbertlee-amd committed
1062
void DeallocateMemory(MemType memType, void* memPtr, size_t const bytes)
Gilbert Lee's avatar
Gilbert Lee committed
1063
1064
1065
{
  if (memType == MEM_CPU || memType == MEM_CPU_FINE)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1066
1067
1068
1069
1070
    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
1071
1072
    HIP_CALL(hipHostFree(memPtr));
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
1073
1074
  else if (memType == MEM_CPU_UNPINNED)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1075
1076
1077
1078
1079
    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
1080
1081
    numa_free(memPtr, bytes);
  }
Gilbert Lee's avatar
Gilbert Lee committed
1082
1083
  else if (memType == MEM_GPU || memType == MEM_GPU_FINE)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1084
1085
1086
1087
1088
    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
1089
1090
1091
1092
1093
1094
1095
1096
1097
1098
1099
1100
1101
1102
1103
1104
1105
1106
1107
1108
1109
1110
1111
1112
1113
1114
1115
1116
1117
1118
1119
1120
1121
1122
1123
1124
1125
1126
1127
1128
1129
1130
    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);
  }
}

1131
void RunTransfer(EnvVars const& ev, int const iteration,
Gilbert Lee's avatar
Gilbert Lee committed
1132
                 ExecutorInfo& exeInfo, int const transferIdx)
Gilbert Lee's avatar
Gilbert Lee committed
1133
{
gilbertlee-amd's avatar
gilbertlee-amd committed
1134
  Transfer* transfer = exeInfo.transfers[transferIdx];
Gilbert Lee's avatar
Gilbert Lee committed
1135

gilbertlee-amd's avatar
gilbertlee-amd committed
1136
  if (transfer->exeType == EXE_GPU_GFX)
Gilbert Lee's avatar
Gilbert Lee committed
1137
1138
  {
    // Switch to executing GPU
gilbertlee-amd's avatar
gilbertlee-amd committed
1139
    int const exeIndex = RemappedIndex(transfer->exeIndex, false);
Gilbert Lee's avatar
Gilbert Lee committed
1140
1141
    HIP_CALL(hipSetDevice(exeIndex));

Gilbert Lee's avatar
Gilbert Lee committed
1142
1143
1144
    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
1145

gilbertlee-amd's avatar
gilbertlee-amd committed
1146
1147
1148
1149
    // 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;
1150
1151
1152
1153
1154
#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
1155
1156
1157
1158
1159
1160
    hipExtLaunchKernelGGL(GpuKernelTable[ev.gpuKernel],
                          dim3(numBlocksToRun, 1, 1),
                          dim3(BLOCKSIZE, 1, 1),
                          ev.sharedMemBytes, stream,
                          startEvent, stopEvent,
                          0, transfer->subExecParamGpuPtr);
1161
#endif
Gilbert Lee's avatar
Gilbert Lee committed
1162
1163
    // 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
1164
    HIP_CALL(hipStreamSynchronize(stream));
Gilbert Lee's avatar
Gilbert Lee committed
1165
1166
1167
1168

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

Gilbert Lee's avatar
Gilbert Lee committed
1172
1173
      if (ev.useSingleStream)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
1174
        // Figure out individual timings for Transfers that were all launched together
gilbertlee-amd's avatar
gilbertlee-amd committed
1175
        for (Transfer* currTransfer : exeInfo.transfers)
Gilbert Lee's avatar
Gilbert Lee committed
1176
        {
gilbertlee-amd's avatar
gilbertlee-amd committed
1177
1178
1179
          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
1180
          {
gilbertlee-amd's avatar
gilbertlee-amd committed
1181
1182
            minStartCycle = std::min(minStartCycle, currTransfer->subExecParamGpuPtr[i].startCycle);
            maxStopCycle  = std::max(maxStopCycle,  currTransfer->subExecParamGpuPtr[i].stopCycle);
Gilbert Lee's avatar
Gilbert Lee committed
1183
          }
Gilbert Lee's avatar
Gilbert Lee committed
1184
1185
          int const wallClockRate = GetWallClockRate(exeIndex);
          double iterationTimeMs = (maxStopCycle - minStartCycle) / (double)(wallClockRate);
gilbertlee-amd's avatar
gilbertlee-amd committed
1186
          currTransfer->transferTime += iterationTimeMs;
Gilbert Lee's avatar
Gilbert Lee committed
1187
        }
Gilbert Lee's avatar
Gilbert Lee committed
1188
1189
1190
1191
        exeInfo.totalTime += gpuDeltaMsec;
      }
      else
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
1192
        transfer->transferTime += gpuDeltaMsec;
Gilbert Lee's avatar
Gilbert Lee committed
1193
1194
1195
      }
    }
  }
gilbertlee-amd's avatar
gilbertlee-amd committed
1196
1197
1198
1199
1200
1201
1202
1203
1204
1205
1206
1207
1208
1209
1210
1211
1212
1213
1214
1215
1216
1217
1218
1219
1220
1221
1222
1223
1224
1225
1226
1227
1228
1229
  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;
    }
  }
  else if (transfer->exeType == EXE_CPU) // CPU execution agent
Gilbert Lee's avatar
Gilbert Lee committed
1230
1231
  {
    // Force this thread and all child threads onto correct NUMA node
gilbertlee-amd's avatar
gilbertlee-amd committed
1232
    int const exeIndex = RemappedIndex(transfer->exeIndex, true);
1233
    if (numa_run_on_node(exeIndex))
Gilbert Lee's avatar
Gilbert Lee committed
1234
    {
1235
      printf("[ERROR] Unable to set CPU to NUMA node %d\n", exeIndex);
Gilbert Lee's avatar
Gilbert Lee committed
1236
1237
1238
1239
1240
1241
1242
      exit(1);
    }

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

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

gilbertlee-amd's avatar
gilbertlee-amd committed
1243
1244
1245
    // 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
1246
1247

    // Wait for child-threads to finish
gilbertlee-amd's avatar
gilbertlee-amd committed
1248
    for (int i = 0; i < transfer->numSubExecs; ++i)
Gilbert Lee's avatar
Gilbert Lee committed
1249
1250
1251
1252
1253
1254
      childThreads[i].join();

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

    // Record time if not a warmup iteration
    if (iteration >= 0)
gilbertlee-amd's avatar
gilbertlee-amd committed
1255
      transfer->transferTime += (std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() * 1000.0);
Gilbert Lee's avatar
Gilbert Lee committed
1256
1257
1258
  }
}

gilbertlee-amd's avatar
gilbertlee-amd committed
1259
void RunPeerToPeerBenchmarks(EnvVars const& ev, size_t N)
Gilbert Lee's avatar
Gilbert Lee committed
1260
{
gilbertlee-amd's avatar
gilbertlee-amd committed
1261
1262
  ev.DisplayP2PBenchmarkEnvVars();

Gilbert Lee's avatar
Gilbert Lee committed
1263
  // Collect the number of available CPUs/GPUs on this machine
gilbertlee-amd's avatar
gilbertlee-amd committed
1264
1265
  int const numCpus    = ev.numCpuDevices;
  int const numGpus    = ev.numGpuDevices;
Gilbert Lee's avatar
Gilbert Lee committed
1266
1267
1268
1269
1270
1271
1272
1273
1274
1275
1276
1277
1278
  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++)
  {
    // Print header
    if (!ev.outputToCsv)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1279
1280
1281
1282
1283
      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");

1284
1285
1286
1287
1288
1289
1290
1291
1292
1293
1294
      if (isBidirectional)
      {
        printf("%12s", "SRC\\DST");
      }
      else
      {
        if (ev.useRemoteRead)
          printf("%12s", "SRC\\EXE+DST");
        else
          printf("%12s", "SRC+EXE\\DST");
      }
gilbertlee-amd's avatar
gilbertlee-amd committed
1295
1296
      for (int i = 0; i < numCpus; i++) printf("%7s %02d", "CPU", i);
      for (int i = 0; i < numGpus; i++) printf("%7s %02d", "GPU", i);
Gilbert Lee's avatar
Gilbert Lee committed
1297
1298
1299
1300
1301
1302
      printf("\n");
    }

    // Loop over all possible src/dst pairs
    for (int src = 0; src < numDevices; src++)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1303
1304
1305
      MemType const srcType  = (src < numCpus ? MEM_CPU : MEM_GPU);
      int     const srcIndex = (srcType == MEM_CPU ? src : src - numCpus);

Gilbert Lee's avatar
Gilbert Lee committed
1306
      if (!ev.outputToCsv)
1307
        printf("%9s %02d", (srcType == MEM_CPU) ? "CPU" : "GPU", srcIndex);
gilbertlee-amd's avatar
gilbertlee-amd committed
1308

Gilbert Lee's avatar
Gilbert Lee committed
1309
1310
      for (int dst = 0; dst < numDevices; dst++)
      {
gilbertlee-amd's avatar
gilbertlee-amd committed
1311
1312
1313
1314
        MemType const dstType  = (dst < numCpus ? MEM_CPU : MEM_GPU);
        int     const dstIndex = (dstType == MEM_CPU ? dst : dst - numCpus);

        double bandwidth = GetPeakBandwidth(ev, N, isBidirectional, srcType, srcIndex, dstType, dstIndex);
Gilbert Lee's avatar
Gilbert Lee committed
1315
1316
1317
1318
1319
1320
1321
1322
1323
        if (!ev.outputToCsv)
        {
          if (bandwidth == 0)
            printf("%10s", "N/A");
          else
            printf("%10.2f", bandwidth);
        }
        else
        {
gilbertlee-amd's avatar
gilbertlee-amd committed
1324
1325
1326
          printf("%s %02d,%s %02d,%s,%s,%s,%.2f,%lu\n",
                 srcType == MEM_CPU ? "CPU" : "GPU", srcIndex,
                 dstType == MEM_CPU ? "CPU" : "GPU", dstIndex,
Gilbert Lee's avatar
Gilbert Lee committed
1327
                 isBidirectional ? "bidirectional" : "unidirectional",
gilbertlee-amd's avatar
gilbertlee-amd committed
1328
1329
                 ev.useRemoteRead ? "Remote" : "Local",
                 ev.useDmaCopy ? "DMA" : "GFX",
Gilbert Lee's avatar
Gilbert Lee committed
1330
1331
1332
1333
1334
1335
1336
1337
1338
1339
1340
                 bandwidth,
                 N * sizeof(float));
        }
        fflush(stdout);
      }
      if (!ev.outputToCsv) printf("\n");
    }
    if (!ev.outputToCsv) printf("\n");
  }
}

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
1371
1372
1373
1374
1375
1376
1377
1378
1379
1380
1381
1382
1383
1384
1385
1386
1387
1388
1389
1390
1391
1392
1393
1394
1395
1396
1397
1398
1399
1400
1401
1402
1403
1404
1405
1406
1407
1408
1409
1410
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
1411
1412
1413
1414
1415
1416
1417
1418
1419
1420
1421
1422
1423
1424
1425
1426
1427
1428
1429
1430
1431
1432
1433
1434
1435
1436
1437
1438
1439
1440
1441
1442
1443
1444
1445
1446
1447
1448
1449
1450
1451
1452
1453
1454
1455
1456
1457
1458
1459
1460
1461
1462
1463
1464
1465
1466
1467
1468
1469
1470
1471
1472
1473
1474
1475
1476
1477
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
1478
double GetPeakBandwidth(EnvVars const& ev, size_t const N,
Gilbert Lee's avatar
Gilbert Lee committed
1479
                        int     const  isBidirectional,
gilbertlee-amd's avatar
gilbertlee-amd committed
1480
1481
                        MemType const  srcType, int const srcIndex,
                        MemType const  dstType, int const dstIndex)
Gilbert Lee's avatar
Gilbert Lee committed
1482
1483
{
  // Skip bidirectional on same device
gilbertlee-amd's avatar
gilbertlee-amd committed
1484
  if (isBidirectional && srcType == dstType && srcIndex == dstIndex) return 0.0f;
Gilbert Lee's avatar
Gilbert Lee committed
1485

Gilbert Lee's avatar
Gilbert Lee committed
1486
  // Prepare Transfers
gilbertlee-amd's avatar
gilbertlee-amd committed
1487
  std::vector<Transfer> transfers(2);
gilbertlee-amd's avatar
gilbertlee-amd committed
1488
1489
1490
1491
1492
1493
1494
1495
1496
1497
1498
1499
1500
1501
1502
  transfers[0].numBytes = transfers[1].numBytes = N * sizeof(float);

  // SRC -> DST
  transfers[0].numSrcs = transfers[0].numDsts = 1;
  transfers[0].srcType.push_back(srcType);
  transfers[0].dstType.push_back(dstType);
  transfers[0].srcIndex.push_back(srcIndex);
  transfers[0].dstIndex.push_back(dstIndex);

  // DST -> SRC
  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);
Gilbert Lee's avatar
Gilbert Lee committed
1503
1504

  // Either perform (local read + remote write), or (remote read + local write)
gilbertlee-amd's avatar
gilbertlee-amd committed
1505
1506
1507
1508
1509
1510
1511
1512
1513
  ExeType gpuExeType = ev.useDmaCopy ? EXE_GPU_DMA : EXE_GPU_GFX;
  transfers[0].exeType = IsGpuType(ev.useRemoteRead ? dstType : srcType) ? gpuExeType : EXE_CPU;
  transfers[1].exeType = IsGpuType(ev.useRemoteRead ? srcType : dstType) ? gpuExeType : EXE_CPU;
  transfers[0].exeIndex = (ev.useRemoteRead ? dstIndex : srcIndex);
  transfers[1].exeIndex = (ev.useRemoteRead ? srcIndex : dstIndex);
  transfers[0].numSubExecs = IsGpuType(transfers[0].exeType) ? ev.numGpuSubExecs : ev.numCpuSubExecs;
  transfers[1].numSubExecs = IsGpuType(transfers[0].exeType) ? ev.numGpuSubExecs : ev.numCpuSubExecs;

  // Remove (DST->SRC) if not bidirectional
gilbertlee-amd's avatar
gilbertlee-amd committed
1514
  transfers.resize(isBidirectional + 1);
Gilbert Lee's avatar
Gilbert Lee committed
1515

1516
1517
1518
  // Abort if executing on NUMA node with no CPUs
  for (int i = 0; i <= isBidirectional; i++)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1519
    if (transfers[i].exeType == EXE_CPU && ev.numCpusPerNuma[transfers[i].exeIndex] == 0)
1520
      return 0;
1521
1522
1523
1524
1525
1526

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

gilbertlee-amd's avatar
gilbertlee-amd committed
1529
  ExecuteTransfers(ev, 0, N, transfers, false);
Gilbert Lee's avatar
Gilbert Lee committed
1530
1531
1532
1533
1534

  // Collect aggregate bandwidth
  double totalBandwidth = 0;
  for (int i = 0; i <= isBidirectional; i++)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1535
    double transferDurationMsec = transfers[i].transferTime / (1.0 * ev.numIterations);
gilbertlee-amd's avatar
gilbertlee-amd committed
1536
    double transferBandwidthGbs = (transfers[i].numBytesActual / 1.0E9) / transferDurationMsec * 1000.0f;
Gilbert Lee's avatar
Gilbert Lee committed
1537
    totalBandwidth += transferBandwidthGbs;
Gilbert Lee's avatar
Gilbert Lee committed
1538
1539
1540
1541
  }
  return totalBandwidth;
}

gilbertlee-amd's avatar
gilbertlee-amd committed
1542
void Transfer::PrepareSubExecParams(EnvVars const& ev)
Gilbert Lee's avatar
Gilbert Lee committed
1543
{
gilbertlee-amd's avatar
gilbertlee-amd committed
1544
1545
1546
1547
1548
1549
1550
  // 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
1551

gilbertlee-amd's avatar
gilbertlee-amd committed
1552
  // In some cases, there may not be enough data for all subExectors
1553
  int const maxSubExecToUse = std::min((size_t)(N + targetMultiple - 1) / targetMultiple, (size_t)this->numSubExecs);
gilbertlee-amd's avatar
gilbertlee-amd committed
1554
1555
  this->subExecParam.clear();
  this->subExecParam.resize(this->numSubExecs);
Gilbert Lee's avatar
Gilbert Lee committed
1556
1557

  size_t assigned = 0;
gilbertlee-amd's avatar
gilbertlee-amd committed
1558
1559
1560
1561
1562
1563
1564
1565
1566
1567
1568
1569
1570
1571
1572
1573
1574
1575
1576
1577
  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
1578

gilbertlee-amd's avatar
gilbertlee-amd committed
1579
1580
1581
    p.startCycle = 0;
    p.stopCycle  = 0;
    assigned += p.N;
Gilbert Lee's avatar
Gilbert Lee committed
1582
1583
  }

Gilbert Lee's avatar
Gilbert Lee committed
1584
  this->transferTime = 0.0;
Gilbert Lee's avatar
Gilbert Lee committed
1585
1586
}

gilbertlee-amd's avatar
gilbertlee-amd committed
1587
1588
1589
1590
1591
1592
1593
1594
1595
1596
1597
1598
1599
1600
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)
1601
        buffer[i] = PrepSrcValue(bufferIdx, i);
gilbertlee-amd's avatar
gilbertlee-amd committed
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
    }
  }
  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];
          }
        }
      }
    }
  }
}

1631
bool Transfer::PrepareSrc(EnvVars const& ev)
gilbertlee-amd's avatar
gilbertlee-amd committed
1632
{
1633
  if (this->numSrcs == 0) return true;
gilbertlee-amd's avatar
gilbertlee-amd committed
1634
1635
1636
1637
1638
1639
  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
1640
    float* srcPtr = this->srcMem[srcIdx] + initOffset;
1641
    PrepareReference(ev, reference, srcIdx);
gilbertlee-amd's avatar
gilbertlee-amd committed
1642
1643
1644

    // Initialize source memory array with reference pattern
    if (IsGpuType(this->srcType[srcIdx]))
1645
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1646
1647
1648
1649
1650
1651
      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));
1652
1653
      HIP_CALL(hipDeviceSynchronize());
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
1654
    else if (IsCpuType(this->srcType[srcIdx]))
1655
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1656
      memcpy(srcPtr, reference.data(), this->numBytesActual);
1657
    }
1658
1659

    // Perform check just to make sure that data has been copied properly
gilbertlee-amd's avatar
gilbertlee-amd committed
1660
    float* srcCheckPtr = srcPtr;
1661
    std::vector<float> srcCopy(N);
gilbertlee-amd's avatar
gilbertlee-amd committed
1662
1663
1664
1665
1666
1667
1668
1669
1670
    if (IsGpuType(this->srcType[srcIdx]))
    {
      if (!ev.validateDirect)
      {
        HIP_CALL(hipMemcpy(srcCopy.data(), srcPtr, this->numBytesActual, hipMemcpyDefault));
        HIP_CALL(hipDeviceSynchronize());
        srcCheckPtr = srcCopy.data();
      }
    }
1671
1672
1673

    for (size_t i = 0; i < N; ++i)
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1674
      if (reference[i] != srcCheckPtr[i])
1675
1676
      {
        printf("\n[ERROR] Unexpected mismatch at index %lu of source array %d:\n", i, srcIdx);
1677
1678
1679
#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
1680
               srcIdx, srcCheckPtr[i], *(unsigned int*)&srcCheckPtr[i], val, *(unsigned int*)&val);
1681
#else
gilbertlee-amd's avatar
gilbertlee-amd committed
1682
        printf("[ERROR] SRC %02d   value: %10.5f [%08X]\n", srcIdx, srcCheckPtr[i], *(unsigned int*)&srcCheckPtr[i]);
1683
#endif
1684
1685
1686
1687
1688
1689
1690
1691
1692
        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);
1693
        return false;
1694
1695
      }
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
1696
  }
1697
  return true;
gilbertlee-amd's avatar
gilbertlee-amd committed
1698
1699
1700
1701
1702
1703
1704
1705
1706
1707
1708
1709
1710
1711
1712
}

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;
1713
    if (IsCpuType(this->dstType[dstIdx]) || ev.validateDirect)
gilbertlee-amd's avatar
gilbertlee-amd committed
1714
1715
1716
1717
1718
    {
      output = this->dstMem[dstIdx] + initOffset;
    }
    else
    {
gilbertlee-amd's avatar
gilbertlee-amd committed
1719
1720
      int const deviceIdx = RemappedIndex(this->dstIndex[dstIdx], false);
      HIP_CALL(hipSetDevice(deviceIdx));
gilbertlee-amd's avatar
gilbertlee-amd committed
1721
      HIP_CALL(hipMemcpy(hostBuffer.data(), this->dstMem[dstIdx] + initOffset, this->numBytesActual, hipMemcpyDefault));
gilbertlee-amd's avatar
gilbertlee-amd committed
1722
      HIP_CALL(hipDeviceSynchronize());
gilbertlee-amd's avatar
gilbertlee-amd committed
1723
1724
1725
1726
1727
1728
1729
      output = hostBuffer.data();
    }

    for (size_t i = 0; i < N; ++i)
    {
      if (reference[i] != output[i])
      {
1730
1731
1732
1733
1734
        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));
1735
1736
1737
1738
1739
#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
1740
          printf("[ERROR] SRC %02d   value: %10.5f [%08X]\n", srcIdx, srcVal, *(unsigned int*)&srcVal);
1741
#endif
1742
        }
1743
        printf("[ERROR] EXPECTED value: %10.5f [%08X]\n", reference[i], *(unsigned int*)&reference[i]);
1744
1745
1746
1747
1748
#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
1749
        printf("[ERROR] DST %02d   value: %10.5f [%08X]\n", dstIdx, output[i], *(unsigned int*)&output[i]);
1750
#endif
gilbertlee-amd's avatar
gilbertlee-amd committed
1751
1752
1753
1754
1755
1756
        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());
1757
1758
        if (!ev.continueOnError)
          exit(1);
1759
1760
        else
          break;
gilbertlee-amd's avatar
gilbertlee-amd committed
1761
1762
1763
1764
1765
1766
1767
1768
1769
1770
1771
1772
1773
1774
1775
1776
1777
1778
1779
1780
1781
1782
1783
      }
    }
  }
}

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
1784
1785
1786
1787
1788
1789
1790
1791
1792
1793
1794
1795
1796
// 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++)
    {
1797
1798
1799
1800
1801
1802
#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
1803
1804
1805
1806
1807
      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
1808
      case 940: case 941: case 942: value = 100000; break;
Gilbert Lee's avatar
Gilbert Lee committed
1809
1810
1811
      default:
        printf("Unrecognized GCN arch %d\n", prop.gcnArch);
      }
1812
#endif
Gilbert Lee's avatar
Gilbert Lee committed
1813
1814
1815
1816
1817
      wallClockPerDeviceMhz[i] = value;
    }
  }
  return wallClockPerDeviceMhz[deviceId];
}
Gilbert Lee's avatar
Gilbert Lee committed
1818

gilbertlee-amd's avatar
gilbertlee-amd committed
1819
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
1820
1821
1822
1823
{
  ev.DisplaySweepEnvVars();

  // Compute how many possible Transfers are permitted (unique SRC/EXE/DST triplets)
gilbertlee-amd's avatar
gilbertlee-amd committed
1824
  std::vector<std::pair<ExeType, int>> exeList;
Gilbert Lee's avatar
Gilbert Lee committed
1825
1826
  for (auto exe : ev.sweepExe)
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1827
1828
    ExeType const exeType = CharToExeType(exe);
    if (IsGpuType(exeType))
Gilbert Lee's avatar
Gilbert Lee committed
1829
    {
1830
      for (int exeIndex = 0; exeIndex < ev.numGpuDevices; ++exeIndex)
gilbertlee-amd's avatar
gilbertlee-amd committed
1831
        exeList.push_back(std::make_pair(exeType, exeIndex));
Gilbert Lee's avatar
Gilbert Lee committed
1832
    }
gilbertlee-amd's avatar
gilbertlee-amd committed
1833
    else if (IsCpuType(exeType))
Gilbert Lee's avatar
Gilbert Lee committed
1834
    {
1835
1836
1837
1838
      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
1839
        exeList.push_back(std::make_pair(exeType, exeIndex));
1840
      }
Gilbert Lee's avatar
Gilbert Lee committed
1841
1842
    }
  }
1843
  int numExes = exeList.size();
Gilbert Lee's avatar
Gilbert Lee committed
1844
1845
1846
1847

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

Gilbert Lee's avatar
Gilbert Lee committed
1851
    for (int srcIndex = 0; srcIndex < numDevices; ++srcIndex)
gilbertlee-amd's avatar
gilbertlee-amd committed
1852
      srcList.push_back(std::make_pair(srcType, srcIndex));
Gilbert Lee's avatar
Gilbert Lee committed
1853
1854
1855
1856
1857
1858
1859
  }
  int numSrcs = srcList.size();


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

    for (int dstIndex = 0; dstIndex < numDevices; ++dstIndex)
gilbertlee-amd's avatar
gilbertlee-amd committed
1864
      dstList.push_back(std::make_pair(dstType, dstIndex));
Gilbert Lee's avatar
Gilbert Lee committed
1865
1866
1867
  }
  int numDsts = dstList.size();

1868
1869
  // Build array of possibilities, respecting any additional restrictions (e.g. XGMI hop count)
  struct TransferInfo
Gilbert Lee's avatar
Gilbert Lee committed
1870
  {
gilbertlee-amd's avatar
gilbertlee-amd committed
1871
1872
1873
    MemType srcType; int srcIndex;
    ExeType exeType; int exeIndex;
    MemType dstType; int dstIndex;
1874
1875
1876
1877
1878
1879
1880
1881
  };

  // 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
1882
  {
1883
1884
    // Skip CPU executors if XGMI link must be used
    if (useXgmiOnly && !IsGpuType(exeList[i].first)) continue;
gilbertlee-amd's avatar
gilbertlee-amd committed
1885
1886
    tinfo.exeType  = exeList[i].first;
    tinfo.exeIndex = exeList[i].second;
1887

gilbertlee-amd's avatar
gilbertlee-amd committed
1888
    bool isXgmiSrc  = false;
1889
1890
1891
1892
1893
1894
1895
    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)
        {
1896
1897
1898
#if defined(__NVCC__)
          isXgmiSrc = false;
#else
1899
          uint32_t exeToSrcLinkType, exeToSrcHopCount;
gilbertlee-amd's avatar
gilbertlee-amd committed
1900
1901
          HIP_CALL(hipExtGetLinkTypeAndHopCount(RemappedIndex(exeList[i].second, false),
                                                RemappedIndex(srcList[j].second, false),
1902
1903
1904
1905
                                                &exeToSrcLinkType,
                                                &exeToSrcHopCount));
          isXgmiSrc = (exeToSrcLinkType == HSA_AMD_LINK_INFO_TYPE_XGMI);
          if (isXgmiSrc) numHopsSrc = exeToSrcHopCount;
1906
#endif
1907
1908
1909
1910
1911
1912
1913
1914
1915
1916
1917
1918
1919
1920
1921
        }
        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
1922
1923
      tinfo.srcType  = srcList[j].first;
      tinfo.srcIndex = srcList[j].second;
1924
1925
1926
1927
1928
1929
1930
1931
1932

      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)
          {
1933
1934
1935
#if defined(__NVCC__)
            isXgmiSrc = false;
#else
1936
            uint32_t exeToDstLinkType, exeToDstHopCount;
gilbertlee-amd's avatar
gilbertlee-amd committed
1937
1938
            HIP_CALL(hipExtGetLinkTypeAndHopCount(RemappedIndex(exeList[i].second, false),
                                                  RemappedIndex(dstList[k].second, false),
1939
1940
1941
1942
                                                  &exeToDstLinkType,
                                                  &exeToDstHopCount));
            isXgmiDst = (exeToDstLinkType == HSA_AMD_LINK_INFO_TYPE_XGMI);
            if (isXgmiDst) numHopsDst = exeToDstHopCount;
1943
#endif
1944
1945
1946
1947
1948
1949
1950
1951
1952
1953
1954
1955
1956
1957
1958
1959
1960
          }
          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;

1961
1962
1963
1964
1965
1966
#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
1967
1968
        tinfo.dstType  = dstList[k].first;
        tinfo.dstIndex = dstList[k].second;
1969
1970
1971
1972

        possibleTransfers.push_back(tinfo);
      }
    }
Gilbert Lee's avatar
Gilbert Lee committed
1973
1974
  }

1975
1976
1977
  int const numPossible = (int)possibleTransfers.size();
  int maxParallelTransfers = (ev.sweepMax == 0 ? numPossible : ev.sweepMax);

Gilbert Lee's avatar
Gilbert Lee committed
1978
1979
1980
1981
1982
1983
  if (ev.sweepMin > numPossible)
  {
    printf("No valid test configurations exist\n");
    return;
  }

1984
1985
1986
1987
1988
1989
  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
1990
1991
  int numTestsRun = 0;
  int M = ev.sweepMin;
gilbertlee-amd's avatar
gilbertlee-amd committed
1992
1993
1994
1995
1996
1997
1998
1999
2000
2001
2002
  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
2003
2004
2005
2006
2007
2008
2009
2010
2011
  // 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
2012
      M = distribution(*ev.generator);
Gilbert Lee's avatar
Gilbert Lee committed
2013
2014
2015
2016

      // Generate a random bitmask
      for (int i = 0; i < numPossible; i++)
        bitmask[i] = (i < M) ? 1 : 0;
2017
      std::shuffle(bitmask.begin(), bitmask.end(), *ev.generator);
Gilbert Lee's avatar
Gilbert Lee committed
2018
2019
2020
2021
2022
2023
2024
2025
2026
2027
    }

    // 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
2028
2029
2030
2031
2032
        transfer.numSrcs        = 1;
        transfer.numDsts        = 1;
        transfer.srcType        = {possibleTransfers[value].srcType};
        transfer.srcIndex       = {possibleTransfers[value].srcIndex};
        transfer.exeType        = possibleTransfers[value].exeType;
2033
        transfer.exeIndex       = possibleTransfers[value].exeIndex;
gilbertlee-amd's avatar
gilbertlee-amd committed
2034
2035
2036
        transfer.dstType        = {possibleTransfers[value].dstType};
        transfer.dstIndex       = {possibleTransfers[value].dstIndex};
        transfer.numSubExecs    = IsGpuType(transfer.exeType) ? numGpuSubExecs : numCpuSubExecs;
gilbertlee-amd's avatar
gilbertlee-amd committed
2037
        transfer.numBytes       = ev.sweepRandBytes ? randSize(*ev.generator) * sizeof(float) : 0;
Gilbert Lee's avatar
Gilbert Lee committed
2038
2039
2040
2041
        transfers.push_back(transfer);
      }
    }

gilbertlee-amd's avatar
gilbertlee-amd committed
2042
2043
    LogTransfers(fp, ++numTestsRun, transfers);
    ExecuteTransfers(ev, numTestsRun, numBytesPerTransfer / sizeof(float), transfers);
Gilbert Lee's avatar
Gilbert Lee committed
2044
2045
2046
2047
2048
2049
2050
2051
2052
2053
2054
2055
2056
2057
2058
2059
2060
2061
2062
2063
2064
2065
2066
2067
2068
2069
2070
2071
2072
2073
2074

    // 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
2075
2076
2077
2078
2079
2080
2081
2082
2083
2084
  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
2085
2086
2087
2088
            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
2089
2090
2091
2092
            transfer.numBytes);
  }
  fprintf(fp, "\n");
  fflush(fp);
Gilbert Lee's avatar
Gilbert Lee committed
2093
}
gilbertlee-amd's avatar
gilbertlee-amd committed
2094
2095
2096
2097
2098
2099
2100
2101
2102
2103
2104

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