Intel® VTune™ Profiler

User Guide

Download PDF
ID 766319
Date 6/24/2024
Public

A newer version of this document is available. Customers should click here to go to the newest version.

Document Table of Contents
Document Table of Contents x
User Guide
User Guide x
Introduction Install Intel® VTune™ Profiler Open Intel® VTune™ Profiler Set Up Project Set Up Analysis Target Analyze Performance Intel® VTune™ Profiler Command Line Interface API Support Troubleshooting Reference Notices and Disclaimers
Introduction x
Tuning Methodology Tutorials and Samples Notational Conventions Get Help Product Website and Support Related Information
Install Intel® VTune™ Profiler x
Sampling Drivers Set Up System for GPU Analysis Verify Intel® VTune™ Profiler Installation Install VTune Profiler Server Security Best Practices
Set Up System for GPU Analysis x
Rebuild and Install the Kernel for GPU Analysis Rebuild and Install Module i915 for GPU Analysis on CentOS* Rebuild and Install Module i915 for GPU Analysis on Ubuntu*
Install VTune Profiler Server x
Set Up Transport Security Configure User Authentication/Authorization
Open Intel® VTune™ Profiler x
Get Started with Intel® VTune™ Profiler Intel® VTune™ Profiler Graphical User Interface Web Server Interface Microsoft Visual Studio* Integration Eclipse* and Intel System Studio IDE Integration Containerization Support
Containerization Support x
Run VTune Profiler in a Container Profile Container Targets from the Host
Set Up Project x
WHERE: Analysis System WHAT: Analysis Target HOW: Analysis Types Search Directories
WHERE: Analysis System x
Analysis System Options
WHAT: Analysis Target x
Analysis Target Options
Search Directories x
Search Order
Set Up Analysis Target x
Prepare Application for Analysis Windows* Targets Linux* Targets Embedded Linux* Targets FreeBSD* Targets QNX* Targets Managed Code Targets Android* Targets Intel® Xeon Phi™ Processor Targets Targets in Virtualized Environments Targets in a Cloud Environment Arbitrary Targets Embedded System Targets
Windows* Targets x
Install the Sampling Drivers for Windows* Targets Debug Information for Windows* Application Binaries Compiler Switches for Performance Analysis on Windows* Targets Debug Information for Windows* System Libraries Add Administrative Privileges
Linux* Targets x
Build and Install the Sampling Drivers for Linux* Targets Debug Information for Linux* Application Binaries Compiler Switches for Performance Analysis on Linux* Targets Enable Linux* Kernel Analysis Resolution of Symbol Names for Linux-Loadable Kernel Modules Analyze Statically Linked Binaries on Linux* Targets Set Up Remote Linux* Target
Set Up Remote Linux* Target x
Set Up Linux* System for Remote Analysis Configure SSH Access for Remote Collection Search Directories for Remote Linux* Targets Temporary Directory for Performance Results on Linux* Targets
Embedded Linux* Targets x
Configure Yocto Project* and VTune Profiler with the Integration Layer Configure Yocto Project*/Wind River* Linux* and Intel® VTune™ Profiler with the Intel System Studio Integration Layer Configure Yocto Project* and Intel® VTune™ Profiler with the Linux* Target Package
FreeBSD* Targets x
Set Up FreeBSD* System
Managed Code Targets x
.NET* Targets Windows Store Application Targets Go* Application Targets
Android* Targets x
Build and Install Sampling Drivers for Android* Targets Set Up Android* System Enable Java* Analysis on Android* System Prepare an Android* Application for Analysis Analyze Unplugged Devices Search Directories for Android* Targets
Targets in Virtualized Environments x
Profile Targets on a VMware* Guest System Profile Targets on a Parallels* Guest System Profile Targets on a KVM* Guest System Profile Targets on a Xen* Virtualization Platform Profile Targets in the Hyper-V* Environment
Profile Targets on a KVM* Guest System x
Profile KVM Kernel Modules from the Host Profile KVM Kernel and User Space on the KVM System Profile KVM Kernel and User Space from the Host
Analyze Performance x
User-Mode Sampling and Tracing Collection Hardware Event-based Sampling Collection Performance Snapshot Algorithm Group Microarchitecture Analysis Group Parallelism Analysis Group Input and Output Analysis Accelerators Analysis Group Platform Analysis Group Platform Analysis Hybrid CPU Analysis Source Code Analysis Custom Analysis Energy Analysis Code Profiling Scenarios Control Data Collection Manage Data Views Manage Result Files
Hardware Event-based Sampling Collection x
Allow Multiple Runs or Multiplex Events Hardware Event-based Sampling Collection with Stacks
Algorithm Group x
Hotspots Analysis for CPU Usage Issues Anomaly Detection Analysis (preview) Memory Consumption Analysis
Hotspots Analysis for CPU Usage Issues x
Hotspots View
Anomaly Detection Analysis (preview) x
Anomaly Detection View
Memory Consumption Analysis x
Memory Consumption and Allocations View
Microarchitecture Analysis Group x
Microarchitecture Exploration Analysis for Hardware Issues Memory Access Analysis for Cache Misses and High Bandwidth Issues
Microarchitecture Exploration Analysis for Hardware Issues x
Microarchitecture Exploration View Microarchitecture Pipe
Memory Access Analysis for Cache Misses and High Bandwidth Issues x
Memory Usage View
Parallelism Analysis Group x
Threading Analysis HPC Performance Characterization Analysis
Threading Analysis x
Threading Efficiency View
HPC Performance Characterization Analysis x
HPC Performance Characterization View
Input and Output Analysis x
Analyze Platform Performance Analyze DPDK Applications Analyze SPDK Applications Analyze Linux Kernel I/O
Accelerators Analysis Group x
GPU Offload Analysis GPU Compute/Media Hotspots Analysis (Preview) CPU/FPGA Interaction Analysis NPU Exploration Analysis
GPU Compute/Media Hotspots Analysis (Preview) x
GPU Compute/Media Hotspots View
CPU/FPGA Interaction Analysis x
CPU/FPGA Interaction View
NPU Exploration Analysis x
NPU Exploration View
Platform Analysis Group x
System Overview Analysis
System Overview Analysis x
Analyze Interrupts Analyze Latency Issues
Custom Analysis x
Custom Analysis Options
Custom Analysis Options x
Highly Accurate CPU Time Data Collection Hardware Event List Linux* and Android* Kernel Analysis Sampling Interval Sample After Value
Hardware Event List x
Hardware Event Skid Instructions Retired Event Precise Events
Energy Analysis x
Run Energy Analysis View Energy Analysis Data with Intel® VTune™ Profiler Interpret Energy Analysis Data with Intel® VTune™ Profiler
Code Profiling Scenarios x
Java* Code Analysis Python* Code Analysis Intel® Threading Building Blocks Code Analysis MPI Code Analysis OpenSHMEM* Code Analysis with Fabric Profiler GPU Application Analysis on Intel® HD Graphics and Intel® Iris® Graphics Frame Data Analysis Task Analysis
GPU Application Analysis on Intel® HD Graphics and Intel® Iris® Graphics x
GPU OpenCL™ Application Analysis Explore Summary Statistics Analyze Hot GPU OpenCL Kernels Correlate OpenCL Kernels Data with GPU Metrics Explore the Computing Queue Analyze Source and Assembly Code See Also Intel® Media SDK Program Analysis
Control Data Collection x
Finalization Pause Data Collection Limit Data Collection Generate Command Line Configuration from GUI Minimize Collection Overhead Import External Data
Import External Data x
Use a Custom Collector Create a CSV File with External Data Import Linux Perf* Trace with VTune Profiler Metrics Examples of CSV Format and Imported Data
Manage Data Views x
Switch Viewpoints Control Window Synchronization View Stacks Manage Grid Views Manage Timeline View Change Threshold Values Choose Data Format Group and Filter Data View Data on Inline Functions Analyze Loops Stitch Stacks for Intel® oneAPI Threading Building Blocks or OpenMP* Analysis Search for Data
View Stacks x
Call Stack Mode Metrics Distribution Over Call Stacks
Manage Result Files x
VTune Profiler Filenames and Locations Import Results and Traces into VTune Profiler GUI Compare Results
Compare Results x
Compare Source Code View Comparison Data
View Comparison Data x
Comparison Summary Bottom-up Comparison Top-down Tree Comparison
Intel® VTune™ Profiler Command Line Interface x
vtune Command Syntax vtune Actions Run Command Line Analysis Work with Results from Command Line Generate Command Line Reports Command Line Usage Scenarios Command Line Interface Reference Report Problems from Command Line
Run Command Line Analysis x
performance-snapshot Command Line Analysis hotspots Command Line Analysis anomaly-detection Command Line Analysis threading Command Line Analysis memory-consumption Command Line Analysis hpc-performance Command Line Analysis uarch-exploration Command Line Analysis memory-access Command Line Analysis tsx-exploration Command Line Analysis tsx-hotspots Command Line Analysis sgx-hotspots Command Line Analysis gpu-hotspots Command Line Analysis gpu-offload Command Line Analysis npu graphics-rendering Command Line Analysis fpga-interaction Command Line Analysis io Command Line Analysis system-overview Command Line Analysis runsa/runss Custom Command Line Analysis Configure Analysis Options from Command Line
Configure Analysis Options from Command Line x
Collect System-Wide Data from Command Line Collect Data on Remote Linux* Systems from Command Line Configure GPU Analysis from Command Line Specify Search Directories from Command Line Specify Result Directory from Command Line Pause Collection from Command Line Manage Analysis Duration from Command Line Limit Data Collection from Command Line
Work with Results from Command Line x
View Command Line Results in the GUI Import Results from Command Line Re-finalize Results from Command Line
Generate Command Line Reports x
Summary Report Hotspots Report Hardware Events Report Callstacks Report Timeline Report Top-down Report GPU Compute/Media Hotspots Report gprof-cc Report Difference Report View Source Objects from Command Line Save and Format Command Line Reports Filter and Group Command Line Reports
Command Line Usage Scenarios x
Use VTune Profiler Server in Containers Android* Target Analysis from the Command Line OpenMP* Analysis from the Command Line Java* Code Analysis from the Command Line
Command Line Interface Reference x
Option Descriptions and General Rules allow-multiple-runs analyze-kvm-guest analyze-system app-working-dir archive call-stack-mode collect collect-with column command cpu-mask csv-delimiter cumulative-threshold-percent custom-collector data-limit discard-raw-data duration filter finalization-mode finalize format group-by help import inline-mode knob kvm-guest-kallsyms kvm-guest-modules limit loop-mode mrte-mode no-follow-child no-summary no-unplugged-mode quiet report report-knob report-output report-width result-dir resume-after return-app-exitcode ring-buffer search-dir show-as sort-asc sort-desc source-object source-search-dir stack-size start-paused strategy target-install-dir target-system target-tmp-dir target-duration-type target-pid target-process time-filter trace-mpi user-data-dir verbose version
API Support x
Support for Instrumentation and Tracing Technology API (ITT API) JIT Profiling API System APIs Supported by Intel® VTune™ Profiler
Support for Instrumentation and Tracing Technology API (ITT API) x
Basic Usage and Configuration Instrumentation and Tracing Technology API Reference
Basic Usage and Configuration x
Configure Your Build System Attach ITT APIs to a Launched Application Instrument Your Application Minimize ITT API Overhead View Instrumentation and Tracing Technology (ITT) API Task Data in Intel® VTune™ Profiler
Instrumentation and Tracing Technology API Reference x
Domain API String Handle API Collection Control API Thread Naming API Task API Frame API Histogram API User-Defined Synchronization API Event API Counter API Context Metadata API Load Module API Memory Allocation APIs
JIT Profiling API x
Using JIT Profiling API JIT Profiling API Reference
JIT Profiling API Reference x
iJIT_NotifyEvent iJIT_IsProfilingActive iJIT_ GetNewMethodID
Troubleshooting x
Best Practices: Resolve Intel® VTune™ Profiler BSODs, Crashes, and Hangs in Windows* OS Error Message: Application Sets Its Own Handler for Signal Error Message: Cannot Enable Event-Based Sampling Collection Error Message: Cannot Collect GPU Hardware Metrics Error Message: Cannot Load Data File Error Message: Cannot Locate Debugging Information Error Message: Cannot Open Data Error Message: Client Is Not Authorized to Connect to Server Error Message: Root Privileges Required for Processor Graphics Events Error Message: No Pre-built Driver Exists for This System Error Message: Not All OpenCL™ API Profiling Callbacks Are Received Error Message: Problem Accessing the Sampling Driver Error Message: Required Key Not Available Error Message: Scope of ptrace System Call Is Limited Error Message: Stack Size Is Too Small Error Message: Symbol File Is Not Found Problem: Analysis of the .NET* Application Fails Problem: Cannot Access VTune Profiler Documentation Problem: CPU time for Hotspots or Threading Analysis is Too Low Problem: 'Events= Sample After Value (SAV) * Samples' Is Not True If Multiple Runs Are Disabled Problem: Guessed Stack Frames Problem: GUI Hangs or Crashes Problem: Inaccurate Sum in the Grid Problem: Information Collected via ITT API Is Not Available When Attaching to a Process Problem: No GPU Utilization Data Is Collected Problem: Same Functions Are Compared As Different Instances Problem: Skipped Stack Frames Problem: Stack in the Top-Down Tree Window Is Incorrect Problem: Stacks in Call Stack and Bottom-Up Panes Are Different Problem: System Functions Appear in the User Functions Only Mode Problem: Intel® VTune™ Profiler is Slow to Respond When Collecting or Displaying Data Problem: Intel® VTune™ Profiler is Slow on X-Servers with SSH Connection Problem: Unexpected Paused Time Problem: {Unknown Timer} in the Platform Power Analysis Viewpoint Problem: Unknown Critical Error Due to Disabled Loopback Interface Problem: Unknown Frames Problem: Unsupported Microsoft* Windows* OS Warnings about Accurate CPU Time Collection
Reference x
User Interface Reference CPU Metrics Reference GPU Metrics Reference OpenCL™ Kernel Analysis Metrics Reference Energy Analysis Metrics Reference Intel Processor Events Reference
User Interface Reference x
Context Menu: Grid Context Menus: Call Stack Pane Context Menus: Project Navigator Context Menus: Source/Assembly Window Dialog Box: Binary/Symbol Search Dialog Box: Source Search Hot Keys Menu: Customize Grouping Menu: Intel VTune Profiler Pane: Call Stack Pane: Options - General Pane: Options - Result Location Pane: Options - Source/Assembly Project Navigator Pane: Timeline Toolbar: Configure Analysis Toolbar: Filter Toolbar: Source/Assembly Toolbar: Intel VTune Profiler Window: Bandwidth - Platform Power Analysis Window: Bottom-up Window: Caller/Callee Window: Cannot Find <file type> File Window: Collection Log Window: Compare Results Window: Configure Analysis Window: Core Wake-ups - Platform Power Analysis Window: Correlate Metrics - Platform Power Analysis Window: CPU C/P States - Platform Power Analysis Window: Debug Window: Event Count - Hardware Events Window: Flame Graph Window: Graphics - GPU Compute/Media Hotspots Window: Graphics C/P States - Platform Power Analysis Window: NC Device States - Platform Power Analysis Window: Platform Window: Platform Power Analysis Window: Sample Count - Hardware Events Window: SC Device States - Platform Power Analysis Window: Summary Window: System Sleep States - Platform Power Analysis Window: Temperature/Thermal Sample - Platform Power Analysis Window: Timer Resolution - Platform Power Analysis Window: Top-down Tree Window: Uncore Event Count - Hardware Events Window: Wakelocks - Platform Power Analysis
Window: Summary x
Window: Summary - Input and Output Summary Window: Summary - Microarchitecture Exploration Window: Summary - GPU Analysis Window: Summary - Hardware Events Window: Summary - Hotspots by CPU Utilization Window: Summary - HPC Performance Characterization Window: Summary - Memory Consumption Window: Summary - Memory Usage Window: Summary - Platform Power Analysis
GPU Metrics Reference x
ALU0 Active ALU0 Instructions ALU1 Active ALU1 Instructions ALU2 Active ALU2 Instructions ALU0 and ALU1 Active ALU0 and ALU2 Active Average Time Computing Threads Started Computing Threads Started, Threads/sec CPU Time EU 2 FPU Pipelines Active EU Array Active EU Array Idle EU Array Stalled/Idle EU Array Stalled EU IPC Rate EU Send pipeline active EU Threads Occupancy Host to GPU Memory Read Bandwidth Host-to-GPU Memory Write Bandwidth Global GPU EU Array Usage GPU L3 Bound GPU L3 Miss Ratio GPU L3 Misses GPU L3 Misses, Misses/sec GPU Memory Read Bandwidth, GB/sec GPU Memory Texture Read Bandwidth, GB/sec GPU Memory Write Bandwidth, GB/sec GPU Texel Quads Count, Count/sec GPU Utilization Instance Count L3 Read Bandwidth L3 Write Bandwidth L3 Sampler Bandwidth, GB/sec L3 Shader Bandwidth, GB/sec LLC Miss Rate due GPU Lookups LLC Miss Ratio due GPU Lookups Local Maximum GPU Utilization Occupancy PS EU Active % PS EU Stall % Ratio to Max Bandwidth, % Ratio to Max Bandwidth, % Ratio to Max Bandwidth, % Render/GPGPU Command Streamer Loaded Samples Blended Samples Killed in PS, pixels Samples Written Sampler Busy Sampler Is Bottleneck Shared Local Memory Read Bandwidth, GB/sec Shared Local Memory Write Bandwidth, GB/sec SIMD Width Stack-to-stack Incoming Bandwidth Stack-to-stack Outgoing Bandwidth System Memory Read Bandwidth System Memory Write Bandwidth Size Total, GB/sec Total Time Typed Memory Read Bandwidth, GB/sec Typed Memory Write Bandwidth, GB/sec Typed Reads Coalescence Typed Writes Coalescence Untyped Memory Read Bandwidth, GB/sec Untyped Memory Write Bandwidth, GB/sec Untyped Reads Coalescence Untyped Writes Coalescence VS EU Active VS EU Stall
OpenCL™ Kernel Analysis Metrics Reference x
Computing Task Total Time Instance Count SIMD Width SIMD Utilization Work Size
Energy Analysis Metrics Reference x
Available Core Time C-State D0ix States DRAM Self Refresh Energy Consumed (mJ) Idle Wake-ups P-State S0ix States Temperature Timer Resolution Total Time in C0 State Total Time in Non-C0 States Total Time in S0 State Total Wake-up Count Wake-ups Wake-ups/sec per Core
User Guide

