intel / intel/llvm

SYCL runtime: Severe host overhead in sycl::get_kernel_bundle

Open
#15,824 3 comments 0 reactions 0 assignees View on GitHub
confirmed performance
Dominant language
LLVM
Stars
1.5k
Forks
854
Avg merge
3d 17h
Merged PRs (30d)
137

Description

For platform compatible, we didn't use device max work group size to launch kernel, and switch to query specific max work group size for kernel by SYCL API. following is our code example
```
auto kid = ::sycl::get_kernel_id();
auto kbundle = ::sycl::get_kernel_bundle<::sycl::bundle_state::executable>(
ctx, {dev}, {kid});
::sycl::kernel k = kbundle.get_kernel(kid);
int max_work_group_size = k.get_info<::sycl::info::kernel_device_specific::work_group_size>(dev);
```
We found this usage takes much host overhead in application. we measured one kernel CPU performance here, each API name in table maps example code:

<!--table
{mso-displayed-decimal-separator:"\.";
mso-displayed-thousand-separator:"\,";}
@page
{margin:.75in .7in .75in .7in;
mso-header-margin:.3in;
mso-footer-margin:.3in;}
tr
{mso-height-source:auto;}
col
{mso-width-source:auto;}
br
{mso-data-placement:same-cell;}
td
{padding-top:1px;
padding-right:1px;
padding-left:1px;
mso-ignore:padding;
color:black;
font-size:11.0pt;
font-weight:400;
font-style:normal;
text-decoration:none;
font-family:"Aptos Narrow", sans-serif;
mso-font-charset:0;
mso-number-format:General;
text-align:general;
vertical-align:bottom;
border:none;
mso-background-source:auto;
mso-pattern:auto;
mso-protection:locked visible;
white-space:nowrap;
mso-rotate:0;}
-->

API | get_kernel_id | get_kernel_bundle | get_kernel | get_info
-- | -- | -- | -- | --
time (us) | 0.434 | 42.481 | 4.241 | 1.125

We also file internal jira to track this issue. Can you help evaluate this slow performance.

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.