GPU Analysis with VTuneTM Profiler
VTune
TM
Profiler is a performance analysis tool for serial and
multi-threaded applications. It helps you analyze algorithm choices
and identify where and how your application can benefit from available
hardware resources. Use it to locate or determine:- Sections of code that don’t effectively utilize available processor resources
- The best sections of code to optimize for both sequential and threaded performance
- Synchronization objects that affect the application performance
- Whether, where, and why your application spends time on input/output operations
- Whether your application is CPU-bound or GPU-bound and how effectively it offloads code to the GPU
- The performance impact of different synchronization methods, different numbers of threads, or different algorithms
- Thread activity and transitions
- Hardware-related issues in your code such as data sharing, cache misses, branch misprediction, and others
- Profiling a SYCL application running on a GPU
The tool also has new features to support GPU analysis:
- GPU Offload Analysis (technical preview)
- GPU Compute/Media Hotspots Analysis (technical preview)
GPU Offload Analysis
(Preview)Use this analysis type to analyze code execution on the CPU and GPU
cores of your platform, correlate CPU and GPU activity, and identify
whether your application is GPU-bound or CPU-bound. The tool
infrastructure automatically aligns clocks across all cores in the
system so you can analyze some CPU-based workloads together with
GPU-based workloads within a unified time domain. This analysis lets
you:
- Identify how effectively your application uses SYCL or OpenCLTMkernels.
- Analyze execution of Intel Media SDK tasks over time (for Linux targets only).
- Explore GPU usage and analyze a software queue for GPU engines at each moment in time.
GPU Compute/Media Hotspots Analysis
(Preview)Use this tool to analyze the most time-consuming GPU kernels,
characterize GPU usage based on GPU hardware metrics, identify
performance issues caused by memory latency or inefficient kernel
algorithms, and analyze GPU instruction frequency for certain
instruction types. The GPU Compute/Media Hotspots analysis lets you:
- Explore GPU kernels with high GPU utilization, estimate the efficiency of this utilization, and identify possible reasons for stalls or low occupancy.
- Explore the performance of your application per selected GPU metrics over time.
- Analyze the hottest SYCL or OpenCL kernels for inefficient kernel code algorithms or incorrect work item configuration.
- Run GPU Offload Analysis on a SYCL Application.
Using VTune Profiler to Analyze GPU Applications
- Launch VTune Profiler and from the Welcome page, clickNew Project. The Create a Project dialog box opens.
- Specify a project name and a location for your project and clickCreate Project. The Configure Analysis window opens.
- Make sure the Local Host is selected in the WHERE pane.
- In the WHAT pane, make sure the Launch Application target is selected and specify thematrix_multiplybinary as an Application to profile.
- In the HOW pane, select the GPU Offload analysis type from the Accelerators group.
- ClickStartto launch the analysis.
This is the least intrusive analysis for applications running on
platforms with Intel Graphics as well as third-party GPUs
supported by VTune Profiler.
Run Analysis from the Command Line
To run the analysis from the command line:
On Linux* OS:
- Set VTune Profiler environment variables by sourcing the script:
source <*install_dir*>/env/vars.sh
- Run the analysis command:
vtune -collect gpu-offload -- ./matrix.dpcpp
On Windows* OS:
- Set VTune Profiler environment variables by running the batch file:
export <*install_dir*>\env\vars.bat
- Run the analysis command:
vtune.exe -collect gpu-offload -- matrix_multiply.exe
Analyzing Collected Data
Start your analysis with the GPU Offload viewpoint. In the Summary
window, you can see see statistics on CPU and GPU resource usage to
determine if your application is GPU-bound, CPU-bound, or not
effectively utilizing the compute capabilities of the system. In this
example, the application should use the GPU for intensive
computation. However, the result summary indicates that GPU usage is
actually low.
GPU Offload viewpoint

Switch to the Platform window. Here, you can see basic CPU and GPU metrics that
help analyze GPU usage on a software queue. This data is correlated with CPU
usage on the timeline. The information in the Platform window can help you
make some inferences.
GPU Utilization

Most applications may not present obvious situations as described
above. A detailed analysis is important to understand all
dependencies. For example, GPU engines that are responsible for video
processing and rendering are loaded in turns. In this case, they are
used in a serial manner. When the application code runs on the CPU,
this can cause ineffective scheduling on the GPU, which can
lead you to mistakenly interpret the application as GPU bound.
Identify the GPU execution phase based on the computing task reference
and GPU Utilization metrics. Then, you can define the overhead for
creating the task and placing it in a queue.
GPU Offload Platform window