GPU OpenCL™ Application Analysis

If you identified with the Intel® VTune™ Profiler that your application is GPU-bound and your application uses OpenCL™ software technology, you may enable the Trace GPU Programming APIs configuration option for your custom analysis to identify how effectively your application uses OpenCL kernels. By default, this option is enabled for the GPU Compute/Media Hotspots and GPU Offload analyses. To explore the performance of your OpenCL application, use the GPU Compute/Media Hotspots viewpoint.

Follow these steps to explore the data provided by the VTune Profiler for OpenCL application analysis:

  1. Explore summary statistic:

    • Analyze GPU usage.

    • Identify why execution units (EUs) were stalled or idle.

    • Identify OpenCL kernels overutilizing both Floating Point Units (FPUs).

  2. Analyze hot GPU OpenCL kernels.

  3. Correlate OpenCL kernels data with GPU metrics.

  4. Explore the computing queue.

  5. Analyze source and assembly code.

Explore Summary Statistics

Start your data analysis with the Summary window that provides application-level performance statistics. Typically, you focus on the primary baseline, which is the Elapsed Time metric that shows the total time your target ran:

You can correlate this data with the GPU Time used by GPU engines while your application was running:

If the GPU Time takes a significant portion of the Elapsed Time (95.6%), it clearly indicates that the application is GPU-bound. You see that 94.4% of the GPU Time was spent on the OpenCL kernel execution.

