gdb: add "catch hiperr" to break on HIP API errors

[amd-shahab]

This change relies on "Add __hipOnError debugger hook for HIP API failures" PR to be functional. Although, without it, ROCgdb won't be broken. Please read the commit message for the description of this change itself.

--

What

This change adds a new feature to GDB, when it is targeted for
amdgcn-amd-amdhsa, that enables it to stop whenever HIP API
errors are returned. An abridged sample of execution looks
like this:

(gdb) catch hiperr
(gdb) run
Thread 1 "hip" hit Catchpoint 1 (hip error)
hipErrorInvalidDevice (101): invalid device ordinal

Moreover, a $_hiperr convenience variable is added that holds
the error code. For further details, please see the updated
documentation in this change.

Design

Most of the implementation is in the newly added module called
break-catch-hiperr.c. The naming and its context follows the
break-catch-throw.c module.

The whole mechanism work by setting a breakpoint (catchpoint)
on a __hipOnError() symbol, that must be provided by the CLR
library. Enabling the catchpoint when there's no __hipOnError
symbol, simply has no effect.

There's a architect dependent section that is responsible for
extracting the error parameters (code, name, description) from
a single argument passed to the __hipOnError() function as a
structure pointer. The expected format of this struct is:

uint32_t version; // at this point, only 1 is supported
uint32_t code; // the numerical error code
const char *name; // name of the error
const char *desc; // a phrase about the error

The code is devised in such a way that if the __hipOnError()
exists, but for whatever reason the parameters cannot be deduced,
it won't interfere with the catchpoint being triggered. It will
affect the reported data of the catchpoint though.

Test

A new test, gdb.rocm/hip-catch-errors, is added to the mix to
ensure that:

1. (Pending) catchpoints work.
2. Multiple catchpoints in one execution are reported accordingly.
3. Errors of indirect HIP calls like "kenel<<<...>>>()" are caught.
4. The convenience variable can be used for conditional situations.
5. Conditional and temporary catchpoints work as expected.

ROCM-1548

[lumachad]

@amd-shahab If this shouldn't be tested yet and isn't mergeable until the dependent PR, feel free to label it "ci:skip".

[lancesix]

First quick first pass of comments.

One question regarding the test: what kind of guarantee do we have that the error code and error string are going to be stable over time? Let's imagine that the hip runtime decides to localize the error string (not the name). Couldn't that be a legitimate thing to do, but it would break the test, would it?

[amd-shahab]

• error string not part of API

• __hipOnError -> @samp{__hipOnError}

• "int main" -> "int\nmain"

• gdb_test -> gdb_test_on_output

• target_read_stack -> target_read_memory

• drop the "ignored" param name

• Else -> Otherwise

• NULL -> nullptr

• removed dangling "backtrace" command from the documentation

• !arg -> arg != nullptr

• binay -> binary

• relvant -> relevant

• bool print_one (const bp_location **) -> bool print_one (const bp_location **last_loc)

• rename variable "x" to "dontcare"

• comment on true/false return of amd64_fetch_hiperr_parameters()

[amd-shahab]

• emit a warning when the "version" is not expected.

[amd-shahab]

• make the invocation of the test fully argument dependent
• add comments for every method in break-catch-hiperr.c

[amd-shahab]

• check for the existence of __hipOnError symbol. If not found, skip the test.

[amd-shahab]

• check for the right frame (__hipOnError)

[amd-shahab]

• Use separate hooks for Linux and Windows implementation of amd64 targets
• The test comes up with its own reference value instead of having them hardcoded.
• Better string output. Also, tell user about $_hiperr and backtrace.

[aktemur]

If I set the feature with "catch hiperr" and then do an infcall with "print hipSetDevice(98765)", what's GDB's behavior? Can you please check?

[lancesix]

If I set the feature with "catch hiperr" and then do an infcall with "print hipSetDevice(98765)", what's GDB's behavior? Can you please check?

I have not tested, but am highly confident that this will trigger the catch point.

[aktemur]

If I set the feature with "catch hiperr" and then do an infcall with "print hipSetDevice(98765)", what's GDB's behavior? Can you please check?

I have not tested, but am highly confident that this will trigger the catch point.

I also think the catchpoint will be triggered but just in case I wanted make sure GDB can handle the infcall interruption correctly. It would be great to test with and without set unwind-on-signal on. (I cannot test because I don't have the hipOnError symbol.

[amd-shahab]

If I set the feature with "catch hiperr" and then do an infcall with "print hipSetDevice(98765)", what's GDB's behavior? Can you please check?

As Lancelot predicted:

(gdb) catch hiperr
Catchpoint 2 (HIP error)
(gdb) p hipSetDevice(98765)
[New Thread 0x7ffff4fff6c0 (LWP 8965) (id 2)]
[New Thread 0x7ffff47fe6c0 (LWP 8966) (id 3)]
[New Thread 0x7ffff3cff6c0 (LWP 8967) (id 4)]
[Thread 0x7ffff3cff6c0 (LWP 8967) (id 4) exited]

Thread 1 "hip-catch-error" hit Catchpoint 2 (HIP error)
HIP API call failed with error hipErrorInvalidDevice (101): invalid device ordinal

The $_hiperr holds the error number.
A "backtrace" will show the call stack.
The program being debugged stopped while in a function called from GDB.
Evaluation of the expression containing the function
(hipSetDevice(int)) will be abandoned.
When the function is done executing, GDB will silently stop.


(gdb) bt
#0  __hipOnError (err_info=0x7fffffffd7a0) at /src/clr/hipamd/src/hip_error.cpp:11
#1  0x00007ffff62863e9 in hip::hipSetDevice (device=98765) at /src/clr/hipamd/src/hip_device_runtime.cpp:784
#2  0x00007ffff668440f in hipSetDevice (deviceId=98765) at /src/clr/hipamd/src/hip_table_interface.cpp:2133
#3  <function called from gdb>
#4  main (argc=1, argv=0x7fffffffe108) at /src/rocgdb/gdb/testsuite/gdb.rocm/hip-catch-errors.cpp:26

With set unwind-on-signal on (default is off), we get the same output. "Signal"s have nothing to do with these errors.

[amd-shahab]

• In catchpoint output, refer to "$_hiperr" as a "convenience variable"
• In catchpoint output, change the "backtrace" section to the suggested sentence
• gdb.texinfo updated to reflect the latest output.
• Use "$::hip_" instead of "global hip_"
• Use ATTRIBUTE_UNUSED
• rename the fetch_hiperr parameter names to err_{no,name,str}
• a few typos are fixed

[amd-shahab]

• test: Add an implicit and faulty kernel<<<1, big_size>>> call to HIP API as test case
• test: Make the routine for getting reference values and causing the errors the same