Kernels not running on GPU: cl::sycl::gpu_selector selector not work

[babouFomb]

Kernels not running on GPU: cl::sycl::gpu_selector selector not work

[DuncanMcBain]

Hi,
the reason that you are not seeing kernels run on the GPU is because we output kernels to be executed on the GPU in the SPIR 1.2 or SPIR-V formats (that said, the SPIR-V backend is very experimental, and for OpenCL only). Currently, nvidia's OpenCL implementation does not support either SPIR 1.2 or SPIR-V, and therefore we can't run on nvidia.

On a more technical level, what is happening is the default selectors circulated with ComputeCpp query the devices on the system for SPIR support (i.e. look for the extension "cl_khr_spir"). If no devices on the system report that they support it, as is the case here, ComputeCpp falls back to the host device. When you explicitly request a GPU, it fails to find such a device, so an exception is thrown. We decided we wanted the SDK to run under most circumstances we could think of (and it helps with travis testing, too).

I can recommend trying out the pocl implementation. Recently, they added support for nvidia devices, though it is somewhat incomplete. We don't officially support it, but I have managed to run some of the samples on my home PC that has a 780 in it. There is one issue: pocl reports that it doesn't support the spir extension either, but they definitely do! This requires modification, but you can force selection of any device you like (for example, by requesting all devices on the system, then filtering on name or similar). There is also a custom device selection interface as described in the "custom device selector" sample.

As a last note, we recently updated ComputeCpp again, you should now be able to download 0.3.1 from the website. I don't believe there are any changes required to the SDK for this release, even though there is no 0.3.1 tag yet.

If you've any other questions or problems, update this issue and I'll do my best to help.

[DuncanMcBain]

I did little more digging this week, as it turns out it's relatively simple to turn on SPIR support in pocl - though I should check that what I'm doing is valid, as I am experiencing some failures and would like to check that it's not me first! If I verify that using pocl with a system clang is basically equivalent to using the SPIR generator from Khronos, then I will post my (very simple) patch here and let pocl know. If it turns out that the clang SPIR generator is very necessary for proper operation, then back to the drawing board I suppose!

[DuncanMcBain]

Hi @babouFomb,
Did you manage to make any progress on this?

[keryell]

@DuncanMcBain haha the PoCL path seems just insane. :-)

I would love having feedback on this at the next Khronos meeting during SYCL and/or SPIR sesions.

[DuncanMcBain]

It's honestly not so bad! I've not yet made an issue, but my opinion is that pocl should unconditionally report SPIR support, since they are always capable of consuming it. If that were the case then pocl/nvidia would "just work" with the SDK.

The problem is that ComputeCpp discards devices it cannot use in the default selector, so unless you make some changes the samples won't find any devices. I don't think we should change the SDK for this case though.

What sort of feedback were you looking for in the next meeting?

[keryell]

it think it would be interesting to have a few slides describing this interesting ComputeCpp+PoCL use case, even if it requires selecting explictly the device.

[keryell]

I was thinking to the ComputeCPP→PoCL→CUDA/PTX use case.
There is otherwise some support for SPIR in PoCL, even I had to do some change to support SPIR 2.0 encoded in Clang 3.9 for example (SPIR "de-facto"...).

[keryell]

I proposed my SPIR "de-facto" ingestion in pocl/pocl#529
I do not know if it solves your issue or not.

[DuncanMcBain]

That might make things easier for sure, though I was ill yesterday and haven't yet tested the clang spir issue. Hopefully I will get to that today. Thanks!

[DuncanMcBain]

Closing due to lack of action.