For OpenCL applications, the VTune Profiler provides a list of OpenCL kernels with the highest execution time on the GPU:

Mouse over the flagged kernels to learn what kind of performance problems were identified during their execution. Clicking such a kernel name in the list opens the Graphics window grouped by computing tasks, sorted by the Total Time, and with this kernel selected in the grid.

Depending on the GPU hardware events preset you used during the analysis configuration, the VTune Profiler explores potential reasons for stalled/idle GPU execution units and provides them in the Summary. For example, for the Compute Basic preset, you may analyze GPU L3 Bandwidth Bound issues:

Or potential occupancy issues:

In this example, EU stalls are caused by GPU L3 high bandwidth. You may click the hottest kernels in the list to switch to the Graphics view, drill down to the Source or Assembly views of the selected kernel to identify possible options for cache reuse.

If your application execution takes more than 80% of collection time heavily utilizing floating point units, the VTune Profiler highlights such a value as an issue and lists the kernels that overutilized the FPUs:

You can switch to the Timeline pane on the Graphics tab and explore the distribution of the GPU EU Instructions metric that shows the FPU usage during the analysis run:

Analyze Hot GPU OpenCL Kernels

To view detailed information about all OpenCL kernels running on the GPU, switch to the Graphics window. By default, the grid data is grouped by Computing Task / Instance that shows Compute tasks only. Data collected for program units outside any OpenCL computing tasks are attributed to the [Outside any task] entry.