To investigate a computing task, switch to the Graphics window to
examine the type of work (rendering or computation) running on the GPU
per thread. Select the Computing Task grouping and use the table to
study the performance characterization of your task. To further
analyze your computing task, run the GPU Compute/Media Hotspots
analysis type.
Computing Task grouping

Run GPU Compute/Media Hotspots Analysis
To run the analysis:
- In the Accelerators group, select the GPU Compute/Media Hotspots analysis type.
- Configure analysis options as described in the previous section.
- ClickStartto run the analysis.
GPU Compute/Media Hotspots analysis

Run Analysis from the Command line
On Linux OS:
vtune -collect gpu-hotspots -- ./matrix.dpcpp
On Windows OS:
vtune.exe -collect gpu-hotspots -- matrix_multiply.exe
Analyze Your Compute Tasks
Characterization profile

The default analysis configuration invokes the Characterization
profile with the Overview metric set. In addition to individual
compute task characterization available through the GPU
Offload analysis, VTune Profiler provides memory bandwidth metrics
that are categorized by different levels of GPU memory hierarchy.
VTune Profiler memory bandwidth metrics

You can analyze compute tasks at source code level too. For example,
to count GPU clock cycles spent on a particular task or due to memory latency,
use the Source Analysis option.
GPU Compute/Media Hotspots analysis, Source Analysis

Use our matrix multiply example in SYCL:
// Basic matrix multiply
void multiply1(int msize, int tidx, int numt, TYPE a[][NUM], TYPE b[][NUM],
TYPE c[][NUM], TYPE t[][NUM]) {
int i, j, k;
// Declare a deviceQueue
sycl::default_selector device;
sycl::queue q(device, exception_handler);
cout << "Running on " << q.get_device().get_info<sycl::info::device::name>()
<< "\n";
// Declare a 2 dimensional range
sycl::range<2> matrix_range{NUM, NUM};
// Declare 3 buffers and Initialize them
sycl::buffer<TYPE, 2> bufferA((TYPE *)a, matrix_range);
sycl::buffer<TYPE, 2> bufferB((TYPE *)b, matrix_range);
sycl::buffer<TYPE, 2> bufferC((TYPE *)c, matrix_range);
// Submit our job to the queue
q.submit([&](auto &h) {
// Declare 3 accessors to our buffers. The first 2 read and the last
// read_write
sycl::accessor accessorA(bufferA, h, sycl::read_only);
sycl::accessor accessorB(bufferB, h, sycl::read_only);
sycl::accessor accessorC(bufferC, h);
// Execute matrix multiply in parallel over our matrix_range
// ind is an index into this range
h.parallel_for(matrix_range, [=](sycl::id<2> ind) {
int k;
for (k = 0; k < NUM; k++) {
// Perform computation ind[0] is row, ind[1] is col
accessorC[ind[0]][ind[1]] +=
accessorA[ind[0]][k] * accessorB[k][ind[1]];
}
});
}).wait_and_throw();
} // multiply1
Analyzing the GPU-offload report from the command-line gives
detailed recommendations on how to optimize the application.
Elapsed Time: 2.805s
GPU Utilization: 3.3%
GPU utilization is low. Consider offloading more work to the GPU to
increase overall application performance.
GPU Utilization
GPU Engine Packet Type GPU Time GPU Utilization(%)
---------------- ----------- -------- ---------------------
Render and GPGPU Unknown 0.091s 3.3%
Hottest GPU Computing Tasks
Computing Task Total Time Execution Time % of Execution(%) Instance Count
-------------- ---------- -------------- ----------------- --------------
Matrix1<float> 0.183s 0.086s 47.0% 1
Recommendations:
GPU Utilization: 3.3%
GPU utilization is low. Switch to the Graphics View for in-depth analysis
of host activity. Poor GPU utilization can prevent the application from
offloading effectively.
Transfer Time: 0.097s
Execution time on the device is less than memory transfer time. Make sure
your offload schema is optimal. Use the Intel Advisor tool to get an
insight into possible causes of inefficient offload.
We can also examine how efficiently our GPU kernel is running using
GPU-hotspots. How often our execution units are stalled can be a good
indication of GPU performance. Another important metric is whether we
are L3 bandwidth bound. In our case VTune is indicating that our L3
bandwidth was high when VEs were stalled.
Elapsed Time: 1.849s
GPU Time: 0.090s
EU Array Stalled/Idle: 6.2%
GPU L3 Bandwidth Bound: 65.2%
L3 bandwidth was high when VEs were stalled or idle.
Consider improving cache reuse.
FPU Utilization: 76.4%
For more ways to optimize GPU performance using VTune Profiler,
see Software Optimization for Intel® GPUs
in the
Intel® VTune™ Profiler Performance Analysis Cookbook and
Optimize Applications for Intel® GPUs with Intel® VTune™ Profiler .