While PTI_VIEW_DEVICE_GPU_KERNEL is enabled, SYCL event profiling returns zero kernel durations

Abierto
  • ¿Algo no está bien?
#129 0 comentarios 0 reacciones 0 asignados Ver en GitHub

Nadie ha tomado este issue todavía.

Evaluación

Dificultad
4/5
Tiempo estimado
3-5 días
Aptitud para principiantes
55/100
Tipo de issue
Error
Claridad
Bien especificado
Estado de actividad
Activo
Stack tecnológico
cpp, linux
Área
observability, performance

Línea de trabajo

Start by building and running the provided reproducer with SYCL event profiling enabled, then compare the event durations with PTI kernel records across the three phases. The issue reports zero SYCL durations only while PTI_VIEW_DEVICE_GPU_KERNEL is enabled, plus increased host submission time. Done means identifying and fixing the interaction so SYCL event profiling remains correct while the view is collecting.

Escrito por el modelo de indexación a partir del texto del issue.

Descripción

Describe the bug

With the GPU kernel view enabled, command_end - command_start from SYCL event profiling (a queue created with property::queue::enable_profiling) is 0 for every kernel the program submits. PTI's own kernel records are correct: their durations match what SYCL event profiling gives with PTI off. After ptiViewDisable, SYCL event profiling works again. An application that times its own kernels with SYCL events gets zeros while a profiler built on PTI is collecting.

Also visible in the same runs: each launch takes 7-15 us more host time while the view is enabled.

To reproduce

// Kernel durations from PTI's GPU kernel view against SYCL event profiling of the same kernel.
// One small kernel (y = a*x + y over 2^20 floats) is submitted 200 times back to back, first with
// PTI off (durations from SYCL event profiling), then with PTI_VIEW_DEVICE_GPU_KERNEL enabled
// (durations from PTI's records, and SYCL events alongside). Prints medians and wall time per call.
#include <sycl/sycl.hpp>
#include <pti/pti_view.h>
#include <algorithm>
#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <vector>

static std::vector<double> pti_us;
static void requested(unsigned char** buf, size_t* size) {
  *size = 1 << 22;
  *buf = static_cast<unsigned char*>(std::aligned_alloc(64, *size));
}
static void completed(unsigned char* buf, size_t, size_t used) {
  pti_view_record_base* rec = nullptr;
  while (ptiViewGetNextRecord(buf, used, &rec) == PTI_SUCCESS && rec) {
    if (rec->_view_kind == PTI_VIEW_DEVICE_GPU_KERNEL) {
      auto* k = reinterpret_cast<pti_view_record_kernel*>(rec);
      if (std::strstr(k->_name, "axpy")) pti_us.push_back((k->_end_timestamp - k->_start_timestamp) / 1000.0);
    }
  }
  std::free(buf);
}
static double median(std::vector<double> v) {
  std::sort(v.begin(), v.end());
  return v.empty() ? 0 : v[v.size() / 2];
}
class axpy;