In the Computing Task columns explore the overall time a kernel ran on the GPU and the average time for a single invocation (corresponding to one call of clEnqueueNDRangeKernel ), working group sizes, as well as averaged GPU hardware metrics collected for a kernel. Hover over a metric column header to read the metric description. If a metric value for a computing task exceeds a threshold set up by Intel architects for the metric, this value is highlighted in pink, which signals a performance issue. Hover over such a value to read the issue description.

In the example below, the Accelerator_Intersect kernel took the most time to execute (53.398s). The GPU metrics collected for this workload show high L3 Bandwidth usage spent in stalls when executing this kernel. For compute bound code it indicates that the performance might be limited by cache usage.

Analyze and optimize hot kernels with the longest Total Time values first. These include kernels characterized by long average time values and kernels whose average time values are not long, but they are invoked more frequently than the others. Both groups deserve attention.

If a kernel instance used the OpenCL 2.0 Shared Virtual Memory (SVM), the VTune Profiler detects it and, depending on your hardware, displays the SVM usage type as follows:

  • Coarse-Grained Buffer SVM: Sharing occurs at the granularity of regions of OpenCL buffer memory objects. Cross-device atomics are not supported.

  • Fine-Grained Buffer SVM: Sharing occurs at the granularity of individual loads and stores within OpenCL buffer memory objects. Cross-device atomics are optional.

  • Fine-Grained System SVM: Sharing occurs at the granularity of individual loads/stores occurring anywhere within the host memory. Cross-device atomics are optional.

