Using VTune Profiler to Analyze GPU Applications

Launch VTune Profiler and from the Welcome page, click
New Project
. The Create a Project dialog box opens.
Specify a project name and a location for your project and click
Create 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 the
matrix_multiply
binary as an Application to profile.
In the HOW pane, select the GPU Offload analysis type from the Accelerators group.
Click
Start
to launch the analysis.

Run Analysis from the Command Line

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

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

Run GPU Compute/Media Hotspots Analysis

In the Accelerators group, select the GPU Compute/Media Hotspots analysis type.
Configure analysis options as described in the previous section.
Click
Start
to run the analysis.

vtune -collect gpu-hotspots -- ./matrix.dpcpp

vtune.exe -collect gpu-hotspots -- matrix_multiply.exe

Analyze Your Compute Tasks

// 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

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.

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%