ROCgdb - HackMD

# Offsite training * Get the initial material `cp -r /Shared/HPCTrainingExamples/ .` ## ROCgdb * Save the following code in a file called let's say `saxpy.hip` or use the file from ` HPCTrainingExamples/HIP/saxpy/` ``` #include <hip/hip_runtime.h> __constant__ float a = 1.0f; const float a_h = 1.0f; void init (int n, float *x, float *y) { for (std::size_t i = 0; i < n; ++i) { x[i] = 1.0; y[i] = 0.0; } } __global__ void saxpy (int n, float const* x, int incx, float* y, int incy) { int i = blockDim.x * blockIdx.x + threadIdx.x; if (i < n) y[i] += a*x[i]; } int main() { int n = 256; std::size_t size = sizeof(float)*n; float *h_x = new float [n]; float *h_y = new float [n]; init(n, h_x, h_y); float* d_x; float *d_y; hipMalloc(&d_x, size); hipMalloc(&d_y, size); hipMemcpy(d_x, h_x, size, hipMemcpyHostToDevice); hipMemcpy(d_y, h_y, size, hipMemcpyHostToDevice); hipMemcpy(d_x, h_x, size, hipMemcpyHostToDevice); hipMemcpy(d_y, h_y, size, hipMemcpyHostToDevice); int num_groups = 2; int group_size = 128; saxpy<<<num_groups, group_size>>>(n, d_x, 1, d_y, 1); hipDeviceSynchronize(); hipMemcpy(h_y, d_y, size, hipMemcpyDeviceToHost); delete [] h_x; delete [] h_y; hipFree(d_x); hipFree(d_y); return 0; } ``` * Compile `hipcc --offload-arch=gfx90a -o saxpy saxpy.hip` * Declare your GPU `export HIP_VISIBLE_DEVICES=X` * Execute `./saxpy` * Execute `saxpy` with rocgdb `rocgdb saxpy` * Inside rocgdb ``` (gdb) run ``` * Compile with debugging `hipcc -ggdb --offload-arch=gfx90a -o saxpy saxpy.hip` * Start rocgdb with the saxpy `rocgdb saxpy` * Inside rocgdb ``` (gdb) run (gdb) i th (gdb) t 1 (gdb) where (gdb) b saxpy.cpp:22 ``` * Compile for -O0 `hipcc -ggdb -O0 --offload-arch=gfx90a -o saxpy saxpy.cpp` * Inside rocgdb ``` (gdb) r (gdb) i th (gdb) show architecture (gdb) b saxpy (gdb) c (gdb) show architecture (gdb) info agents (gdb) info queues (gdb) info dispatches ``` ## Omnitrace * Load Omniperf ``` module load rocm module load omnitrace ``` * Before you execute any Omnitrace call, select a specific GPU: `export HIP_VISIBLE_DEVICES=X` * Check the various options and their values and also a second command for description `omnitrace-avail --categories omnitrace --brief --description` * Create an Omnitrace configuration file with description per option `omnitrace-avail -G omnitrace.cfg --all` * Declare to use this configuration file: `export OMNITRACE_CONFIG_FILE=/path/omnitrace.cfg` * Activate the TIMEMORY option in the `omnitrace.cfg`, edit the file, find the `OMNITRACE_USE_TIMEMORY` and declare it equal to true `OMNITRACE_USE_TIMEMORY = true` * Get the training examples: `git clone https://github.com/amd/HPCTrainingExamples.git` * Compile and execute Jacobi * `cd HPCTrainingExamples/HIP/jacobi` * `make clean;make` * Binary `jacobi_hip`  * Now execute the binary * `time mpirun -np 1 Jacobi_hip -g 1 1` * Check the duration  ### Binary rewriting (to be used with MPI codes and decreases overhead) * Binary rewriting: `omnitrace-instrument -o jacobi.inst -- ./Jacobi_hip` * We created a new instrumented binary called jacobi.inst * Executing the new instrumented binary: `time mpirun -n 1 omnitrace-run -- ./jacobi.inst -g 1 1` and check the duration * See the list of the instrumented GPU calls: `cat omnitrace-jacobi.inst-output/TIMESTAMP/roctracer-0.txt` * See the list of the instrumented CPU calls: `cat omnitrace-jacobi.inst-output/TIMESTAMP/wallclock-0.txt` or wallclock-1.txt * Check the MPI calls ### Visualization * Copy the `perfetto-trace.proto` to your laptop, open the web page https://ui.perfetto.dev/ click to open the trace and select the file perfetto-trace-0.proto or perfetto-trace-1.proto. * Where are all the MPI calls? * If MPI call is not called from the main call-stack then you need to profile the call-stack (default 64 layers) ### Call-stack Edit your omnitrace.cfg: ``` OMNITRACE_USE_SAMPLING = true; OMNITRACE_SAMPLING_FREQ = 100 ``` Execute again the instrumented binary and now you can see the call-stack when you visualize with perfetto. ### Hardware counters * See a list of all the counters: `omnitrace-avail --all` * Declare in your configuration file: `OMNITRACE_ROCM_EVENTS = GPUBusy,Wavefronts,VALUBusy,L2CacheHit,MemUnitBusy` * Execute: `omnitrace-run -- ./jacobi.inst -g 1 1` and copy the perfetto file and visualize ### Kernel timings * Open the file `omnitrace-binary-output/timestamp/wall_clock.txt` (replace binary and timestamp with your information) * In order to see the kernels gathered in your configuration file, make sure that `OMNITRACE_USE_TIMEMORY = true` and `OMNITRACE_FLAT_PROFILE = true`, execute the code and open again the file `omnitrace-binary-output/timestamp/wall_clock.txt` ## Omniperf * Add this in your path: `export PATH=/opt/conda/bin/:$PATH` * Load Omniperf ``` module load rocm module load omniperf ``` * Go to the Omniperf examples: ``` cd HPCTrainingExamples/OmniperfExamples/ ``` * Enter each directory and read the instructions even from a web page: https://github.com/amd/HPCTrainingExamples/tree/main/OmniperfExamples * Before you execute any Omniper call, select a specific GPU: `export HIP_VISIBLE_DEVICES=X`