Skip to content

Commit b61b50a

Browse files
mconcasmconcas
committed
Improve copy kernel
1 parent 29643bc commit b61b50a

4 files changed

Lines changed: 39 additions & 22 deletions

File tree

‎GPU/GPUbenchmark/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,7 @@ if(HIP_ENABLED)
3333
add_custom_command(
3434
OUTPUT ${HIP_KERNEL_PATH}
3535
COMMAND ${HIPIFY_EXECUTABLE} --quiet-warnings ${CU_KERNEL} | sed '1{/\#include \"hip\\/hip_runtime.h\"/d}' > ${HIP_KERNEL_PATH}
36+
DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/cuda/Kernels.cu
3637
)
3738
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
3839
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})

‎GPU/GPUbenchmark/Shared/Utils.h‎

Lines changed: 18 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, int nLaunches);
142+
void storeBenchmarkEntry(Test test, 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,25 @@ 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, int nLaunches)
173+
inline void ResultWriter::storeBenchmarkEntry(Test test, int chunk, float entry, float chunkSizeGB, int nLaunches)
174174
{
175+
// https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html
176+
// Eff_bandwidth (GB/s) = (B_r + B_w) / (~1e9 * Time (s))
175177
mTimeResults[chunk] = entry;
176-
mThroughputResults[chunk] = 1e3 * chunkSizeGB * nLaunches / entry;
178+
switch (test) {
179+
case Test::Read: {
180+
mThroughputResults[chunk] = 1e3 * chunkSizeGB * nLaunches / entry;
181+
break;
182+
}
183+
case Test::Write: {
184+
mThroughputResults[chunk] = 1e3 * chunkSizeGB * nLaunches / entry;
185+
break;
186+
}
187+
case Test::Copy: {
188+
mThroughputResults[chunk] = 2 * 1e3 * chunkSizeGB * nLaunches / entry;
189+
break;
190+
}
191+
}
177192
}
178193

179194
inline void ResultWriter::snapshotBenchmark()

‎GPU/GPUbenchmark/cuda/Kernels.cu‎

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -75,11 +75,11 @@ __global__ void readChunkSBKernel(
7575
chunk_t* results,
7676
size_t chunkSize)
7777
{
78-
chunk_t sink{0};
78+
chunk_t sink{0}; // local memory -> excluded from bandwidth accounting
7979
for (size_t i = threadIdx.x; i < chunkSize; i += blockDim.x) {
80-
sink += chunkPtr[i];
80+
sink += chunkPtr[i]; // 1 read operation, performed "chunkSize" times
8181
}
82-
results[chunkId] = sink;
82+
results[chunkId] = sink; // writing done once
8383
}
8484

8585
template <class chunk_t>
@@ -89,7 +89,7 @@ __global__ void readChunkMBKernel(
8989
chunk_t* results,
9090
size_t chunkSize)
9191
{
92-
chunk_t sink{0};
92+
chunk_t sink{0}; // local memory -> excluded from bandwidth accounting
9393
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < chunkSize; i += blockDim.x * gridDim.x) {
9494
sink += chunkPtr[i];
9595
}
@@ -130,7 +130,7 @@ __global__ void copyChunkSBKernel(
130130
size_t chunkSize)
131131
{
132132
for (size_t i = threadIdx.x; i < chunkSize; i += blockDim.x) {
133-
chunkPtr[i] = inputs[chunkId];
133+
chunkPtr[chunkSize - i - 1] = chunkPtr[i];
134134
}
135135
}
136136

@@ -142,7 +142,7 @@ __global__ void copyChunkMBKernel(
142142
size_t chunkSize)
143143
{
144144
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < chunkSize; i += blockDim.x * gridDim.x) {
145-
chunkPtr[i] = inputs[chunkId];
145+
chunkPtr[chunkSize - i - 1] = chunkPtr[i];
146146
}
147147
}
148148

