Conversation
PeterTh
force-pushed
the
khr_work_item_queries
branch
from
92330e9 to
9342b37
Compare
fknorr
left a comment
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. |
|
Switched to thread locals now. |
PeterTh
force-pushed
the
khr_work_item_queries
branch
from
44599fa to
81e0323
Compare
fknorr
left a comment
fknorr
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.
Is this a clang-format change?
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 ;))
2 participants
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.