New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
portable device-side abort()
or throw
?
#2258
Comments
I talked with @psychocoderHPC . I did a mistake. I was not alpaka, it was PIC which I saw it and it use native traps. |
IMO this is interesting. A cool feature would be could handle such throws to avoid an application crash but as I know this is not possible. @fwyzard For HIP we used the Clang intrinsic |
With CUDA 12.3 and __global__ void test() {
__trap();
} compiles to
while __global__ void test() {
assert(0);
} compiles to
CUDA does not have With ROCm 5.7 and __global__ void test() {
__trap();
} compiles to test1(): ; @test1()
s_trap 2 while __global__ void test() {
assert(0);
} compiles to test2(): ; @test2()
s_add_u32 flat_scratch_lo, s6, s9
s_addc_u32 flat_scratch_hi, s7, 0
s_add_u32 s0, s0, s9
s_addc_u32 s1, s1, 0
s_mov_b64 s[8:9], s[4:5]
s_mov_b32 s32, 0
s_getpc_b64 s[6:7]
s_add_u32 s6, s6, __assert_fail@rel32@lo+4
s_addc_u32 s7, s7, __assert_fail@rel32@hi+12
s_swappc_b64 s[30:31], s[6:7]``` plus a very long definition for |
In a few places in the CMS code we would like to signal an error from the device side to the host code.
Something similar to a device-side
assert(false)
, but with a more meaningful error message and always available (ALPAKA_ASSERT_ACC
may not be implemented on some back-ends).For example, in one case we have something like
Does Alpaka already provide something like this ?
If not, would it be interesting to have it in Alpaka ?
The text was updated successfully, but these errors were encountered: