Instrumentation to trace Global Memory operations in CUDA applications for sake of communication analysis. Consists of Clang front-end plugin for AST Manipulation and a LLVM pass for trace instrumentation. Applications instrumented with a tracing, generate a trace file in the /tmp folder.
The application creates one trace file for each CUDA stream.
Each file has the following format:
<size of one record, in bytes> //stored in a single byte
<kernel name> // name of kernel, followed by LF
<records> // all trace records, as one binary blob
000000000000000000000000 // EOT
<kernel name>
<records>
000000000000000000000000
<EOF>
Each record has the same structure, and looks like this:
|32 bit |4 bit |28 bit |64 bit |64 bit |
|191 160|159 156|155 128|127 64|63 32|31 16|15 0|
|SMID |Type |Size |Address|CTA.x |CTA.y |CTA.z|
- Clang/LLVM v. 5.0 or higher
- CMake
$ cd memtrace-pass/build
$ cmake ..
$ make
Instrumenting an application at compile time happens in two steps. The first step is the instrumenting all source file containing:
-
Kernel Calls
-
Kernel and
__device__
function definitions -
Kernel
__device__
functoin declarations
Each file needs to be parsed separately. For each file, a new file with
the prefix ’augmented-’ is created. This is then used for the further
build process. Compile units bear some problems with kernels defined and
declared in .cu
files, which are then included in other files. The
easiest workaround is to copy the kernel into the file including it, and
then only use the including file for instrumentation and build.
The following snipped is an example how to augment a kernel file.
clang++ -Xclang -load -Xclang "clang-plugin/libMemtrace-AA.so"
-Xclang -plugin -Xclang cuda-aug -Xclang -plugin-arg-cuda-aug
-Xclang -f -Xclang -plugin-arg-cuda-aug -Xclang ./augmented-kernel.cu
-I$CUDASYS/samples/common/inc -I. -I../utils
--cuda-path=$CUDASYS -std=c++11 -E application/kernel.cu
Next, the host and device utils are compiled, for later linking with the application.
// Host
clang++ -c -L$CUDASYS/lib64 --cuda-path=$CUDASYS
-I$CUDASYS/samples/common/inc -O1 --cuda-gpu-arch=sm_30
-I$CUDASYS/include -o hutils.o --std=c++11 -I. -I../utils
../utils/TraceUtils.cpp
// Device
clang++ -c -L$CUDASYS/lib64 --cuda-path=$CUDASYS
-I$CUDASYS/samples/common/inc -O1 --cuda-gpu-arch=sm_30
-I$CUDASYS/include -o dutils.o --std=c++11 -I. -I../utils
../utils/DeviceUtils.cu
Next, all the augmented kernel files are compiled at once.
clang++ -c -Xclang -load -Xclang $LLVMPLUGIN --cuda-path=$CUDASYS
-I$CUDASYS/samples/common/inc --cuda-gpu-arch=sm_30 -L$CUDASYS/lib64 -O1
-lcudart_static -m64 --std=c++11 -I. -I../utils -I<app-includes>
./augmented-kernel.cu ./augmented-kernel2.cu
Finally, all compiled files are linked together.
clang++ --cuda-path=$CUDASYS -I$CUDASYS/samples/common/inc --cuda-gpu-arch=sm_30 -L$CUDASYS/lib64 -O1
-lcudart -ldl -lrt -L. -m64 --std=c++11 -I. -I../utils -I<app-includes>
-o application dutils.o hutils.o augmented-kernel.o augmented-kernel2.o
During the execution one or more files are generated in the /tmp
folder. The files are named MemTrace-pipe-<n>
, one for each stream.
Any software related to what that happens after the analysis, is stored in ’memtrace-pass/post-processing’. The trace files can be parsed with the ’extract-subset.py’ python3 script. From the original full trace, the communication subset is extracted. Three data structures are generated by ’extract-subset.py’. All structures are stored using pickle. The data structures are of type ’AutoDict()’, which is defined in ’TraceInc.py’
-
Load/Store volumes, both total and during communication, for Kernels, CTAs and SMs. Separated by kernel and superstep.
"KCV" : { // Kernel Communication Volume <kernel>: { <superstep>: { "Load" : <int> "Store": <int> } } } "KDV" : { // Total Kernel Data Volume <kernel>: { <superstep>: { "Load" : <int> "Store": <int> } } } "CCV" : { // CTA Communication Volume <kernel>: { <superstep>: { <CTA> : { "Load" : <int> "Store": <int> } } } } "CDV" : { // Total CTA Data Volume <kernel>: { <superstep>: { <CTA> : { "Load" : <int> "Store": <int> } } } } "SCV" : { // SM Communication Volume <kernel>: { <superstep>: { <SM> : { "Load" : <int> "Store": <int> } } } } "SDV" : { // Total SM Data Volume <kernel>: { <superstep>: { <SM> : { "Load" : <int> "Store": <int> } } } }
-
Map of transfers, ordered by kernel, CTA and superstep.
<source-kernels> : { <source-cta> : { <source-superstep> : { <recv-kernel>: { <recv-cta>: { <recv-superstep>: { 'Size' : <int> 'cnt' : <int> } } } } } }
-
All addresses involved in communication, with all operations performed on this address, in superstep order.
{ <address> : [ { // record "kernel": <kernel> "it" : <superstep> "cta" : <CTA-ID, xyz order> "addr" : <address> "smid" : <SM id> "size" : <operation size in bytes> "type" : <type of MOp> } ] }