@@ -423,7 +423,7 @@ void GPUbenchmark<chunk_t>::readSequential(SplitLevel sl)
423423
nThreads, // args...
424424
mState.deviceReadResultsPtr,
425425
capacity);
426-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
426+
mResultWriter.get()->storeBenchmarkEntry(Test::Read, iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
427427
}
428428
mResultWriter.get()->snapshotBenchmark();
429429
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -447,7 +447,7 @@ void GPUbenchmark<chunk_t>::readSequential(SplitLevel sl)
447447
nThreads, // args...
448448
mState.deviceReadResultsPtr,
449449
capacity);
450-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
450+
mResultWriter.get()->storeBenchmarkEntry(Test::Read, iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
451451
}
452452
mResultWriter.get()->snapshotBenchmark();
453453
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -478,7 +478,7 @@ void GPUbenchmark<chunk_t>::readConcurrent(SplitLevel sl, int nRegions)
478478
mState.deviceReadResultsPtr, // kernel arguments (chunkId is passed by wrapper)
479479
capacity);
480480
for (auto iResult{0}; iResult < results.size(); ++iResult) {
481-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
481+
mResultWriter.get()->storeBenchmarkEntry(Test::Read, iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
482482
}
483483
mResultWriter.get()->snapshotBenchmark();
484484
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -502,7 +502,7 @@ void GPUbenchmark<chunk_t>::readConcurrent(SplitLevel sl, int nRegions)
502502
mState.deviceReadResultsPtr, // kernel arguments (chunkId is passed by wrapper)
503503
capacity);
504504
for (auto iResult{0}; iResult < results.size(); ++iResult) {
505-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
505+
mResultWriter.get()->storeBenchmarkEntry(Test::Read, iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
506506
}
507507
mResultWriter.get()->snapshotBenchmark();
508508
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -550,7 +550,7 @@ void GPUbenchmark<chunk_t>::writeSequential(SplitLevel sl)
550550
nThreads,
551551
mState.deviceWriteResultsPtr,
552552
capacity);
553-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
553+
mResultWriter.get()->storeBenchmarkEntry(Test::Write, iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
554554
}
555555
mResultWriter.get()->snapshotBenchmark();
556556
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -574,7 +574,7 @@ void GPUbenchmark<chunk_t>::writeSequential(SplitLevel sl)
574574
nThreads,
575575
mState.deviceWriteResultsPtr,
576576
capacity);
577-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
577+
mResultWriter.get()->storeBenchmarkEntry(Test::Write, iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
578578
}
579579
mResultWriter.get()->snapshotBenchmark();
580580
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -605,7 +605,7 @@ void GPUbenchmark<chunk_t>::writeConcurrent(SplitLevel sl, int nRegions)
605605
mState.deviceWriteResultsPtr, // kernel arguments (chunkId is passed by wrapper)
606606
capacity);
607607
for (auto iResult{0}; iResult < results.size(); ++iResult) {
608-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
608+
mResultWriter.get()->storeBenchmarkEntry(Test::Write, iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
609609
}
610610
mResultWriter.get()->snapshotBenchmark();
611611
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -629,7 +629,7 @@ void GPUbenchmark<chunk_t>::writeConcurrent(SplitLevel sl, int nRegions)
629629
mState.deviceWriteResultsPtr, // kernel arguments (chunkId is passed by wrapper)
630630
capacity);
631631
for (auto iResult{0}; iResult < results.size(); ++iResult) {
632-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
632+
mResultWriter.get()->storeBenchmarkEntry(Test::Write, iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
633633
}
634634
mResultWriter.get()->snapshotBenchmark();
635635
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -679,7 +679,7 @@ void GPUbenchmark<chunk_t>::copySequential(SplitLevel sl)
679679
nThreads,
680680
mState.deviceCopyInputsPtr,
681681
capacity);
682-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
682+
mResultWriter.get()->storeBenchmarkEntry(Test::Copy, iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
683683
}
684684
mResultWriter.get()->snapshotBenchmark();
685685
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -703,7 +703,7 @@ void GPUbenchmark<chunk_t>::copySequential(SplitLevel sl)
703703
nThreads,
704704
mState.deviceCopyInputsPtr,
705705
capacity);
706-
mResultWriter.get()->storeBenchmarkEntry(iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
706+
mResultWriter.get()->storeBenchmarkEntry(Test::Copy, iChunk, result, mState.chunkReservedGB, mState.getNKernelLaunches());
707707
}
708708
mResultWriter.get()->snapshotBenchmark();
709709
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -734,7 +734,7 @@ void GPUbenchmark<chunk_t>::copyConcurrent(SplitLevel sl, int nRegions)
734734
mState.deviceCopyInputsPtr, // kernel arguments (chunkId is passed by wrapper)
735735
capacity);
736736
for (auto iResult{0}; iResult < results.size(); ++iResult) {
737-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
737+
mResultWriter.get()->storeBenchmarkEntry(Test::Copy, iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
738738
}
739739
mResultWriter.get()->snapshotBenchmark();
740740
std::cout << "\033[1;32m complete\033[0m" << std::endl;
@@ -758,7 +758,7 @@ void GPUbenchmark<chunk_t>::copyConcurrent(SplitLevel sl, int nRegions)
758758
mState.deviceCopyInputsPtr, // kernel arguments (chunkId is passed by wrapper)
759759
capacity);
760760
for (auto iResult{0}; iResult < results.size(); ++iResult) {
761-
mResultWriter.get()->storeBenchmarkEntry(iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
761+
mResultWriter.get()->storeBenchmarkEntry(Test::Copy, iResult, results[iResult], mState.chunkReservedGB, mState.getNKernelLaunches());
762762
}
763763
mResultWriter.get()->snapshotBenchmark();
764764
std::cout << "\033[1;32m complete\033[0m" << std::endl;

‎GPU/GPUbenchmark/macro/showBenchmarks.C‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -75,7 +75,8 @@ void showBenchmarks(const bool times = false, const TString fileName = "0_benchm
7575
g->GetYaxis()->SetRangeUser(0, 5000);
7676
g->GetXaxis()->SetRangeUser(-2.f, nChunk);
7777
g->SetTitle(Form("%s, N_{test}=%d;chunk_id;elapsed (GB/s)", keyPair.first.c_str(), (int)keyPair.second->GetEntriesFast()));
78-
g->SetFillColor(40);
78+
g->SetFillColor(kBlue);
79+
g->SetFillStyle(3335);
7980
g->Draw("AB");
8081
}
8182
} else { // TP plots //

0 commit comments

Comments
 (0)