Skip to content

Commit 08d7c75

Browse files
mconcasmconcas
committed
Improve macro + improve measurement
1 parent 13daad5 commit 08d7c75

3 files changed

Lines changed: 86 additions & 47 deletions

File tree

‎GPU/GPUbenchmark/Shared/Utils.h‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -139,7 +139,7 @@ class ResultWriter
139139
public:
140140
explicit ResultWriter(const std::string resultsTreeFilename = "benchmark_results.root");
141141
~ResultWriter() = default;
142-
void storeBenchmarkEntry(int chunk, float entry, float chunkSizeGB);
142+
void storeBenchmarkEntry(int chunk, float entry, float chunkSizeGB, int nLaunches);
143143
void addBenchmarkEntry(const std::string bName, const std::string type, const int nChunks);
144144
void snapshotBenchmark();
145145
void saveToFile();
@@ -170,10 +170,10 @@ inline void ResultWriter::addBenchmarkEntry(const std::string bName, const std::
170170
mThroughputTrees.back()->Branch("throughput", &mThroughputResults);
171171
}
172172

173-
inline void ResultWriter::storeBenchmarkEntry(int chunk, float entry, float chunkSizeGB)
173+
inline void ResultWriter::storeBenchmarkEntry(int chunk, float entry, float chunkSizeGB, int nLaunches)
174174
{
175175
mTimeResults[chunk] = entry;
176-
mThroughputResults[chunk] = 1e3 * chunkSizeGB / entry;
176+
mThroughputResults[chunk] = 1e3 * chunkSizeGB * nLaunches / entry;
177177
}
178178

179179
inline void ResultWriter::snapshotBenchmark()

