Throw error from device to host with a msg #2258

[mehmetyusufoglu]

Macros to make signalling an error from the device side to the host code.

In case of Cuda, Hip, Sycl a user defined message is added to the abort. For other backends std::runtime_error exception with the user message is thrown.

Testing:
-Tests could be done for Accs: CpuThreads and CpuBlocks by catching the runtime_error exceptions thrown during exec.
-Aborts can not be catched from Cuda, Hip, Sycl as we call exec. (Only tested by running a temporary fail test at CI)
But for Cuda; __trap triggers runtime_error can be catched during the wait(queue).
-For sycl assert(false) is used. std::abort generated a runtime error at HAL no compile error but not "aborted"
-I turned out to be OpenMP specification mandates std::runtime errors should be handled by the same thread otherwise it is converted to abort. I checked with a signal handler and SIGABRT is fired. (Therefore openmp cases could not be tested other than a one time fail test at CI.)

Issue

fix #2258

[fwyzard]

Are you sure we can call std::abort() from inside a SYCL kernel running e.g. on a GPU ?

[mehmetyusufoglu]

We only support one-api and it is tested on CI and HAL. . I wanted to do assert(false) but assert turned out to be calling std::abort from kernel in sycl (url)

[mehmetyusufoglu]

I've tried looking into the specification of SYCL and the implementation of oneAPI, and I suspect the only portable solution may be assert(false).

Hover

* on my laptop integrated GPU that actually does not do anything :-(

* on a Flax 170 datacenter GPU it causes the program to abort in a way that cannot be caught.

I'll bring it up with Intel at the next occasion - but for the time being, that seems better than nothing.

Ok I have added assert(false) , thanks a lot. I can not catch the exception at HAL computer so i did not test it in the test code.

1
2
AssertHandler::printMessage
test_abort.cpp:27: auto main()::(anonymous class)::operator()(sycl::handler &)::(anonymous class)::operator()() const: global id: [0,0,0], local id: [0,0,0] Assertion `false` failed.

[fwyzard]

can not catch the exception at HAL computer so i did not test it in the test code.

Unrelated to the test, but what is the HAL computer ?

[fwyzard]

Does this need to be a macro ?
Can we use a device function instead ?

[mehmetyusufoglu]

Will not make much difference since only 2 lines are inlined. On the other hand Isn't it better to call the aborts or exceptions directly ?

[fwyzard]

@psychocoderHPC what do you prefer here, functions or macros ?

[psychocoderHPC]

Functions are always prefered above macros. We should only use macros if we need __LINE__, ...

[fwyzard]

I think a message like

%s [ALPAKA_DEVICE_THROW: Calling __trap(). Cuda backend is enabled.]\n

is fine for debugging the implementation - but once this goes in production can we print just the user-supplied message ?

Or, if we do want to add more information, something like

alpaka encountered a user-defined error condition while running on the ___ back-end:\n
%s

[mehmetyusufoglu]

ok, changed. Thanks.

[mehmetyusufoglu]

can not catch the exception at HAL computer so i did not test it in the test code.

Unrelated to the test, but what is the HAL computer ?

A HPC system at HZDR :) Has different kind of gpus and easily configurable.

[psychocoderHPC]

I was on the way to say change it to a C++ function and than this should be deviceThrow() within the alpaka namespace.
BUT than it should be a trait which gets the acc as argument with the coresponding class equal to the trait name which is than specialized for each acc. Device functions always take the acc as first argument.

Never the less I would handle it equal to ALPAKA_ASSERT_ACC and would add the line and error message to the printf.
To follow the naming schema this macro should be named ALPAKA_THROW_ACC

[mehmetyusufoglu]

Hence; according to you it should be a macro called ALPAKA_THROW_ACC without an Acc argument?

[psychocoderHPC]

yes this will follow our macro ALPAKA_ASSERT_ACC and avoid introducing a new trait. IMO this function/macro will be used in seldem cases therefore I would like to avoid increasing the code base because of this tiny throw.

[mehmetyusufoglu]

Ok, then now we have a macro which is taking a string, throwing or aborting depending on backend.

[mehmetyusufoglu]

macro name changed to ALPAKA_THROW_ACC. Thanks.

[psychocoderHPC]

@mehmetyusufoglu please rebase against develop branch to fix the CI issues