Skip to content

sycl: fix performance regression by caching getenv #23381

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.

Sign up for GitHub

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

Closed
iOptimizeThings wants to merge 1 commit into from
Closed

sycl: fix performance regression by caching getenv #23381

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion ggml/src/ggml-sycl/common.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -86,7 +86,8 @@ int64_t downsample_sycl_global_range(int64_t accumulate_block_num, int64_t block

#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO
static bool ggml_sycl_use_level_zero_device_alloc(sycl::queue &q) {
return ggml_sycl_get_env("GGML_SYCL_ENABLE_LEVEL_ZERO", 1) &&
static const int enable_level_zero = ggml_sycl_get_env("GGML_SYCL_ENABLE_LEVEL_ZERO", 1);
return enable_level_zero &&
Copy link
Copy Markdown
Contributor

@sanmai sanmai May 20, 2026 •
edited
Loading

Choose a reason for hiding this comment

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

It might be better to check for env var in ggml_sycl_init() once #23097 goes through (I'm hesitant to tweak it more as it is approved and merge ready)

Copy link
Copy Markdown
Contributor

@arthw arthw May 20, 2026

Choose a reason for hiding this comment

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

Yes.
This function ggml_sycl_use_level_zero_device_alloc(sycl::queue &q) can be replaced by g_ggml_sycl_enable_level_zero which has been set.

Copy link
Copy Markdown
Contributor

@arthw arthw May 20, 2026

Choose a reason for hiding this comment

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

The original code should be updated as this PR or suggestion.
But even if use my suggestion, it impacts the performance less.

Copy link
Copy Markdown
Author

@iOptimizeThings iOptimizeThings May 20, 2026

Choose a reason for hiding this comment

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

Makes sense using g_ggml_sycl_enable_level_zero has even less overhead. I'll update ggml_sycl_use_level_zero_device_alloc to use the global. Should I rebase on #23097 first since it's merge-ready, or is it fine to land this independently as the small fix?

Copy link
Copy Markdown
Author

@iOptimizeThings iOptimizeThings May 20, 2026

Choose a reason for hiding this comment

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

actually scratch that, g_ggml_sycl_enable_level_zero is less overhead but it starts at 0 and only gets set inside ggml_check_sycl(), and the comment under this function notes large buffers can be allocated before that runs, so gating on it directly would make those early allocs skip the level zero path and fall back to the DMA-buf/TTM staging this was added to avoid. i checked #23097 too and it keeps the assignment in ggml_check_sycl(), so that early window doesn't change. the static const cache keeps it init-order independent while still dropping the per-alloc getenv, so i think it covers both — good to keep that approach, or is the global guaranteed set before those early allocs somewhere i'm not seeing?

Copy link
Copy Markdown
Contributor

@sanmai sanmai May 20, 2026 •
edited
Loading

Choose a reason for hiding this comment

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

I went with #23399 (rebase-pending, so the diff is a mess) but there could be a better way.

Copy link
Copy Markdown
Author

@iOptimizeThings iOptimizeThings May 20, 2026

Choose a reason for hiding this comment

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

Makes sense, checking it during ggml_sycl_init() avoids the early allocation window entirely, cleaner than the static cache. Happy to close this in favor of #23399 if that's the direction.

Copy link
Copy Markdown
Contributor

@arthw arthw May 21, 2026

Choose a reason for hiding this comment

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

@iOptimizeThings
OK, it's great if you close this PR.
We could focus on the review of #23399. The comments will be moved to it.

Thank you for your sharing of this issue and solution!

q.get_device().is_gpu() &&
q.get_backend() == sycl::backend::ext_oneapi_level_zero;
}
Expand Down