‎GPU/GPUbenchmark/cuda/Kernels.cu‎

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -326,14 +326,14 @@ std::vector<float> GPUbenchmark<chunk_type>::benchmarkAsync(void (*kernel)(int,
326326
// Warm up on every stream
327327
for (auto iStream{0}; iStream < nStreams; ++iStream) {
328328
chunk_type* chunkPtr = getPartPtr<chunk_type>(mState.scratchPtr, mState.chunkReservedGB, iStream);
329-
(*kernel)<<<blocks, threads, 0, streams[iStream]>>>(iStream, chunkPtr, args...);
329+
(*kernel)<<<blocks / mState.getMaxChunks(), threads, 0, streams[iStream]>>>(iStream, chunkPtr, args...);
330330
}
331331

332332
for (auto iStream{0}; iStream < nStreams; ++iStream) {
333333
chunk_type* chunkPtr = getPartPtr<chunk_type>(mState.scratchPtr, mState.chunkReservedGB, iStream);
334334
GPUCHECK(cudaEventRecord(starts[iStream], streams[iStream]));
335335
for (auto iLaunch{0}; iLaunch < nLaunches; ++iLaunch) {
336-
(*kernel)<<<blocks, threads, 0, streams[iStream]>>>(iStream, chunkPtr, args...);
336+
(*kernel)<<<blocks / mState.getMaxChunks(), threads, 0, streams[iStream]>>>(iStream, chunkPtr, args...);
337337
}
338338
GPUCHECK(cudaEventRecord(stops[iStream], streams[iStream]));
339339
}
@@ -424,7 +424,7 @@ void GPUbenchmark<chunk_type>::readSequential(SplitLevel sl)
424424
nThreads, // args...
425425
mState.deviceReadResultsPtr,
426426
capacity);
427-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB);
427+
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
428428
}
429429
mResultWriter.get()->snapshotBenchmark();
430430
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -448,7 +448,7 @@ void GPUbenchmark<chunk_type>::readSequential(SplitLevel sl)
448448
nThreads, // args...
449449
mState.deviceReadResultsPtr,
450450
capacity);
451-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB);
451+
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
452452
}
453453
mResultWriter.get()->snapshotBenchmark();
454454
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -479,7 +479,7 @@ void GPUbenchmark<chunk_type>::readConcurrent(SplitLevel sl, int nRegions)
479479
mState.deviceReadResultsPtr, // kernel arguments (chunkId is passed by wrapper)
480480
capacity);
481481
for (auto iResult{0}; iResult < results.size(); ++iResult) {
482-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB);
482+
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
483483
}
484484
mResultWriter.get()->snapshotBenchmark();
485485
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -503,7 +503,7 @@ void GPUbenchmark<chunk_type>::readConcurrent(SplitLevel sl, int nRegions)
503503
mState.deviceReadResultsPtr, // kernel arguments (chunkId is passed by wrapper)
504504
capacity);
505505
for (auto iResult{0}; iResult < results.size(); ++iResult) {
506-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB);
506+
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
507507
}
508508
mResultWriter.get()->snapshotBenchmark();
509509
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -551,7 +551,7 @@ void GPUbenchmark<chunk_type>::writeSequential(SplitLevel sl)
551551
nThreads,
552552
mState.deviceWriteResultsPtr,
553553
capacity);
554-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB);
554+
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
555555
}
556556
mResultWriter.get()->snapshotBenchmark();
557557
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -575,7 +575,7 @@ void GPUbenchmark<chunk_type>::writeSequential(SplitLevel sl)
575575
nThreads,
576576
mState.deviceWriteResultsPtr,
577577
capacity);
578-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB);
578+
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
579579
}
580580
mResultWriter.get()->snapshotBenchmark();
581581
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -606,7 +606,7 @@ void GPUbenchmark<chunk_type>::writeConcurrent(SplitLevel sl, int nRegions)
606606
mState.deviceWriteResultsPtr, // kernel arguments (chunkId is passed by wrapper)
607607
capacity);
608608
for (auto iResult{0}; iResult < results.size(); ++iResult) {
609-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB);
609+
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
610610
}
611611
mResultWriter.get()->snapshotBenchmark();
612612
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -630,7 +630,7 @@ void GPUbenchmark<chunk_type>::writeConcurrent(SplitLevel sl, int nRegions)
630630
mState.deviceWriteResultsPtr, // kernel arguments (chunkId is passed by wrapper)
631631
capacity);
632632
for (auto iResult{0}; iResult < results.size(); ++iResult) {
633-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB);
633+
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
634634
}
635635
mResultWriter.get()->snapshotBenchmark();
636636
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -680,7 +680,7 @@ void GPUbenchmark<chunk_type>::copySequential(SplitLevel sl)
680680
nThreads,
681681
mState.deviceCopyInputsPtr,
682682
capacity);
683-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB);
683+
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
684684
}
685685
mResultWriter.get()->snapshotBenchmark();
686686
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -704,7 +704,7 @@ void GPUbenchmark<chunk_type>::copySequential(SplitLevel sl)
704704
nThreads,
705705
mState.deviceCopyInputsPtr,
706706
capacity);
707-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB);
707+
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
708708
}
709709
mResultWriter.get()->snapshotBenchmark();
710710
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -735,7 +735,7 @@ void GPUbenchmark<chunk_type>::copyConcurrent(SplitLevel sl, int nRegions)
735735
mState.deviceCopyInputsPtr, // kernel arguments (chunkId is passed by wrapper)
736736
capacity);
737737
for (auto iResult{0}; iResult < results.size(); ++iResult) {
738-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB);
738+
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
739739
}
740740
mResultWriter.get()->snapshotBenchmark();
741741
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -759,7 +759,7 @@ void GPUbenchmark<chunk_type>::copyConcurrent(SplitLevel sl, int nRegions)
759759
mState.deviceCopyInputsPtr, // kernel arguments (chunkId is passed by wrapper)
760760
capacity);
761761
for (auto iResult{0}; iResult < results.size(); ++iResult) {
762-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB);
762+
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
763763
}
764764
mResultWriter.get()->snapshotBenchmark();
765765
std::cout << "\033[1;32m complete\033[0m" << std::endl;