Every clCreateKernel results in a line in the Compute category. If two different kernels with the same name (even from the same source) were created with two clCreateKernel calls (and then invoked through two or more clEnqueueNDRangeKernel ), two lines with the same kernel name appear in the table. If they are enqueued twice with a different global or local size or different sets of SVM arguments, they are also listed separately in the grid.

Correlate OpenCL Kernels Data with GPU Metrics

In the Graphics window, explore the Timeline pane > Platform tab to analyze OpenCL kernels execution over time.

OpenCL APIs (for example, clWaitForEvents) show up on the Thread area as tasks:

Correlate GPU metrics and OpenCL kernels data:

NOTE:

GPU hardware metrics are available if you enabled the Analyze Processor Graphics events option for Intel® HD Graphics or Intel® Iris® Graphics. To collect these metrics, make sure to set up your system for GPU analysis.

You may find it easier to analyze your OpenCL application by exploring the GPU hardware metrics per GPU architecture blocks. To do this, choose the Computing Task grouping level in the Graphics window, select an OpenCL kernel of interest and click the Memory Hierarchy Diagram tab in the Timeline pane. VTune Profiler updates the architecture diagram for your platform with performance data per GPU hardware metrics for the time range the selected kernel was executed.

Currently this feature is available starting with the 4th generation Intel® Core™ processors and the Intel® Core™ M processor, with a wider scope of metrics presented for the latter one.

