Flaky CI failure on SYCL :: Assert/assert_in_simultaneous_kernels.cpp on CUDA

[aelovikov-intel]

See https://github.com/intel/llvm/runs/7240670323?check_suite_focus=true. The relevant part of the log is

$ ":" "RUN: at line 13"
note: command had no output on stdout or stderr
$ "env" "SYCL_PI_LEVEL_ZERO_TRACK_INDIRECT_ACCESS_MEMORY=1" "env" "SYCL_DEVICE_FILTER=ext_oneapi_cuda:gpu,host" "SYCL_PI_CUDA_ENABLE_IMAGE_SUPPORT=1" "/__w/llvm/llvm/build/SYCL/Assert/Output/assert_in_simultaneous_kernels.cpp.tmp.out"
# redirected output from '/__w/llvm/llvm/build/SYCL/Assert/Output/assert_in_simultaneous_kernels.cpp.tmp.txt':
/__w/llvm/llvm/llvm_test_suite/SYCL/Assert/assert_in_simultaneous_kernels.hpp:13: void assertFunc(): block: [9
PI CUDA ERROR:,7
	Value:           ,0710], thread: [0
	Name:            ,0CUDA_ERROR_ASSERT,0
	Description:     ] Assertion `false && "from assert statement"device-side assert triggered
	Function:        record
	Source Location: /__w/llvm/llvm/src/sycl/plugins/cuda/pi_cuda.cpp:570` failed.
terminate called after throwing an instance of 'cl::sycl::runtime_error'
  what():  Native API failed. Native API returns: -999 (Unknown PI error) -999 (Unknown PI error)
note: command had no output on stdout or stderr
error: command failed with exit status: -6
$ "true"

later

$ "env" "SYCL_DEVICE_FILTER=ext_oneapi_cuda:gpu,host" "SYCL_PI_CUDA_ENABLE_IMAGE_SUPPORT=1" "FileCheck" "/__w/llvm/llvm/llvm_test_suite/SYCL/Assert/assert_in_simultaneous_kernels.cpp" "--input-file" "/__w/llvm/llvm/build/SYCL/Assert/Output/assert_in_simultaneous_kernels.cpp.tmp.txt"
# command stderr:
/__w/llvm/llvm/llvm_test_suite/SYCL/Assert/assert_in_simultaneous_kernels.cpp:19:11: error: CHECK: expected string not found in input
// CHECK: {{.*}}assert_in_simultaneous_kernels.hpp:13: void assertFunc(): {{.*}}[9,7,0], {{.*}}[0,0,0]
          ^
