To profile the Intel GPU kernels with the SYCL event profiling time stamp instead of barrier time diff.

[chengjunlu]

It is more accurate to use the SYCL event profiling time stamp after the immediate command list has been enabled.

To use the SYCL event profiling to capture the Intel GPU kernel time.

[ESI-SYD]

I think we can enable Triton-benchmark CI, especially benchmarks files changed, to ensure the basic function.
https://github.com/intel/intel-xpu-backend-for-triton/blob/llvm-target/.github/workflows/triton-benchmarks.yml#L3

on:
  workflow_dispatch:
    inputs:
      runner_label:
        description: Runner label, keep empty for default
        type: string
        default: ""
  pull_request:
    branches:
      - llvm-target
    paths:
      - 'benchmarks/**'  
  schedule:
    - cron: "5 23 * * *"

[chengjunlu]

I think we can enable Triton-benchmark CI, especially benchmarks files changed, to ensure the basic function. https://github.com/intel/intel-xpu-backend-for-triton/blob/llvm-target/.github/workflows/triton-benchmarks.yml#L3

on:
  workflow_dispatch:
    inputs:
      runner_label:
        description: Runner label, keep empty for default
        type: string
        default: ""
  pull_request:
    branches:
      - llvm-target
    paths:
      - 'benchmarks/**'  
  schedule:
    - cron: "5 23 * * *"

We need Pavel to review the side effect in running the micro benchmark for each PR. And let's use a sperate PR for review that.

[etiotto]

@chengjunlu why can't we use the same do_bench as Triton uses to compute the benchmark timing ?
What is the difference between triton.testing.do_bench and benchmark_suit.do_bench. Note that the former now uses event profiling torch.xpu.Event now.

[chengjunlu]

The torch.xpu.Event is based on the SYCL barrier. And we measure the approximate kernel time by diff the start barrier time stamp and the end barrier time stamp.
The method is less accurate after the SYCL runtime default uses the immediate command list to reduce the latency. In the experimental, we found there could be 10% difference between the approximate time and the actual time for the kernel.

In the micro-benchmark tools, we use the SYCL event profiling time returned by the SYCL launcher interface which is the actual time for the kernel for short term.

I think later we need to support the Triton Proton to profile the Triton XPU kernel for long term.

[etiotto]

@chengjunlu why do we keep commented out lines?