NOTE:

You can right-click the Memory Hierarchy Diagram, select Show Data As and choose a format of metric data representation:

  • Total Size

  • Bandwidth (default)

  • Percent of Bandwidth Maximum Value

Explore the Computing Queue

To view details on OpenCL kernels submission, in particular distinguish the order of submission and execution, and analyze the time spent in the queue, zoom in and explore the Computing Queue data in the Timeline pane. You can click a kernel task to highlight the whole queue to the execution displayed at the top layer. Kernels with the same name and size show up in the same color.

VTune Profiler displays kernels with the same name and size in the same color. Synchronization tasks are marked with vertical hatching . Data transfers, OpenCL routines responsible for transferring data from the host system to a GPU, are marked with cross-diagonal hatching .

NOTE:

In the Attach mode if you attached to a process when the computing queue is already created, VTune Profiler will not display data for the OpenCL kernels in this queue.

Analyze Source and Assembly Code

You may select a computing task of interest in the grid view, double-click it to open the Source/Assembly window and analyze the code for the selected kernel (with source files available).

Analyze the assembler code provided by your compiler for the OpenCL kernel, estimate its complexity, identify issues, match the critical assembly lines with the affected source code, and optimize, if possible. For example, if you see that some code lines were compiled into a high number of assembly instructions, consider simplifying the source code to decrease the number of assembly lines and make the code more cache-friendly.

