SYCL Device Selector can return an invalid device

[Michoumichmich]

Hello,

My question concern device selectors. It seems to me that clang allows you to construct a device and a queue for a 'target' without having compiled for it.

Let's consider the following code :

#include <sycl/sycl.hpp>

class cuda_selector : public sycl::device_selector {
public:
    int operator()(const sycl::device &device) const override {
        return device.get_platform().get_backend() == sycl::backend::cuda && device.get_info<sycl::info::device::is_available>() ? 1 : -1;
    }
};

/**
 * Tries to get a queue from a selector else returns the host device
 */
static inline sycl::queue try_get_queue(const sycl::device_selector &selector) {
    auto exception_handler = [](const sycl::exception_list &exceptions) {
        for (std::exception_ptr const &e : exceptions) {
            try {
                std::rethrow_exception(e);
            }
            catch (sycl::exception const &e) {
                std::cout << "Caught asynchronous SYCL exception: " << e.what() << std::endl;
            }
            catch (std::exception const &e) {
                std::cout << "Caught asynchronous STL exception: " << e.what() << std::endl;
            }
        }
    };

    sycl::device dev;
    sycl::queue q;
    try {
        dev = sycl::device(selector);
        q = sycl::queue(dev, exception_handler);
        //if (dev.is_cpu() || dev.is_gpu()) q.single_task([]() {}).wait_and_throw(); //workaround
    }
    catch (...) {
        dev = sycl::device(sycl::host_selector());
        q = sycl::queue(dev, exception_handler);
        std::cout << "Warning: Expected device not found! Fall back on: " << dev.get_info<sycl::info::device::name>() << std::endl;
    }
    return q;
}

static inline void probe_queue(sycl::queue q){
    try{
        q.single_task([]() {}).wait_and_throw();
        std::cout << "Successfully ran on " << q.get_device().get_info<sycl::info::device::name>() << std::endl; 
    } catch (...){
        std::cout << "Something went wrong running on " << q.get_device().get_info<sycl::info::device::name>() << std::endl; 
    }
}

int main(){
    auto cuda_q = try_get_queue(cuda_selector{});
    auto cpu_q = try_get_queue(sycl::cpu_selector{});
    auto host_q = try_get_queue(sycl::host_selector{});
    probe_queue(cuda_q);
    probe_queue(cpu_q);
    probe_queue(host_q);
}

And the "makefile"

all:
	clang++ -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice,nvptx64-nvidia-cuda-sycldevice -sycl-std=2020 -std=c++20 -fsycl-unnamed-lambda main.cpp
cuda:
	clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice -sycl-std=2020 -std=c++20 -fsycl-unnamed-lambda main.cpp
cpu:
	clang++ -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice -sycl-std=2020 -std=c++20 -fsycl-unnamed-lambda main.cpp
host:
	clang++ -fsycl -sycl-std=2020 -std=c++20 -fsycl-unnamed-lambda main.cpp

Expected behaviour
We would expect try_get_queue(cuda_selector{}) to return a device that satisfies sycl::info::device::is_available and has a cuda backend. Turns out that compiling that code without nvptx64-nvidia-cuda-sycldevice still returns a queue attached to a Nvidia GPU. Running probing function on this queue results in CL_INVALID_BINARY (caught here) as there is no PTX code embedded in the binary, of course. The same behaviour happens when using the sycl::cpu_selector{} and not providing the target.

The program shown should only print Successfully ran on "Insert device compiled for".

Finally, there is also a bug which occurs with the target all. We should expect to run on the Host, the CPU and GPU, but it's not the case.

Workaround
Trying to run, in a try catch block, a dummy kernel such as

q.single_task([]() {}).wait_and_throw();

on the queue allows to check whether the device is "usable". See the commented line in the previous code.
This solution is not perfect if the device really is usable as it will block the caller. Furthermore, the time during which the caller will be blocked is potentially unknown and can depend on the device's and runtime's schedulers.

[Open] Solutions
The norm is vague about is_available, it states that it Returns true if the SYCL device is available and returns false if the device is not avail­able. One solution could use get_info<sycl::info::device::is_available>() and check the binary to see if at least some code is inside, before using the SYCL runtime to detect the devices. One couldn't imagine having the platform not listing the devices that cannot possibly be used as it would conflict with sycl-ls, I guess.

Any suggestions ?

[alexbatashev]

@Michoumichmich I think the observed behavior is correct. Consider the following situation:

You are writing an application, that doesn't really need much of SYCL. All you do is create a buffer, fill it with data, create a queue from some sort of device (doesn't really matter what device), and then pass the queue and the buffer to a library (oneMKL or something). In this scenario your application will not contain any kind of device image. But it is still legal for you to call host APIs on different devices. The library, on the other hand, may contain device image, compatible with your device, and will execute the necessary workload without an error.

Another case is when you want to use SYCL as a convenient C++ wrapper around low-level API. In that case you'll use the interoperability APIs, and it'll be your responsibility to provide necessary device binary.

One thing, that SYCL spec could improve though, is provide some sort of test for device image compatibility with target device for the device_image class. Tagging @keryell, @jbrodman, @gmlueck for that matter.

[gmlueck]

One thing, that SYCL spec could improve though, is provide some sort of test for device image compatibility with target device for the device_image class. Tagging @keryell, @jbrodman, @gmlueck for that matter.

An application can use is_compatible)() to determine if a kernel is compatible with a device. I think it would make sense for the implementation to return false if the kernel has no device image that could be used on the queried device.

It looks like there is a larger question here about whether the -fsycl-targets option should cause the runtime to hide devices for targets that were not compiled for. I think the answer has to be "no". Consider a case where the application consists of several source file, where some files were compiled with -fsycl-targets and others were not. I think it would be weird in a case like this to hide devices based on the -fsycl-targets option.

[AGindinson]

Consider a case where the application consists of several source file, where some files were compiled with -fsycl-targets and others were not.

-fsycl-targets is, in essence, used for every compilation - it's just that its implicit value is the generic spir64 target.

A larger concern with this example is the practical impossibility to compile any application in a way that would have some modules compiled for a different set of targets. Considering the following build scenario:

clang++ -fsycl module1.cpp -c-o module1.o
clang++ -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice module2.cpp -o module2.o

The object files would actually contain LLVM BC for the device code that would be yet to undergo the device linkage (sycl-post-link processing/SPIR-V Translator/AOT compilation if desired). To have kernels from module2 AOT-compiled for CPU device, the link command would have to look like:

clang++ -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice module1.o module2.o -o app.out

...which would result in both modules being AOT-compiled for CPU. I can hardly envision a compilation flow which would lead to the opposite (maybe some sort of an AOT-compiled shared library flow? please correct me if I'm missing such a scenario in https://github.com/intel/llvm/blob/sycl/sycl/doc/SharedLibraries.md).

There can be other reasons not to use compile-time information for exclusion of devices (far beyond my competence area), but I thought it best to elaborate on this particular example so that we don't delve into a misunderstanding of some sort.

[gmlueck]

The object files would actually contain LLVM BC for the device code that would be yet to undergo the device linkage (sycl-post-link processing/SPIR-V Translator/AOT compilation if desired). To have kernels from module2 AOT-compiled for CPU device, the link command would have to look like:

I guess you are saying that the -fsycl-targets option is ignored when creating an object file, and it is only meaningful on the link line? Answering your question about shared libraries, though, yes I think you could have a case where the main application and a shared library had different device image types (e.g. main library AOT compiled for CPU and shared lib has SPIR-V, etc.)

[AGindinson]

I guess you are saying that the -fsycl-targets option is ignored when creating an object file, and it is only meaningful on the link line?

In short, yes :)

[Ruyk]

I guess you are saying that the -fsycl-targets option is ignored when creating an object file, and it is only meaningful on the link line?

In short, yes :)

Probably going off-topic on the issue, but I am working on an example (codeplaysoftware/SYCL-For-CUDA-Examples#16) where I mix CUDA Fortran with SYCL code. Because nvfortran does its own magic at linking, I cannot use clang for linking the final binary. The only way I could put together the two is if I create a library with the SYCL code and link that with the Fortran binary.
@AGindinson is there any way to trigger the "sycl-post-link processing/SPIR-V Translator/AOT compilation" without creating a library or an executable?

To come back to the device selection topic: Would it be a valid use case to load a different dynamic library depending on the device you have selected? You would not know at device selection which device images are there or even which targets you have built for...

[AGindinson]

@Ruyk, I've reposted your question here: #3952 so that we can discuss that in isolation from the main thread.

[romanovvlad]

Based on the comments above it seems this is not feasible. Closing. Please feel free to reopen if any concerns.

[Michoumichmich]

Based on the comments above it seems this is not feasible. Closing. Please feel free to reopen if any concerns.

Yes, thanks! I forgot to close it