ComputeCpp Professional Edition has built-in support to output a JSON file that contains profiling information of the runtime and the underlying OpenCL implementation. This JSON file is conformant with the Chrome Trace Events format and can be loaded in any Chrome installation by accessing the URL chrome://tracing.

The objective of this document is to describe how to activate profiling for a ComputeCpp application using the configuration file.

A configuration file may be used to configure how the ComputeCpp runtime will perform certain operations.

In order to enable profiling of the runtime, create a file and add the option enable_profiling = true. Point an environment variable called COMPUTECPP_CONFIGURATION_FILE to this file and run the application. The following is a single command to create a file called sycl_config.txt and export the environment variable to enable profiling:

• Windows:

echo "enable_profiling = true" > sycl_config.txt; set COMPUTECPP_CONFIGURATION_FILE=sycl_config.txt

• Linux:

echo "enable_profiling = true" > sycl_config.txt; export COMPUTECPP_CONFIGURATION_FILE=sycl_config.txt

Note that if COMPUTECPP_CONFIGURATION_FILE points to a file that doesn't exists, the runtime will fail to initialize.

The following is a list of all the configuration options available to change how the profiler backend behaves:

By default, when the application finishes, the runtime will write the JSON file in the current working directory, usually the same directory as the binary of the application, in the format [executable_name]_[current_date].json.

This behaviour can be changed by setting the environment variable COMPUTECPP_PROFILING_OUTPUT. If this is set, the runtime will use the value of this variable as the output file. The file doesn't need to exist but the application must have permissions to create and write it.

When running ComputeCpp applications in Intel GPU's that match one of the following architectures, the profiler can display performance counters for the ComputeBasic metrics set.

• Intel(R) Processors with Gen11 graphics devices (formerly Icelake),
• Intel(R) Processors with Gen9 graphics devices (formerly Skylake, Kaby Lake, Apollo Lake/Broxton, Gemini Lake, Coffee Lake),
• Intel(R) Processors with Gen8 graphics devices (formerly Broadwell),
• Intel(R) Processors with Gen7.5 graphics devices (formerly Haswell).

No extra configuration is required, as long as the system supports Intel's Metrics Discovery API.

When running ComputeCpp applications on an ARM Mali device with either the Midgard, Valhall, or Bifrost architecture, the profiler can display a set of performance counters.

The following is an example of the Google Chrome Visualizer that displays a execution of BabelStream in an Intel GPU.

Note at the bottom the performances counters. Many more counters are available and they can be accessed by scrolling the window.

This is a Professional Edition feature. Using codeplay::profiling::profiling_zone has no effect in ComputeCpp CE.

ComputeCpp supports user-defined profiling zones. This feature allows developers to annotate the source code and display extra information in the profiler.

This can be used to mark sections of interest to be displayed as API calls.

To use this feature, just include SYCL/codeplay/profiling_user.h and start annotating the source code. Annotations can be created anywhere using the cl::sycl::codeplay::profiling::profiling_zone object. Refer to example for more information on the usage.

Here is a simple example:

#include <CL/sycl.hpp>

#include <SYCL/codeplay/profiling_user.h>

#include <array>
#include <iostream>

using namespace cl::sycl;

constexpr access::mode sycl_read = access::mode::read;
constexpr access::mode sycl_write = access::mode::write;

template <typename T>
class SimpleVadd;

template <typename T, size_t N>
void simple_vadd(const std::array<T, N>& VA, const std::array<T, N>& VB,
                 std::array<T, N>& VC) {
  // Creates a zone that spawns the execution of the simple_vadd function
  codeplay::profiling::profiling_zone simpleVAddZone("simple_vadd");
  queue deviceQueue;
  range<1> numOfItems{N};
  buffer<T> bufferA(VA.data(), numOfItems);
  buffer<T> bufferB(VB.data(), numOfItems);
  buffer<T> bufferC(VC.data(), numOfItems);

  deviceQueue.submit([&](cl::sycl::handler& cgh) {
    // This allows you to know exactly when the submit happened
    codeplay::profiling::profiling_zone submitVAddZone("submit simple_vadd");
    auto accessorA = bufferA.template get_access<sycl_read>(cgh);
    auto accessorB = bufferB.template get_access<sycl_read>(cgh);
    auto accessorC = bufferC.template get_access<sycl_write>(cgh);

    auto kern = [=](id<1> wiID) {
      accessorC[wiID] = accessorA[wiID] + accessorB[wiID];
    };
    cgh.parallel_for<class SimpleVadd<T>>(numOfItems, kern);
  });
}

int main() {
  // Creates a zone that spawns the whole application
  codeplay::profiling::profiling_zone mainZone("main");
  const size_t array_size = 4;
  std::array<cl_int, array_size> A = {{1, 2, 3, 4}}, B = {{1, 2, 3, 4}}, C;
  std::array<cl_float, array_size> D = {{1.f, 2.f, 3.f, 4.f}},
                                   E = {{1.f, 2.f, 3.f, 4.f}}, F;
  simple_vadd(A, B, C);
  simple_vadd(D, E, F);
}

Running this application with the JSON profiler enabled will produce the following output.

WARNING: User defined zones inside kernels are not supported at the moment.

Note how the user-defined zones are naturally blended with other API calls. This allows grouping regions of interested.

A zone is created with an object of class codeplay::profiling::profiling_zone is created and closed when the object is destroyed.

User zones can be nested. The following example creates four user-defined zones, however, zones 2 and 4 are nested under zone 1 and zone 3 is nested under zone 2.

profiling_zone zone1("zone1");
{
  profiling_zone zone2("zone2");
  {
    profiling_zone zone3("zone3");
  }
  profiling_zone zone4("zone4");
}

All profilers in ComputeCpp support this feature.