Explore GPU metrics data per computing task in the Graphics window and drill down to the Source/Assembly view to explore instructions that may have contributed to the detected issues. For example, if you identified the Sampler Busy or Stalls issues in the Graphics window, you may search for the send instructions in the Assembly pane and analyze their usage since these instructions often cause frequent stalls and overload the sampler. Each send/sends instruction is annotated with comments in square brackets that show a purpose of the instruction, such as data reads/writes (for example, Typed/Untyped Surface Read), accesses to various architecture units (Sampler, Video Motion Estimation), end of a thread (Thread Spawner), and so on. For example, this sends instruction is used to access the Sampler unit:

0x408 260 sends (8|M0) r10:d r100 r8 0x82 0x24A7000 [Sampler, msg-length:1, resp-length:4, header:yes, func-control:27000]

NOTE:

  • Source/Assembly support is available for OpenCL programs with sources and for kernels created with IL (intermediate language), if the intermediate SPIR-V binary was built with the -gline-tables-only -s <cl_source_file_name> option.

  • The Source/Assembly analysis is not supported for the source code using the #line directive.

  • If your OpenCL kernels use inline functions, you can enable the Inline Mode filter bar option to view inline functions in the grid and analyze them in the Source view.

Parent topic: GPU Application Analysis on Intel® HD Graphics and Intel® Iris® Graphics
Level Two Title