Obtaining Runtime Error Codes

Runtime APIs typically return an error code. The meaning of return codes and the method for obtaining errors vary by API type.

Asynchronous API: The API returns before the device task is complete. The return code only indicates whether task dispatch on the host is successful and cannot reflect the actual execution error on the device.
Synchronous API: The return code directly indicates the execution result. Typical scenarios include the following: An invalid parameter (such as a null pointer or an out-of-bounds write size) is passed, an unsupported hardware function is requested, or memory allocation fails.

Error Handling for Asynchronous APIs

To check for asynchronous errors immediately after an asynchronous function is called, the only way is to explicitly call a synchronous API (such as aclrtSynchronizeDevice) to block the host thread and check the error code returned by the synchronous API.

The Runtime module maintains an error variable for each host thread. This variable is initialized to ACL_RT_SUCCESS and is overwritten by the error code each time an error (whether it is a parameter verification error or an asynchronous error) occurs. The aclrtPeekAtLastError API returns only the value of this variable. The aclrtGetLastError API also returns the value of this variable, but resets it to ACL_RT_SUCCESS.

// Specify a device.
aclError error = aclrtSetDevice(0);

// Create a stream.
aclrtStream stream;
error = aclrtCreateStream(&stream);

// Set the ACL_STOP_ON_FAILURE mode.
error = aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);

// Dispatch the task to the stream. The return code only indicates whether the task is successfully dispatched. Generally, it indicates the parameter verification error on the host and cannot indicate the actual execution error on the device.
error = aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, stream);

// No return value is available after the kernel launch symbol <<<>>> is used. If only the host-side error needs to be detected, aclrtPeekAtLastError/aclrtGetLastError can be used to obtain the error.
myKernel<<<8, nullptr, stream>>>();

// Obtain the latest error of the current thread.
error = aclrtPeekAtLastError(ACL_RT_THREAD_LEVEL);
......

// Obtain the latest error of the current thread and reset the status to ACL_RT_SUCCESS.
error = aclrtGetLastError(ACL_RT_THREAD_LEVEL);
......

error = aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize, ACL_MEMCPY_DEVICE_TO_HOST, stream);

// [Key difference] The asynchronous API must call a stream synchronization API to obtain the asynchronous error on the device.
error = aclrtSynchronizeDevice();
if (error != ACL_RT_SUCCESS) {
    // Obtain ErrorMsg.
    char *errMsg = aclGetRecentErrMsg();
    // Output to logs.
    printf("Error: %s\n", errMsg);
    ......
}

// Destroy resources.
error = aclrtDestroyStream(stream);
error = aclrtResetDevice(0);

Note: In continuing upon failure mode (for details, see Task Stopping Upon Failure), if an exception occurs during the execution of a task in a stream, other tasks that have not been executed in the stream can still be executed, and new tasks can still be dispatched to the stream and other streams in the same context. In this case, error information returned by aclrtPeekAtLastError and aclrtGetLastError may not be about the first error.

Error Handling for Synchronous APIs

It is unnecessary to call the stream synchronization API. Instead, directly check the return value.

// Specify a device.
aclError error = aclrtSetDevice(0);

// [Key difference] You do not need to call APIs such as the stream synchronization API. Instead, directly verify the return value.
error = aclrtMemcpy(devPtr, devSize, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE);
if (error != ACL_RT_SUCCESS) {
    // Get ErrorMsg.
    char *errMsg = aclGetRecentErrMsg();
    // Output to logs.
    printf("Error: %s\n", errMsg);
    ......
}

// Destroy resources.
error = aclrtResetDevice(0);

Differences Between Asynchronous and Synchronous Error Handling

Item	Asynchronous API	Synchronous API
Return code meaning	Indicates only whether host-side dispatch was successful.	Directly indicates the execution result.
Whether stream synchronization is required	Required. Asynchronous errors are captured through synchronous APIs such as aclrtSynchronizeDevice.	Not required. The return code contains the execution result.
Error obtaining time	After stream synchronization.	When the API returns a result.

Recommendations

For any Runtime API call, it is necessary to check the return value to detect synchronous errors as early as possible. For asynchronous operations, errors must be checked again at key synchronization points (such as aclrtSynchronizeDevice) to ensure that errors during execution can be captured.

Parent topic: App Debugging