/__w/llvm/llvm/build/SYCL/Assert/Output/assert_in_simultaneous_kernels.cpp.tmp.txt:1:1: note: scanning from here
/__w/llvm/llvm/llvm_test_suite/SYCL/Assert/assert_in_simultaneous_kernels.hpp:13: void assertFunc(): block: [9
^
/__w/llvm/llvm/build/SYCL/Assert/Output/assert_in_simultaneous_kernels.cpp.tmp.txt:1:40: note: possible intended match here
/__w/llvm/llvm/llvm_test_suite/SYCL/Assert/assert_in_simultaneous_kernels.hpp:13: void assertFunc(): block: [9
                                       ^
Input file: /__w/llvm/llvm/build/SYCL/Assert/Output/assert_in_simultaneous_kernels.cpp.tmp.txt
Check file: /__w/llvm/llvm/llvm_test_suite/SYCL/Assert/assert_in_simultaneous_kernels.cpp
-dump-input=help explains the following input dump.
Input was:
<<<<<<
            1: /__w/llvm/llvm/llvm_test_suite/SYCL/Assert/assert_in_simultaneous_kernels.hpp:13: void assertFunc(): block: [9 
check:19'0     X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found
check:19'1                                            ?                                                                        possible intended match
            2: PI CUDA ERROR:,7 
check:19'0     ~~~~~~~~~~~~~~~~~
            3:  Value: ,0710], thread: [0 
check:19'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~
            4:  Name: ,0CUDA_ERROR_ASSERT,0 
check:19'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            5:  Description: ] Assertion `false && "from assert statement"device-side assert triggered 
check:19'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            6:  Function: record 
check:19'0     ~~~~~~~~~~~~~~~~~~
            .
            .
            .
>>>>>>
error: command failed with exit status: 1

It looks like some line-buffering results in output from different places intermixed.

[pvchupin]

Issue faced in CUDA testing but I'm not quite sure if it's CUDA specific.
I think @steffenlarsen also fixed some races in LIT tests... could be related?

[steffenlarsen]

Races were for unittests only. These tests were recently enabled for CUDA.

[jchlanda]

Looks like the output from pi's check_error gets garbled with CUDA's __assert_fail message. Even if check_error was to be guarded by a mutex, I'm not sure we could do anything about the output from CUDA.

[pvchupin]

@jchlanda, does it mean that existing tests are not good to test asserts on CUDA?
Is it possible to modify existing tests?
Does it makes sense to develop a new tests for CUDA and disable current ones?

[bader]

@jchlanda, can we build the string we want to print and use single << operator to print it?
https://github.com/intel/llvm/blob/sycl/sycl/plugins/cuda/pi_cuda.cpp#L140-L146
I hope it will separate the assert and PI error messages.

[jchlanda]

@pvchupin I think the test is well put together, it creates a thread pool that submits multiple asserting kernels, creating a situation in which multiple kernels can asserts simultaneously.

@bader I'm not sure if bundling the strings would fix it entirely, it would improve the situation, as in less "garbling points", but ultimately it would not solve the problem, as both __assert_fail (as provided by CUDA toolchain) and check_error from oneAPI's plugin implementation compete on the same resource. Given the device side nature of __assert_fail I don't think we could do anything to help it without having some sort of kernel level synchronisiation primitives, which I don't think would be worth it. I think calls to std::out should be thread safe, but you are right that multiple uses of << would be treated as separate calls.
Looking at one of incorrect printouts:

	Value:           ,0710], thread: [0

it shows:

Value: ,0 710 ], thread: [0

where it got garbled, along the << operator:

     << "\n\tValue:           " << result

I'm tempted to bundle up the error message, like you suggest, so there is only one point where it can go wrong, and mark it as unsupported for CUDA.

[bader]

As long as the string we match is emitted by a single << operator and it's unique (i.e. we can't match wrong message from __assert_fail or check_error), we should be fine.

[jchlanda]

I don't think it's the case, even with bundled message it would still be possible to have the output of check_error in the middle of __assert_fail's printout. And the test matches on the __assert_fail message.

[bader]

I don't think it's the case, even with bundled message it would still be possible to have the output of check_error in the middle of __assert_fail's printout. And the test matches on the __assert_fail message.

Could you point the part of the log with an example of such case, if they are in the logs, please? I think the case from this comment demonstrates garbling between the << operator calls.

[jchlanda]

If you look at the begging of the output log, it says:

/__w/llvm/llvm/llvm_test_suite/SYCL/Assert/assert_in_simultaneous_kernels.hpp:13: void assertFunc(): block: [9
**********__assert_fail*****************************************************************************************
PI CUDA ERROR:,7
^^^^^^^^^^^^^^** - assert_fail
| check_error

So if we bundle up, the difference will be that instead of partial error message (PI CUDA ERROR:) we'd get the full, but there is no guarantee that __assert_fail's message would not get interleaved.

[bader]

I think I get the problem now. Test matches the string that is not explicitly printed neither by the test nor by the DPC++ runtime. We are matching the driver output, which we do not control. Not sure whether it's a good idea, but in this case, I would suppress DPC++ runtime outputs to not garble the driver output.

I guess this issue doesn't impact other backends because they use "fallback" implementation, which emulates low-level driver behavior via DPC++ device library, which output is synchronized with DPC++ runtime output.