‎GPU/GPUbenchmark/macro/showBenchmarks.C‎

Lines changed: 69 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,7 @@
1313

1414
int nBins{500};
1515
float minHist{0.f}, maxHist{1e4};
16-
void showBenchmarks(const TString fileName = "0_benchmark_results.root")
16+
void showBenchmarks(const bool times = false, const TString fileName = "0_benchmark_results.root")
1717
{
1818
auto f = TFile::Open(fileName.Data(), "read");
1919
std::unordered_map<std::string, TTree*> um_trees;
@@ -40,40 +40,79 @@ void showBenchmarks(const TString fileName = "0_benchmark_results.root")
4040
if (keyPair.first.find(type) != std::string::npos) {
4141
for (auto& pattern : patterns) {
4242
if (keyPair.first.find(pattern) != std::string::npos) {
43-
// Single Tree entry, we know test, type, mode, pattern
44-
std::vector<float>* measures = 0;
45-
TBranch* elapsed;
46-
keyPair.second->SetBranchAddress("elapsed", &measures, &elapsed);
47-
elapsed->GetEntry(keyPair.second->LoadTree(0));
48-
auto nChunk = measures->size();
49-
histograms.emplace_back(nChunk);
50-
for (int iHist{0}; iHist < (int)nChunk; ++iHist) {
51-
histograms.back()[iHist] = new TH1F(Form("Chunk_%d_%s", iHist, keyPair.first.c_str()), Form("Chunk_%d_%s;ms", iHist, keyPair.first.c_str()), 500, 0, 1e4);
52-
}
53-
for (size_t iEntry(0); iEntry < (size_t)keyPair.second->GetEntriesFast(); ++iEntry) {
54-
auto tentry = keyPair.second->LoadTree(iEntry);
55-
elapsed->GetEntry(tentry);
43+
if ((keyPair.first.find("TP") == std::string::npos)) {
44+
if (times) {
45+
// Single Tree entry, we know test, type, mode, pattern
46+
std::vector<float>* measures = 0;
47+
TBranch* elapsed;
48+
keyPair.second->SetBranchAddress("elapsed", &measures, &elapsed);
49+
elapsed->GetEntry(keyPair.second->LoadTree(0));
50+
auto nChunk = measures->size();
51+
histograms.emplace_back(nChunk);
52+
for (int iHist{0}; iHist < (int)nChunk; ++iHist) {
53+
histograms.back()[iHist] = new TH1F(Form("Chunk_%d_%s", iHist, keyPair.first.c_str()), Form("Chunk_%d_%s;ms", iHist, keyPair.first.c_str()), 1000, 0, 1e4);
54+
}
55+
for (size_t iEntry(0); iEntry < (size_t)keyPair.second->GetEntriesFast(); ++iEntry) {
56+
auto tentry = keyPair.second->LoadTree(iEntry);
57+
elapsed->GetEntry(tentry);
58+
for (int iHist{0}; iHist < (int)nChunk; ++iHist) {
59+
histograms.back()[iHist]->Fill((*measures)[iHist]);
60+
}
61+
}
62+
63+
std::vector<float> xCoord(nChunk), exCoord(nChunk), yCoord(nChunk), eyCoord(nChunk);
64+
65+
for (size_t i{0}; i < nChunk; ++i) {
66+
xCoord[i] = (float)i;
67+
yCoord[i] = histograms.back()[i]->GetMean();
68+
eyCoord[i] = histograms.back()[i]->GetRMS();
69+
exCoord[i] = 0.f;
70+
}
71+
TCanvas* c = new TCanvas(Form("c%s", keyPair.first.c_str()), Form("%s", keyPair.first.c_str()));
72+
c->cd();
73+
TGraphErrors* g = new TGraphErrors(nChunk, xCoord.data(), yCoord.data(), exCoord.data(), eyCoord.data());
74+
g->GetYaxis()->SetRangeUser(0, 5000);
75+
g->GetXaxis()->SetRangeUser(-2.f, nChunk);
76+
g->SetTitle(Form("%s, N_{test}=%d;chunk_id;elapsed (GB/s)", keyPair.first.c_str(), (int)keyPair.second->GetEntriesFast()));
77+
g->SetFillColor(40);
78+
g->Draw("AB");
79+
}
80+
} else { // TP plots //
81+
// Single Tree entry, we know test, type, mode, pattern
82+
std::vector<float>* measures = 0;
83+
TBranch* throughput;
84+
keyPair.second->SetBranchAddress("throughput", &measures, &throughput);
85+
throughput->GetEntry(keyPair.second->LoadTree(0));
86+
auto nChunk = measures->size();
87+
histograms.emplace_back(nChunk);
5688
for (int iHist{0}; iHist < (int)nChunk; ++iHist) {
57-
histograms.back()[iHist]->Fill((*measures)[iHist]);
89+
histograms.back()[iHist] = new TH1F(Form("Chunk_%d_%s", iHist, keyPair.first.c_str()), Form("Chunk_%d_%s;GB/s", iHist, keyPair.first.c_str()), 1000, 0, 1e3);
90+
}
91+
for (size_t iEntry(0); iEntry < (size_t)keyPair.second->GetEntriesFast(); ++iEntry) {
92+
auto tentry = keyPair.second->LoadTree(iEntry);
93+
throughput->GetEntry(tentry);
94+
for (int iHist{0}; iHist < (int)nChunk; ++iHist) {
95+
histograms.back()[iHist]->Fill((*measures)[iHist]);
96+
}
5897
}
59-
}
6098

61-
std::vector<float> xCoord(nChunk), exCoord(nChunk), yCoord(nChunk), eyCoord(nChunk);
99+
std::vector<float> xCoord(nChunk), exCoord(nChunk), yCoord(nChunk), eyCoord(nChunk);
62100

63-
for (size_t i{0}; i < nChunk; ++i) {
64-
xCoord[i] = (float)i;
65-
yCoord[i] = histograms.back()[i]->GetMean();
66-
eyCoord[i] = histograms.back()[i]->GetRMS();
67-
exCoord[i] = 0.f;
101+
for (size_t i{0}; i < nChunk; ++i) {
102+
xCoord[i] = (float)i;
103+
yCoord[i] = histograms.back()[i]->GetMean();
104+
eyCoord[i] = histograms.back()[i]->GetRMS();
105+
exCoord[i] = 0.f;
106+
}
107+
TCanvas* c = new TCanvas(Form("c%s", keyPair.first.c_str()), Form("%s", keyPair.first.c_str()));
108+
c->cd();
109+
TGraphErrors* g = new TGraphErrors(nChunk, xCoord.data(), yCoord.data(), exCoord.data(), eyCoord.data());
110+
g->GetYaxis()->SetRangeUser(0, 500);
111+
g->GetXaxis()->SetRangeUser(-2.f, nChunk);
112+
g->SetTitle(Form("%s, N_{test}=%d;chunk_id;throughput (GB/s)", keyPair.first.c_str(), (int)keyPair.second->GetEntriesFast()));
113+
g->SetFillColor(40);
114+
g->Draw("AB");
68115
}
69-
TCanvas* c = new TCanvas(Form("c%s", keyPair.first.c_str()), Form("%s", keyPair.first.c_str()));
70-
c->cd();
71-
TGraphErrors* g = new TGraphErrors(nChunk, xCoord.data(), yCoord.data(), exCoord.data(), eyCoord.data());
72-
g->GetYaxis()->SetRangeUser(0, 5000);
73-
g->GetXaxis()->SetRangeUser(-2.f, nChunk);
74-
g->SetTitle(Form("%s, N_{test}=%d;chunk_id;elapsed (ms)", keyPair.first.c_str(), (int)keyPair.second->GetEntriesFast()));
75-
g->SetFillColor(40);
76-
g->Draw("AB");
77116
}
78117
}
79118
}

0 commit comments

Comments
 (0)