int main(int argc, char** argv) {
  sycl::queue q{sycl::gpu_selector_v, sycl::property::queue::enable_profiling{}};
  const size_t n = argc > 1 ? std::strtoul(argv[1], nullptr, 0) : 1 << 20;
  float* x = sycl::malloc_device<float>(n, q);
  float* y = sycl::malloc_device<float>(n, q);
  q.fill(x, 1.0f, n).wait();
  q.fill(y, 2.0f, n).wait();
  auto submit = [&] {
    return q.parallel_for<axpy>(sycl::range<1>(n), [=](sycl::id<1> i) { y[i] = 0.5f * x[i] + y[i]; });
  };
  auto phase = [&](const char* name) {
    for (int i = 0; i < 20; ++i) submit();
    q.wait();
    std::vector<sycl::event> ev;
    auto t0 = std::chrono::steady_clock::now();
    for (int i = 0; i < 200; ++i) ev.push_back(submit());
    q.wait();
    double wall = std::chrono::duration<double, std::micro>(std::chrono::steady_clock::now() - t0).count() / 200;
    std::vector<double> d;
    for (auto& e : ev)
      d.push_back((e.get_profiling_info<sycl::info::event_profiling::command_end>() -
                   e.get_profiling_info<sycl::info::event_profiling::command_start>()) / 1000.0);
    std::printf("%-22s wall %.1f us/call, SYCL event median %.1f us", name, wall, median(d));
  };
  std::printf("device: %s\n", q.get_device().get_info<sycl::info::device::name>().c_str());
  phase("PTI off");
  std::printf("\n");
  ptiViewSetCallbacks(requested, completed);
  ptiViewEnable(PTI_VIEW_DEVICE_GPU_KERNEL);
  phase("PTI kernel view on");
  ptiViewDisable(PTI_VIEW_DEVICE_GPU_KERNEL);
  ptiFlushAllViews();
  std::printf(", PTI kernel median %.1f us (%zu records)\n", median(pti_us), pti_us.size());
  phase("after ptiViewDisable");
  std::printf("\n");
}

icpx -fsycl -O2 pti_kernel_time.cpp -lpti_view -o pti_kernel_time
./pti_kernel_time 1048576

Output on an Intel Arc A770, three phases of 200 back-to-back submissions:

PTI off                wall 30.6 us/call, SYCL event median 29.4 us
PTI kernel view on     wall 40.9 us/call, SYCL event median 0.0 us, PTI kernel median 29.6 us (220 records)
after ptiViewDisable   wall 30.9 us/call, SYCL event median 29.2 us

Across kernel sizes (argument = element count), SYCL event median with PTI off / with the kernel view on / PTI's median, and wall time per call off / on:

elements SYCL event, PTI off SYCL event, view on PTI record wall off / on
1,024 2.0 us 0.0 2.0 us 4.7 / 15.0 us
65,536 3.2 us 0.0 3.2 us 4.5 / 15.5 us
262,144 9.4 us 0.0 9.4 us 10.3 / 21.2 us
524,288 16.1 us 0.0 16.2 us 17.2 / 24.2 us
4,194,304 115.3 us 0.0 116.0 us 116.5 / 131.0 us

Expected behavior

SYCL event profiling reports kernel durations whether or not PTI's kernel view is enabled.

Environment

  • PTI 1.2.1 (libpti_view.so.1.2.1 from the intel_pti 1.2.1 wheel; headers from release/pti-1.2 at 25a77ed), with oneAPI 2026.0's SYCL runtime; the numbers above are from this build
  • Also on PTI 1.1.0 (the intel_pti 1.1.0 wheel, headers from tag pti-1.1.0), with oneAPI 2026.0's SYCL runtime and with the intel-sycl-rt 2026.1.2 runtime from the PyTorch XPU wheels (libsycl.so.9, libur_loader.so.0 and libpti_view.so.1 all resolved from the wheels' directory): the same zero durations
  • Built with the Intel oneAPI DPC++/C++ Compiler 2026.0.0
  • Measured on one GPU (Arc A770, xe kernel driver); not tried on other Intel GPUs or on i915
  • Intel Arc A770, Linux 7.2.14 with the xe kernel driver, intel-compute-runtime 26.35.39758.10 (Level Zero driver 1.17.39758), level-zero-loader 1.32.0
Lenguaje dominante
C++
Estrellas
278
Forks
75
Sin PR fusionados en 30 d

Preparar el entorno

Primeros pasos

  1. Lee el issue completo y luego la guía de contribución del proyecto.
  2. Comenta en el issue que vas a ocuparte — evita que dos personas hagan lo mismo.
  3. Haz un fork del repositorio y trabaja en una rama.
  4. Abre un pull request que haga referencia al número del issue.

Más de intel/pti-gpu

Todos los issues de intel/pti-gpu

Issues similares

Más issues de C++

Recibe los nuevos issues en tu correo

Un resumen breve de issues de GitHub para principiantes.