[SYCL] atomic_ref: Using wrong address space causes compiler crash.

[j-stephan]

Reproducer:

#include <CL/sycl.hpp>

#include <cstddef>
#include <cstdlib>

class my_kernel;

auto main() -> int
{
    using namespace cl::sycl;

    auto q = queue{gpu_selector{}};

    q.submit([](handler& cgh)
    {
        auto local_acc = accessor<std::byte, 0, access::mode::read_write, access::target::local>{cgh};

        cgh.parallel_for<my_kernel>(nd_range<1>{range<1>{2}, range<1>{42}}, [=](nd_item<1> my_item)
        {
            auto multi_ptr = local_acc.get_pointer();
            auto raw_ptr = reinterpret_cast<float*>(multi_ptr.get());

            auto ref = ONEAPI::atomic_ref<float, ONEAPI::memory_order_relaxed, ONEAPI::memory_scope::device,
                                          access::address_space::global_device_space>{*raw_ptr};
            ref.fetch_add(3.f);
        });
    });

    return EXIT_SUCCESS;
}

Obviously using global_device_space and/or memory_scope::device for the atomic_ref is wrong for a local accessor. However, this isn't diagnosed as such but instead causes a compiler crash:

$ clang++ -fsycl -sycl-std=2020 main.cpp
llvm-spirv: /home/stepha27/sycl-workspace/llvm/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp:915: SPIRV::SPIRVInstruction* SPIRV::LLVMToSPIRV::transUnaryInst(llvm::UnaryInstruction*, SPIRV::SPIRVBasicBlock*): Assertion `SrcAddrSpace == SPIRAS_Generic' failed.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace.
Stack dump:
0.	Program arguments: /home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv -o /tmp/main-c1e680-ce7ced.spv -spirv-max-version=1.1 -spirv-debug-info-version=legacy -spirv-allow-extra-diexpressions -spirv-ext=+all,-SPV_INTEL_usm_storage_classes /tmp/main-ff0fc8.bc
1.	Running pass 'LLVMToSPIRV' on module '/tmp/main-ff0fc8.bc'.
 #0 0x0000560366ea365d llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x64065d)
 #1 0x0000560366ea1294 llvm::sys::RunSignalHandlers() (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x63e294)
 #2 0x0000560366ea1400 SignalHandler(int) (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x63e400)
 #3 0x00007fdbfa15c0f0 __restore_rt (/usr/lib/libpthread.so.0+0x140f0)
 #4 0x00007fdbf9be1615 raise (/usr/lib/libc.so.6+0x3d615)
 #5 0x00007fdbf9bca862 abort (/usr/lib/libc.so.6+0x26862)
 #6 0x00007fdbf9bca747 _nl_load_domain.cold (/usr/lib/libc.so.6+0x26747)
 #7 0x00007fdbf9bd9bf6 (/usr/lib/libc.so.6+0x35bf6)
 #8 0x0000560366a013e6 SPIRV::LLVMToSPIRV::transUnaryInst(llvm::UnaryInstruction*, SPIRV::SPIRVBasicBlock*) (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x19e3e6)
 #9 0x00005603669fa8b3 SPIRV::LLVMToSPIRV::transValueWithoutDecoration(llvm::Value*, SPIRV::SPIRVBasicBlock*, bool, SPIRV::LLVMToSPIRV::FuncTransMode) (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x1978b3)
#10 0x00005603669fb725 SPIRV::LLVMToSPIRV::transValue(llvm::Value*, SPIRV::SPIRVBasicBlock*, bool, SPIRV::LLVMToSPIRV::FuncTransMode) (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x198725)
#11 0x0000560366a02ab6 SPIRV::LLVMToSPIRV::transFunction(llvm::Function*) (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x19fab6)
#12 0x0000560366a05f40 SPIRV::LLVMToSPIRV::translate() (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x1a2f40)
#13 0x0000560366a06143 SPIRV::LLVMToSPIRV::runOnModule(llvm::Module&) (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x1a3143)
#14 0x0000560366dc8f63 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x565f63)
#15 0x0000560366a067aa llvm::writeSpirv(llvm::Module*, SPIRV::TranslatorOpts const&, std::ostream&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&) (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0x1a37aa)
#16 0x0000560366927ddd main (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0xc4ddd)
#17 0x00007fdbf9bcc152 __libc_start_main (/usr/lib/libc.so.6+0x28152)
#18 0x0000560366938b2e _start (/home/stepha27/sycl-workspace/llvm/build/install/bin/llvm-spirv+0xd5b2e)
llvm-foreach: Aborted (core dumped)
clang-12: error: llvm-spirv command failed with exit code 1 (use -v to see invocation)

Output of clang --version:

clang version 12.0.0 (https://github.com/intel/llvm 5c4df90502995fe0ffb9cd96849302dc1d3d862a)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/stepha27/sycl-workspace/llvm/build/install/bin

[MrSidims]

Thanks for the report, I'll take a look on it.

[bader]

This is UB, which compiler won't be able to diagnose reliably. If you build the compiler w/o assertions, I think it won't crash like this, but I can't predict the outcome in this case.

[MrSidims]

This is UB, which compiler won't be able to diagnose reliably. If you build the compiler w/o assertions, I think it won't crash like this, but I can't predict the outcome in this case.

I think, that we can make some sort of diagnostic in LLVM to SPIR-V translator, in a case when we have invalid casts from/to generic address space.

[bader]

This is UB, which compiler won't be able to diagnose reliably. If you build the compiler w/o assertions, I think it won't crash like this, but I can't predict the outcome in this case.

I think, that we can make some sort of diagnostic in LLVM to SPIR-V translator, in a case when we have invalid casts from/to generic address space.

I think it's already clear from the existing error message:

Assertion `SrcAddrSpace == SPIRAS_Generic' failed

