[SYCL] Added a new test case for sub-group attributes #537
Conversation
| } \ | ||
| }; | ||
|
|
||
| KERNEL_FUNCTOR_WITH_SIZE(1); |
There was a problem hiding this comment.
Overall this test looks good. I have only one concern that it might not work properly on any OpenCL RT with sub-groups support.
The problem is that all kernel functions described here are going to be included into a one device program which will be compiled later by a device compiler - and this might cause errors because apparently device doesn't have to support all sub-group sizes listed in here and it is legal to fail compilation if requested sub-group size is not supported.
Probably it is a bug in our implementation, that build_with_kernel_type actually performs build of all kernels instead of just one, but I don't know SYCL spec so good to say for sure.
There was a problem hiding this comment.
I remembered there's a ticket created by Evgeniy describing the issue of building all kernels.
|
@tianshilei1992, could you take a look at failing lit tests, please? |
Sorry for not solving it. I’ll do it in this week. |
| cl::sycl::program Prog(Queue.get_context()); | ||
|
|
||
| // Store the `cl::sycl::kernel` into a vector because `cl::sycl::kernel` | ||
| // doesn't have default constructor |
There was a problem hiding this comment.
Thank you for the comment.
It is not clear what you are trying to do here...
| throw feature_not_supported("sub-group size is not supported"); | ||
| } | ||
|
|
||
| auto Kernel = TheKernel[0]; |
There was a problem hiding this comment.
I see... You are emulating a C++17 std::optional since you are in C++11... :-(
On the other hand you could make a specification proposal to add a default constructor to a cl::sycl::kernel and a member function to test if the kernel is valid... :-)
There was a problem hiding this comment.
Not sure whether it makes sense to request it. :)
Well, mainly working on fixing an issue that could only be reproduced on GPU. But I don't have GPU system on my hand. Waiting for it. :) |
|
Report the issue at #640. |
|
I disabled the test on GPU intentionally as it fails because all kernels are actually built when calling building function to just one kernel. |
Signed-off-by: Shilei Tian <[email protected]>
Signed-off-by: Shilei Tian [email protected]