Uh oh!
There was an error while loading. Please reload this page.
Implement SYCL_KHR_WORK_ITEM_QUERIES - #47
Conversation
92330e9 to
9342b37Compare
fknorr
left a comment
There was a problem hiding this comment.
I'm not entirely sure about the thread safety around the new g_ globals. We do keep a system lock around every submit, but the remaining global state in schedule.cc are thread locals (although this was the case before the system lock was introduced). If the code is sound as-is, maybe having the new variables be thread-local is still a bit easier to reason about.
You mean as it relates to threading by the user program? In any case, I see no problem with making them thread local, so I'll do that. Since you're here (😛), in the new unit test introduced in this PR there's this. |
PeterTh
commented
Sep 1, 2025
Switched to thread locals now. |
44599fa to
81e0323Comparefknorr
left a comment
There was a problem hiding this comment.
Oh, you're right about the global range, [0] being zero is not useful.
LGTM!
| const auto global_id | ||
| = range.get_offset() + (group_id * sycl::id<Dimensions>(local_range)) + local_id; |
There was a problem hiding this comment.
Yes, that's how it is formatted for me right now. (I wouldn't write that on purpose ;))
This adds support for the SYCL_KHR_WORK_ITEM_QUERIES extension, as per KhronosGroup/SYCL-Docs#682.
Although SimSYCL is a library-only implementation, due to its single-threaded nature this can be implemented via thread locals that are set every time a kernel function context is scheduled in.
The actual implementation is a bit more convoluted than I'd like, but it does provide decent error messages / developer experience.