oneAPI GPU Optimization Guide

Developer Guide

Contents

GPU Analysis with VTune™ Profiler

VTune
™
 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 DPC++ 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 DPC++ or OpenCL
    ™
     kernels.
  • 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 DPC++ or OpenCL kernels for inefficient kernel code algorithms or incorrect work item configuration.
  • Run GPU Offload Analysis on a DPC++ Application.

Using VTune Profiler to Analyze GPU Applications

  1. Launch VTune Profiler and from the Welcome page, click
    New Project
    . The Create a Project dialog box opens.
  2. Specify a project name and a location for your project and click
    Create Project
    . The Configure Analysis window opens.
  3. Make sure the Local Host is selected in the WHERE pane.
  4. In the WHAT pane, make sure the Launch Application target is selected and specify the
    matrix_multiply
     binary as an Application to profile.
  5. In the HOW pane, select the GPU Offload analysis type from the Accelerators group.
  6. Click
    Start
     to 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:
  1. Set VTune Profiler environment variables by sourcing the script:
source <*install_dir*>/env/vars.sh
  1. Run the analysis command:
vtune -collect gpu-offload -- ./matrix.dpcpp
On Windows* OS:
  1. Set VTune Profiler environment variables by running the batch file:
export <*install_dir*>\env\vars.bat
  1. 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
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
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
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
Computing Task grouping

Run GPU Compute/Media Hotspots Analysis

To run the analysis:
  1. In the Accelerators group, select the GPU Compute/Media Hotspots analysis type.
  2. Configure analysis options as described in the previous section.
  3. Click
    Start
     to run the analysis.
GPU Compute/Media Hotspots 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
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
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
GPU Compute\/Media Hotspots analysis, Source Analysis
Use our matrix multiply example in DPC++:
// 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 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 EUs 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 EUs 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 .

Product and Performance Information

1

Performance varies by use, configuration and other factors. Learn more at www.Intel.com/PerformanceIndex.