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

Đang mở
  • Có gì không ổn?
#129 0 bình luận 0 reaction 0 người được giao Xem trên GitHub

Chưa có ai nhận issue này.

Đánh giá

Độ khó
4/5
Thời gian dự kiến
3-5 ngày
Mức phù hợp với người mới
55/100
Loại issue
Lỗi
Độ rõ ràng
Đặc tả rõ ràng
Mức độ hoạt động
Sôi nổi
Công nghệ
cpp, linux
Lĩnh vực
observability, performance

Hướng nghiên cứu

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.

Do mô hình lập chỉ mục viết ra từ nội dung của issue.

Mô tả

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
Ngôn ngữ chính
C++
Star
278
Fork
75
Không có pull request nào được merge trong 30 ngày

Chuẩn bị môi trường

Bắt đầu từ đâu

  1. Đọc hết issue, rồi đọc hướng dẫn đóng góp của dự án.
  2. Bình luận trên issue rằng bạn sẽ nhận — tránh hai người làm cùng một việc.
  3. Fork repository và làm thay đổi trên một nhánh.
  4. Mở pull request có tham chiếu số hiệu của issue.

Issue khác của intel/pti-gpu

Tất cả issue của intel/pti-gpu

Issue tương tự

Thêm issue về C++

Nhận issue mới trong hộp thư của bạn

Bản tóm tắt ngắn những issue GitHub phù hợp với người mới.