Do have something different in mind?

[MrSidims]

Do have something different in mind?

But it comes from an assert, hence we won't get it in release build.
Something like that does look a little bit more user-friendly:

 ./bin/clang++ -fsycl test2.cpp
InvalidModule: Invalid SPIR-V module: Casts from private to global are disallowed


llvm-foreach:
clang-11: error: llvm-spirv command failed with exit code 1 (use -v to see invocation)

Error message can be adjusted.

[j-stephan]

I think it's already clear from the existing error message:

For what it's worth, it wasn't directly clear to me as user. Partly because the original bug occurred in a more convoluted code where it wasn't immediately obvious that this is UB, partly because I've had compilers crash on perfectly valid code before. I don't think that crashing the compiler is a good way to signal the user that there is something wrong with the code.

Couldn't this be handled more gracefully with static_assert and/or SFINAE inside the SYCL headers?

[bader]

Couldn't this be handled more gracefully with static_assert and/or SFINAE inside the SYCL headers?

Unfortunately no. As I said before, compiler is unable to detect this type of issues at compile time.
If you pass raw_ptr to atomic_ref constructor inside if-statement which condition is unknown at compile time, compiler can't claim it's an error.

What happens here: the optimizer is able to reduce the problematic code to a single instruction, which is considered invalid for GPU execution environment. The tool crashed on assertion is able to detect such instructions, but nothing more.

If you compile your example with disabled optimizations -O0, this example will work (at least on CPU device).

[MrSidims]

Assertions in the translator were replaced with error messages in KhronosGroup/SPIRV-LLVM-Translator#842 and this patch was pulled in here. Hope, that following error message helps a bit:

InvalidModule: Invalid SPIR-V module: Casts from private/local/global address space are allowed only to generic
  %1 = addrspacecast i32 addrspace(3)* %0 to i32 addrspace(1)*
llvm-foreach:
clang-11: error: llvm-spirv command failed with exit code 1 (use -v to see invocation)

[j-stephan]

Yes, that is much better and clearer. Thank you!