I’m trying to determine how to figure out if a CUDA error is stiky or not, specifically since I need to know if the host process should be terminated or not.
So, I made some tests to figure out how sticky errors can be investigated, and found the following:
// this kernel will access 'input' or 'output' out of bounds, causing a 'cudaErrorIllegalAddress' (code 700)
a_kernel_with_error<<<1,1>>>(input, n_input, output, n_output);
// call 'cudaGetLastError' to verify whether the kernel launch failed
cudaError_t first_get(::cudaGetLastError());
// 'first_get' is 'cudaSuccess', the "out of bounds" error is asynchronous
// run some host code ...
// get to a synchronization point, e.g. copy 'output' to cpu
std::vector<int> output_copy(n_output);
cudaError_t sync_point(::cudaMemcpy(output_copy.data(), output, count * sizeof(int), cudaMemcpyDeviceToHost));
// 'sync_point' stores 'cudaErrorIllegalAddress', since now cpu and gpu synchronize
// now, there is a CUDA error, and I need to determine whether to kill the host process or not
// so, I start to inquiry the gpu with respect to errors:
cudaError_t new_get(::cudaGetLastError());
// 'new_get' stores 'cudaErrorIllegalAddress', as expected since the gpu got a sticky error,
// but if I call again 'cudaGetLastError' the error has been reset to 'cudaSuccess'
cudaError_t a_second_get(::cudaGetLastError()); // cudaSuccess
// every other call to 'cudaPeekAtLastError' or 'cudaGetLastError' returns now 'cudaSuccess'
// only if I synchronize again, I get 'cudaErrorIllegalAddress', but it can be reset with 'cudaGetLastError'
cudaError_t device_sync(::cudaDeviceSynchronize()); // cudaErrorIllegalAddress
cudaError_t another_get(::cudaGetLastError()); // cudaErrorIllegalAddress
cudaError_t last_get(::cudaGetLastError()); // cudaSuccess
So, it looks like that to determine if an error is sticky, the following lines of code should suffice:
cudaError_t first_sync(::cudaDeviceSynchronize());
cudaError_t try_to_reset(::cudaGetLastError());
cudaError_t second_sync(::cudaDeviceSynchronize());
if (first_sync == second_sync)
// the error is sticky
Can anyone confirm I can go with syncronization-reset-syncronization in order to determine if an error is sticky?
As a side note, tbh I don’t understand why ‘cudaGetLastError’ resets the error even with sticky errors, this creates some confusion.