External allocations fail under high memory pressure

[denizyuret]

This is a memory stability issue I was not able to crack during training of a largish model written by @mg9. To replicate, download test.zip, unzip and run julia in the resulting directory and include("train.jl"). The errors appear within 1-2 minutes. There are two problems as far as I can see:

The first problem is something we have known for a while (see e.g. denizyuret/Knet.jl#556): CUDA.jl grabs almost all GPU memory, CUDNN ends up complaining when it tries to allocate off-pool. At least I think this is what's going on because (1) it typically happens when CUDA.usage[] is close to GPU memory limit during expensive CUDNN calls, (2) it goes away if we use CUDA.usage_limit[]. Interestingly it results in an EXECUTION_FAILED error rather than out of memory:

julia> include("train.jl")
...
iteration 29

Stacktrace:
 [1] throw_api_error(::CUDA.CUDNN.cudnnStatus_t) at /userfiles/dyuret/.julia/dev/CUDA/lib/cudnn/error.jl:19
 [2] macro expansion at /userfiles/dyuret/.julia/dev/CUDA/lib/cudnn/error.jl:30 [inlined]
 [3] cudnnRNNForwardTraining(::Ptr{Nothing}, ::Knet.RD, ::Int64, ::Knet.TDs, ::KnetArray{Float32,3}, ::Ptr{Nothing}, ::CUDA.CuPtr{Nothing}, ::Ptr{Nothing}, ::CUDA.CuPtr{Nothing}, ::Knet.FD, ::KnetArray{Float32,3}, ::Knet.TDs, ::KnetArray{Float32,3}, ::Knet.TD, ::KnetArray{Float32,3}, ::Knet.TD, ::KnetArray{Float32,3}, ::KnetArray{UInt8,1}, ::Int64, ::KnetArray{UInt8,1}, ::Int64) at /userfiles/dyuret/.julia/dev/CUDA/lib/utils/call.jl:93
...
ERROR: LoadError: CUDNNError: CUDNN_STATUS_EXECUTION_FAILED (code 8)

For this problem if my hypothesis is correct we just need to figure out what is trying to grab memory and whether there is a way to give it memory from the pool. If not we just need to have usage_limit[] set below some safety limit at initialization.

Starting with a usage_limit[] solves the first problem and makes the second problem appear, which I could not make any progress on:

julia> using CUDA; CUDA.usage_limit[] = 10^10  # For a 12GB K80 card
julia> include("train.jl")
...
iteration 68

Stacktrace:
 [1] throw_api_error(::CUDA.cudaError_enum) at /userfiles/dyuret/.julia/dev/CUDA/lib/cudadrv/error.jl:103
 [2] macro expansion at /userfiles/dyuret/.julia/dev/CUDA/lib/cudadrv/error.jl:110 [inlined]
 [3] cuMemcpyDtoH_v2(::Ptr{Float32}, ::CuPtr{Nothing}, ::Int64) at /userfiles/dyuret/.julia/dev/CUDA/lib/utils/call.jl:93
 [4] _unsafe_copy!(::Array{Float32,2}, ::Int64, ::KnetArray{Float32,2}, ::Int64, ::Int64) at /dev/shm/dyuret/.julia/packages/Knet/exwCE/src/karray.jl:359
 [5] convert at /dev/shm/dyuret/.julia/packages/Knet/exwCE/src/karray.jl:120 [inlined]
 [6] convert(::Type{Array{Int32,N} where N}, ::KnetArray{Float32,2}) at /dev/shm/dyuret/.julia/packages/Knet/exwCE/src/karray.jl:119
 [7] (::DeepBiaffineGraphDecoder)(::AutoGrad.Result{KnetArray{Float32,3}}, ::Array{Int64,2}, ::Array{Int64,2}, ::Array{Int64,2}) at /scratch/users/dyuret/mrp/test/deep_biaffine_graph_decoder.jl:92
...
ERROR: LoadError: CUDA error: an illegal memory access was encountered (code 700, ERROR_ILLEGAL_ADDRESS)

Unfortunately this error does not always occur in the same iteration or on the same function as is typical with memory errors. So I tried running julia with cuda-memcheck:

$ cuda-memcheck julia
julia> using CUDA; CUDA.usage_limit[] = 10^10  # For a 12GB K80 card
julia> include("train.jl")
...
iteration 2
========= Invalid __global__ read of size 4
...
Stacktrace:
 [1] throw_api_error(::CUDA.cudaError_enum) at /userfiles/dyuret/.julia/dev/CUDA/lib/cudadrv/error.jl:103
 [2] macro expansion at /userfiles/dyuret/.julia/dev/CUDA/lib/cudadrv/error.jl:110 [inlined]
 [3] cuMemAlloc_v2(::Base.RefValue{CuPtr{Nothing}}, ::Int64) at /userfiles/dyuret/.julia/dev/CUDA/lib/utils/call.jl:93
 [4] alloc at /userfiles/dyuret/.julia/dev/CUDA/lib/cudadrv/memory.jl:84 [inlined]
...
ERROR: LoadError: CUDA error: unspecified launch failure (code 719, ERROR_LAUNCH_FAILED)

This one I get fairly consistently with cuda-memcheck. Any ideas?

[maleadt]

For this problem if my hypothesis is correct we just need to figure out what is trying to grab memory and whether there is a way to give it memory from the pool.

That's the ideal solution. If that doesn't work, you can always retry the API call while reclaiming memory (this is done already when the library properly reports ALLOC_FAILED):

CUDA.jl/lib/cusolver/wrappers.jl

Lines 25 to 35 in 77bc62b

	function cusolverSpCreate()
	handle_ref = Ref{cusolverSpHandle_t}()
	res = @retry_reclaim err->isequal(err, CUSOLVER_STATUS_ALLOC_FAILED) ||
	isequal(err, CUSOLVER_STATUS_NOT_INITIALIZED) begin
	unsafe_cusolverSpCreate(handle_ref)
	end
	if res != CUSOLVER_STATUS_SUCCESS
	throw_api_error(res)
	end
	return handle_ref[]
	end

For the other issue, it's possible memory in the pool gets corrupted somehow. Running under cuda-memcheck might reveal something.

[denizyuret]

The last box above was under cuda-memcheck, I just don't understand why a regular alloc would cause a Invalid __global__ read of size 4.

[maleadt]

Ah, I can't read. What's the device frame it reported for that invalid read, within memcpy or asynchronously from a kernel of ours?

[denizyuret]

I could not get the same error today after updating CUDA.jl from master. I went back to the last release for replicability. Here is the error log: https://gist.github.com/denizyuret/3082f805381126b2f0d0d8825f0bf833

Next I will try to replicate the error without using KnetArrays to rule out any weird memory block alignment etc type issues while I am hopping back and forth between CUDA pointers and Knet pointers. For that I need to get RNNs working with CuArrays, which I was planning to do anyway.

In the meantime if you can generate some hypotheses looking at the log I'd be happy to test them...

[maleadt]

I just don't understand why a regular alloc would cause a Invalid __global__ read of size 4.

Looking at the full log now, this seems like a legitimate out-of-bounds access in a Knet kernel (in _addrows_32(int, int, int, int*, float*, float*)); it just gets shown under the call to alloc because that's the first synchronizing API call.

[denizyuret]

It turns out there was a getindex without a bounds check.

[denizyuret]

If that doesn't work, you can always retry the API call while reclaiming memory (this is done already when the library properly reports ALLOC_FAILED):

CUDA.jl/lib/cusolver/wrappers.jl

Lines 25 to 35 in 77bc62b

	function cusolverSpCreate()
	handle_ref = Ref{cusolverSpHandle_t}()
	res = @retry_reclaim err->isequal(err, CUSOLVER_STATUS_ALLOC_FAILED) ||
	isequal(err, CUSOLVER_STATUS_NOT_INITIALIZED) begin
	unsafe_cusolverSpCreate(handle_ref)
	end
	if res != CUSOLVER_STATUS_SUCCESS
	throw_api_error(res)
	end
	return handle_ref[]
	end

This looks great. Why can't we make this the standard behavior for all our calls?

[maleadt]

Why can't we make this the standard behavior for all our calls?

It is like that:

CUDA.jl/lib/cusolver/error.jl

Lines 50 to 59 in 47bc733

	macro check(ex)
	quote
	res = @retry_reclaim isequal(CUSOLVER_STATUS_ALLOC_FAILED) $(esc(ex))
	if res != CUSOLVER_STATUS_SUCCESS
	throw_api_error(res)
	end
	return
	end
	end

Here, the error was misreported.

[denizyuret]

I think the problem is (at least with CUDNN) that the error reported is not always ALLOC_FAILED. This could be fixed if instead of insisting on ALLOC_FAILED, we tried this whenever the returned status is not STATUS_SUCCESS. (e.g. the error I mostly get is CUDNN_STATUS_EXECUTION_FAILED, but I am not sure how consistent this is. At the very least we should add this to the CUDNN @check macro).

[maleadt]

This could be fixed if instead of insisting on ALLOC_FAILED, we tried this whenever the returned status is not STATUS_SUCCESS.

No, that's way too hacky.