Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Introduce optional profiling option for SPIRVRunner #2566

Open
wants to merge 3 commits into
base: main
Choose a base branch
from

Conversation

anmyachev
Copy link
Contributor

@anmyachev anmyachev commented Oct 24, 2024

Closes #2565

Usage: ./build/SPIRVRunner tensor_2 --enable-profiling

Output example:

...
Tensor output: [8192, 8192], Float (268435456 bytes)
Kernel execution time: 4.81872 ms
...

Profiling ref: https://github.com/intel/pti-gpu/blob/master/chapters/device_activity_tracing/DPCXX.md#how-to-use
Profiling ref with Kineto: pytorch/pytorch#52936 (comment)

Example with Kineto
  #include <torch/csrc/autograd/profiler_kineto.h>
  using namespace torch::autograd::profiler;
  ProfilerConfig cfg{
    ProfilerState::KINETO,
      false,
      false,
      false,
      false,
      false
  };
  std::set<torch::autograd::profiler::ActivityType> activities{torch::autograd::profiler::ActivityType::CPU, torch::autograd::profiler::ActivityType::XPU};
  prepareProfiler(cfg, activities);
  enableProfiler(cfg, activities);

  sycl_kernel_launch(stream, kernel, triton_args);

  auto result = disableProfiler();
  result->save("./some_local_file2.json");

@@ -352,8 +364,14 @@ at::Tensor launchKernel(sycl::queue stream, sycl::kernel kernel,
}
}

if (!triton_args.host_outbuffer.defined()) {
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

To avoid getting an incomprehensible error.

@anmyachev anmyachev marked this pull request as ready for review October 24, 2024 17:08
@@ -372,22 +390,43 @@ at::Tensor launchKernel(sycl::queue stream, sycl::kernel kernel,
return triton_args.host_outbuffer;
}

bool check_option_amoung_argv(int argc, char **argv, std::string option) {
Copy link
Contributor

@kballeda kballeda Oct 25, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I believe passing argc/argv to another function this way may not be a good idea.
we could use a parsing library to avoid this manual processing of various options, just as an example something like this (my past work).
https://github.com/Xilinx/Vitis_Libraries/blob/3c8a3d59fe76157ffa0b31fb267961bea2e148f5/data_compression/L3/demos/libzso/src/host.cpp#L271

we could use std::find instead of loop to search for the option

bool get_kernel_time =
check_option_amoung_argv(argc, argv, enable_profiling);
sycl::queue q;
if (get_kernel_time) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

as get_kernel_time is passed across launch(), sycl_launch_kernel() would it be possible to set it as part of constructor of struct ?

@@ -352,8 +364,14 @@ at::Tensor launchKernel(sycl::queue stream, sycl::kernel kernel,
}
}

if (!triton_args.host_outbuffer.defined()) {
std::string message = "Output tensor isn't configurated; \
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

message typo (NITS)

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

Introduce optional profiling option for SPIRVRunner
2 participants