Fix tests after enabling rounding-up optimizations when this_item call is used #699
Conversation
aobolensk
force-pushed
the
rounding_up_tests
branch
from
76cc486 to
7fd86f0
Compare
aobolensk
force-pushed
the
rounding_up_tests
branch
from
7fd86f0 to
d51fdab
Compare
| @@ -48,7 +46,8 @@ int main() { | |||
|
|
|||
| auto that_item = sycl::ext::oneapi::experimental::this_item<1>(); | |||
| results_acc[1] = that_item.get_id() == i; | |||
| results_acc[2] = that_item.get_range() == sycl::range<1>(n); | |||
| results_acc[2] = | |||
| that_item.get_range() == ((sycl::range<1>(n) + (31)) & ~(31)); | |||
There was a problem hiding this comment.
rounding up to the nearest divisible by 32. Formula: (x + 31) div 32
There was a problem hiding this comment.
But why 32? Why not 42?
My understanding is that you are trying to guess work group size here, but you can't do that in a reliable way. Moreover, I don't quite understand why get_range has changed when range rounding was re-enabled for that kernel (due to removed check for this_item call in intel/llvm@e4791d1).
According to 4.9.4.2.2. parallel_for invoke:
Below is an example of invoking a SYCL kernel function with parallel_for using a lambda function and passing a SYCL item parameter. In this case, both the global id and global range are queryable. This variant of parallel_for is designed for when it is necessary to query the global range of the index space being executed across.
free function queries spec doesn't say anything special about item returned through this_item, so I assume that it should return global range.
If our implementation of range rounding alters the range returned by item, then I would say we have a bug in that feature and it shouldn't be enabled by default as it is now. Alternatively, if we can detect cases when that feature can't be enabled, then we should do so. My understanding is that we used to detect presence of this_item to disable the feature in cases where it doesn't work correctly, but by some reason we removed that code.
There was a problem hiding this comment.
Tagging @Pennycook here for awareness as well
|
@aobolensk do we still need the PR? Could you please fix the issues or close it? |
Rounding-up optimizations were enabled if this_item API call is used in kernel in intel/llvm#5178