compute-sanitizer out-of-bounds failure

[maleadt]

We're seeing some compute-sanitizer failures in #772, which I can reproduce locally with the following snippet:

using CUDA

a = CuArray{Float32}(undef, 1)
@show pointer(a)
b = CuArray{Float32}(undef, 1)
@show pointer(b)
c = CuArray{Float64}(undef, 1)
@show pointer(c)

a .^ b
device_synchronize()
a .^ b
device_synchronize()

a .^ c
device_synchronize()

Full API trace.

cuDeviceGet(Base.RefValue{Int32}, 0)cuDriverGetVersion(Base.RefValue{Int32}) = CUDA_SUCCESS
 1: 11020
cuDeviceGetCount(Base.RefValue{Int32}) = CUDA_SUCCESS
 1: 1
cuDeviceGet(Base.RefValue{Int32}, 0) = CUDA_SUCCESS
 1: 0
cuDeviceGet(Base.RefValue{Int32}, 0) = CUDA_SUCCESS
 1: 0
cuDeviceGet(Base.RefValue{Int32}, 0) = CUDA_SUCCESS
 1: 0
cuDeviceGet(Base.RefValue{Int32}, 0) = CUDA_SUCCESS
 1: 0
cuDeviceGet(Base.RefValue{Int32}, 0) = CUDA_SUCCESS
 1: 0
cuDeviceGet(Base.RefValue{Int32}, 0) = CUDA_SUCCESS
 1: 0
cuDeviceGet(Base.RefValue{Int32}, 0) = CUDA_SUCCESS
 1: 0
cuDeviceGetAttribute(Base.RefValue{Int32}, CU_DEVICE_ATTRIBUTE_MEMORY_POOLS_SUPPORTED, CuDevice(0)) = CUDA_SUCCESS
 1: 1
cuDeviceGetAttribute(Base.RefValue{Int32}, CU_DEVICE_ATTRIBUTE_MEMORY_POOLS_SUPPORTED, CuDevice(0)) = CUDA_SUCCESS
 1: 1
 = CUDA_SUCCESS
 1: 0
cuDeviceGetAttribute(Base.RefValue{Int32}, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, CuDevice(0)) = CUDA_SUCCESS
 1: 7
cuDeviceGetAttribute(Base.RefValue{Int32}, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, CuDevice(0)) = CUDA_SUCCESS
 1: 5
cuDevicePrimaryCtxRetain(Base.RefValue{Ptr{Nothing}}, CuDevice(0)) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000026454f0
cuCtxSetCurrent(CuContext(0x00000000026454f0, instance f4adfa9382129458)) = CUDA_SUCCESS
cuCtxGetDevice(Base.RefValue{Int32}) = CUDA_SUCCESS
 1: 0
cuMemGetInfo_v2(Base.RefValue{UInt64}, Base.RefValue{UInt64}) = CUDA_SUCCESS
 1: 16784883712
 2: 16908615680
cuDeviceGetAttribute(Base.RefValue{Int32}, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, CuDevice(0)) = CUDA_SUCCESS
 1: 7
cuDeviceGetAttribute(Base.RefValue{Int32}, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, CuDevice(0)) = CUDA_SUCCESS
 1: 5
cuStreamCreate(Base.RefValue{Ptr{Nothing}}, CU_STREAM_NON_BLOCKING) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000027b5eb0
cuCtxGetCurrent(Base.RefValue{Ptr{Nothing}}) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000026454f0
cuMemAllocAsync(Base.RefValue{CuPtr{Nothing}}, 4, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x0000000302000000)
pointer(a) = CuPtr{Float32}(0x0000000302000000)
cuMemAllocAsync(Base.RefValue{CuPtr{Nothing}}, 4, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x0000000302000200)
pointer(b) = CuPtr{Float32}(0x0000000302000200)
cuMemAllocAsync(Base.RefValue{CuPtr{Nothing}}, 8, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x0000000302000400)
pointer(c) = CuPtr{Float64}(0x0000000302000400)
cuMemAllocAsync(Base.RefValue{CuPtr{Nothing}}, 4, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x0000000302000600)
cuDeviceGetAttribute(Base.RefValue{Int32}, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, CuDevice(0)) = CUDA_SUCCESS
 1: 7
cuDeviceGetAttribute(Base.RefValue{Int32}, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, CuDevice(0)) = CUDA_SUCCESS
 1: 5
cuLinkCreate_v2(3, 3-element Vector{CUDA.CUjit_option_enum}, 3-element Vector{Ptr{Nothing}}, Base.RefValue{Ptr{Nothing}}) = CUDA_SUCCESS
 4: Ptr{Nothing} @0x0000000005c146f0
cuCtxGetCurrent(Base.RefValue{Ptr{Nothing}}) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000026454f0
cuLinkAddFile_v2(CuLink(0x0000000005c146f0, CuContext(0x00000000026454f0, instance f4adfa9382129458)), CU_JIT_INPUT_LIBRARY, /home/tim/Julia/depot/artifacts/53345f55f74b563606017c90cbafd05b61cd1768/lib/libcudadevrt.a, 0, Ptr{Nothing} @0x0000000000000000, Ptr{Nothing} @0x0000000000000000) = CUDA_SUCCESS
cuLinkAddData_v2(CuLink(0x0000000005c146f0, CuContext(0x00000000026454f0, instance f4adfa9382129458)), CU_JIT_INPUT_PTX, Ptr{Int8} @0x0000000005f05548, 16705, _Z27julia_broadcast_kernel_199715CuKernelContext13CuDeviceArrayI7Float32Li1ELi1EE11BroadcastedIv5TupleI5OneToI5Int64EE2__S3_I8ExtrudedIS0_IS1_Li1ELi1EES3_I4BoolES3_IS5_EES7_IS0_IS1_Li1ELi1EES3_IS8_ES3_IS5_EEEES5_, 0, Ptr{Nothing} @0x0000000000000000, Ptr{Nothing} @0x0000000000000000) = CUDA_SUCCESS
cuLinkComplete(CuLink(0x0000000005c146f0, CuContext(0x00000000026454f0, instance f4adfa9382129458)), Base.RefValue{Ptr{Nothing}}, Base.RefValue{UInt64}) = CUDA_SUCCESS
 2: Ptr{Nothing} @0x000000000606ebc8
 3: 210088
cuModuleLoadDataEx(Base.RefValue{Ptr{Nothing}}, Ptr{UInt8} @0x000000000606ebc8, 3, 3-element Vector{CUDA.CUjit_option_enum}, 3-element Vector{Ptr{Nothing}}) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x0000000005dbcc90
cuCtxGetCurrent(Base.RefValue{Ptr{Nothing}}) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000026454f0
cuModuleGetFunction(Base.RefValue{Ptr{Nothing}}, CuModule(Ptr{Nothing} @0x0000000005dbcc90, CuContext(0x00000000026454f0, instance f4adfa9382129458)), _Z27julia_broadcast_kernel_199715CuKernelContext13CuDeviceArrayI7Float32Li1ELi1EE11BroadcastedIv5TupleI5OneToI5Int64EE2__S3_I8ExtrudedIS0_IS1_Li1ELi1EES3_I4BoolES3_IS5_EES7_IS0_IS1_Li1ELi1EES3_IS8_ES3_IS5_EEEES5_) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x0000000005c41750
cuModuleGetGlobal_v2(Base.RefValue{CuPtr{Nothing}}, Base.RefValue{UInt64}, CuModule(Ptr{Nothing} @0x0000000005dbcc90, CuContext(0x00000000026454f0, instance f4adfa9382129458)), exception_flag) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x00007fda06c0f800)
 2: 8
cuMemHostAlloc(Base.RefValue{Ptr{Nothing}}, 8, 2) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00007fda01a00000
cuMemHostGetDevicePointer_v2(Base.RefValue{CuPtr{Nothing}}, Ptr{Nothing} @0x00007fda01a00000, 0) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x00007fda01a00000)
cuMemcpyHtoDAsync_v2(CuGlobal{Ptr{Nothing}}(DeviceBuffer(8 bytes at 0x00007fda06c0f800)), Base.RefValue{Ptr{Nothing}}, 8, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
 2: Ptr{Nothing} @0x00007fda01a00000
cuOccupancyMaxPotentialBlockSize(Base.RefValue{Int32}, Base.RefValue{Int32}, CuFunction(Ptr{Nothing} @0x0000000005c41750, CuModule(Ptr{Nothing} @0x0000000005dbcc90, CuContext(0x00000000026454f0, instance f4adfa9382129458))), Ptr{Nothing} @0x0000000000000000, 0, 256) = CUDA_SUCCESS
 1: 192
 2: 256
cuLaunchKernel(CuFunction(Ptr{Nothing} @0x0000000005c41750, CuModule(Ptr{Nothing} @0x0000000005dbcc90, CuContext(0x00000000026454f0, instance f4adfa9382129458))), 1, 1, 1, 1, 1, 1, 0, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458)), 3-element Vector{Ptr{Nothing}}, Ptr{Nothing} @0x0000000000000000) = CUDA_SUCCESS
cuCtxSynchronize() = CUDA_SUCCESS
cuMemAllocAsync(Base.RefValue{CuPtr{Nothing}}, 4, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x0000000302000800)
cuOccupancyMaxPotentialBlockSize(Base.RefValue{Int32}, Base.RefValue{Int32}, CuFunction(Ptr{Nothing} @0x0000000005c41750, CuModule(Ptr{Nothing} @0x0000000005dbcc90, CuContext(0x00000000026454f0, instance f4adfa9382129458))), Ptr{Nothing} @0x0000000000000000, 0, 256) = CUDA_SUCCESS
 1: 192
 2: 256
cuLaunchKernel(CuFunction(Ptr{Nothing} @0x0000000005c41750, CuModule(Ptr{Nothing} @0x0000000005dbcc90, CuContext(0x00000000026454f0, instance f4adfa9382129458))), 1, 1, 1, 1, 1, 1, 0, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458)), 3-element Vector{Ptr{Nothing}}, Ptr{Nothing} @0x0000000000000000) = CUDA_SUCCESS
cuCtxSynchronize() = CUDA_SUCCESS
cuMemFree_v2(DeviceBuffer(4 bytes at 0x0000000302000800)) = CUDA_SUCCESS
cuCtxPushCurrent_v2(CuContext(0x00000000026454f0, instance f4adfa9382129458)) = CUDA_SUCCESS
cuLinkDestroy(CuLink(0x0000000005c146f0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
cuCtxPopCurrent_v2(Base.RefValue{Ptr{Nothing}}) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000026454f0
cuMemAllocAsync(Base.RefValue{CuPtr{Nothing}}, 8, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x0000000302000800)
cuLinkCreate_v2(3, 3-element Vector{CUDA.CUjit_option_enum}, 3-element Vector{Ptr{Nothing}}, Base.RefValue{Ptr{Nothing}}) = CUDA_SUCCESS
 4: Ptr{Nothing} @0x0000000001da7520
cuCtxGetCurrent(Base.RefValue{Ptr{Nothing}}) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000026454f0
cuLinkAddFile_v2(CuLink(0x0000000001da7520, CuContext(0x00000000026454f0, instance f4adfa9382129458)), CU_JIT_INPUT_LIBRARY, /home/tim/Julia/depot/artifacts/53345f55f74b563606017c90cbafd05b61cd1768/lib/libcudadevrt.a, 0, Ptr{Nothing} @0x0000000000000000, Ptr{Nothing} @0x0000000000000000) = CUDA_SUCCESS
cuLinkAddData_v2(CuLink(0x0000000001da7520, CuContext(0x00000000026454f0, instance f4adfa9382129458)), CU_JIT_INPUT_PTX, Ptr{Int8} @0x0000000006d68248, 21912, _Z27julia_broadcast_kernel_344815CuKernelContext13CuDeviceArrayI7Float64Li1ELi1EE11BroadcastedIv5TupleI5OneToI5Int64EE2__S3_I8ExtrudedIS0_I7Float32Li1ELi1EES3_I4BoolES3_IS5_EES7_IS0_IS1_Li1ELi1EES3_IS9_ES3_IS5_EEEES5_, 0, Ptr{Nothing} @0x0000000000000000, Ptr{Nothing} @0x0000000000000000) = CUDA_SUCCESS
cuLinkComplete(CuLink(0x0000000001da7520, CuContext(0x00000000026454f0, instance f4adfa9382129458)), Base.RefValue{Ptr{Nothing}}, Base.RefValue{UInt64}) = CUDA_SUCCESS
 2: Ptr{Nothing} @0x00000000069f3b68
 3: 220016
cuModuleLoadDataEx(Base.RefValue{Ptr{Nothing}}, Ptr{UInt8} @0x00000000069f3b68, 3, 3-element Vector{CUDA.CUjit_option_enum}, 3-element Vector{Ptr{Nothing}}) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000061f3310
cuCtxGetCurrent(Base.RefValue{Ptr{Nothing}}) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000026454f0
cuModuleGetFunction(Base.RefValue{Ptr{Nothing}}, CuModule(Ptr{Nothing} @0x00000000061f3310, CuContext(0x00000000026454f0, instance f4adfa9382129458)), _Z27julia_broadcast_kernel_344815CuKernelContext13CuDeviceArrayI7Float64Li1ELi1EE11BroadcastedIv5TupleI5OneToI5Int64EE2__S3_I8ExtrudedIS0_I7Float32Li1ELi1EES3_I4BoolES3_IS5_EES7_IS0_IS1_Li1ELi1EES3_IS9_ES3_IS5_EEEES5_) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000072a2f80
cuModuleGetGlobal_v2(Base.RefValue{CuPtr{Nothing}}, Base.RefValue{UInt64}, CuModule(Ptr{Nothing} @0x00000000061f3310, CuContext(0x00000000026454f0, instance f4adfa9382129458)), exception_flag) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x00007fda06c1fa00)
 2: 8
cuMemHostAlloc(Base.RefValue{Ptr{Nothing}}, 8, 2) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00007fda01a00200
cuMemHostGetDevicePointer_v2(Base.RefValue{CuPtr{Nothing}}, Ptr{Nothing} @0x00007fda01a00000, 0) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x00007fda01a00000)
cuMemcpyHtoDAsync_v2(CuGlobal{Ptr{Nothing}}(DeviceBuffer(8 bytes at 0x00007fda06c1fa00)), Base.RefValue{Ptr{Nothing}}, 8, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
 2: Ptr{Nothing} @0x00007fda01a00000
cuMemFree_v2(DeviceBuffer(4 bytes at 0x0000000302000600)) = CUDA_SUCCESS
cuOccupancyMaxPotentialBlockSize(Base.RefValue{Int32}, Base.RefValue{Int32}, CuFunction(Ptr{Nothing} @0x00000000072a2f80, CuModule(Ptr{Nothing} @0x00000000061f3310, CuContext(0x00000000026454f0, instance f4adfa9382129458))), Ptr{Nothing} @0x0000000000000000, 0, 256) = CUDA_SUCCESS
 1: 192
 2: 256
cuLaunchKernel(CuFunction(Ptr{Nothing} @0x00000000072a2f80, CuModule(Ptr{Nothing} @0x00000000061f3310, CuContext(0x00000000026454f0, instance f4adfa9382129458))), 1, 1, 1, 1, 1, 1, 0, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458)), 3-element Vector{Ptr{Nothing}}, Ptr{Nothing} @0x0000000000000000) = CUDA_SUCCESS
cuCtxSynchronize() = CUDA_ERROR_LAUNCH_FAILED

========= Invalid __global__ write of size 8 bytes
=========     at 0x9a0 in julia_broadcast_kernel_3448(CuKernelContext,CuDeviceArray<Float64,int=1,int=1>,Broadcasted<void,Tuple<OneTo<Int64>>,__,Broadcasted<Extruded<CuDeviceArray<Float32,int=1,int=1>,Broadcasted<Bool>,Broadcasted<OneTo>>,OneTo<Int64,CuDeviceArray<Float64,int=1,int=1>,Broadcasted<__>,Broadcasted<OneTo>>>>,OneTo)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x302000800 is out of bounds
========= 

Relevant parts:

cuMemAllocAsync(Base.RefValue{CuPtr{Nothing}}, 4, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x0000000302000800)
cuLaunchKernel(CuFunction(Ptr{Nothing} @0x0000000005c41750, CuModule(Ptr{Nothing} @0x0000000005dbcc90, CuContext(0x00000000026454f0, instance f4adfa9382129458))), 1, 1, 1, 1, 1, 1, 0, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458)), 3-element Vector{Ptr{Nothing}}, Ptr{Nothing} @0x0000000000000000) = CUDA_SUCCESS
cuCtxSynchronize() = CUDA_SUCCESS
cuMemFree_v2(DeviceBuffer(4 bytes at 0x0000000302000800)) = CUDA_SUCCESS

cuMemAllocAsync(Base.RefValue{CuPtr{Nothing}}, 8, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458))) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x0000000302000800)
cuLaunchKernel(CuFunction(Ptr{Nothing} @0x00000000072a2f80, CuModule(Ptr{Nothing} @0x00000000061f3310, CuContext(0x00000000026454f0, instance f4adfa9382129458))), 1, 1, 1, 1, 1, 1, 0, CuStream(0x00000000027b5eb0, CuContext(0x00000000026454f0, instance f4adfa9382129458)), 3-element Vector{Ptr{Nothing}}, Ptr{Nothing} @0x0000000000000000) = CUDA_SUCCESS

cuCtxSynchronize() = CUDA_ERROR_LAUNCH_FAILED
========= Invalid __global__ write of size 8 bytes
=========     at 0x9a0 in julia_broadcast_kernel_3448(CuKernelContext,CuDeviceArray<Float64,int=1,int=1>,Broadcasted<void,Tuple<OneTo<Int64>>,__,Broadcasted<Extruded<CuDeviceArray<Float32,int=1,int=1>,Broadcasted<Bool>,Broadcasted<OneTo>>,OneTo<Int64,CuDeviceArray<Float64,int=1,int=1>,Broadcasted<__>,Broadcasted<OneTo>>>>,OneTo)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x302000800 is out of bounds
========= 

So we're allocating and freeing a buffer at 0x302000800, after which a new allocation requests returns 0x302000800 again. All of this happens asynchronously on the same stream, but compute-sanitizer seems to think the pointer is invalid. I think this is a compute-sanitizer bug?

[maleadt]

Reported as NVIDIA bug 3279048.

[aladram]

As a temporary workaround, feel free to replace cuMemFree with cuMemFreeAsync on the NULL stream followed by a synchronization (works on my side).

diff --git a/src/array.jl b/src/array.jl
index 4294004a..535db339 100644
--- a/src/array.jl
+++ b/src/array.jl
@@ -50,7 +50,10 @@ function unsafe_free!(xs::CuArray, finalizer::Bool=false)
     throw(ArgumentError("Cannot free an unmanaged buffer."))
   end

-  @context! skip_destroyed=true xs.ctx free(xs.baseptr; stream_ordered=!finalizer)
+  @context! skip_destroyed=true xs.ctx free(xs.baseptr)
+  if finalizer
+    synchronize()
+  end
   xs.state = ARRAY_FREED

   # the object is dead, so we can also wipe the pointer

[maleadt]

Thanks. Is this only limited to compute-sanitizer reporting something incorrectly, or would it also impact regular execution? I already run compute-sanitizer CI tests with the stream-ordered allocator disabled (i.e. always using cuMemFree, not cuMemFreeAsync).

[aladram]

It should not impact regular execution (compute-sanitizer bug). FYI the problem is the following: when an allocation is created with cuMemAllocAsync, subsequent calls to cuMemFree on this allocation are incorrectly reported to the tool.

I already run compute-sanitizer CI tests with the stream-ordered allocator disabled (i.e. always using cuMemFree, not cuMemFreeAsync).

Doing this does not disable the stream-ordered allocator: an allocation created with cuMemFreeAsync will be issued from the stream's pool, and as soon as cuMemFreeAsync is called, it is queued to be released. You can find some clear explanations on how it works on this post from our developer blog. However, if you call cuMemFree instead, the allocation is released immediately, therefore it is your responsibility to make sure all tasks potentially using this allocation are completed (i.e. synchronization if required) prior to calling cuMemFree. That is why, in the patch I sent earlier, I added a synchronization: I wanted to make sure the allocation would be released immediately (since this patch* is a replacement for cuMemFree). I freely admit this behavior is poorly documented (I'm not even sure it is...). When you're not sure whether it is OK or not to call cuMemFree (e.g. no prior synchronization), prefer cuMemFreeAsync.

By the way, thank you for this bug report: the API trace and the summary are greatly appreciated!

[maleadt]

Thanks for the elaboration. I meant to say that on CI, when running under compute-sanitizer, we also don't use cuMemAllocAsync, so I don't run into that problem there.

But this information is good to know, because we do otherwise call cuMemFree on allocations that were created with cuMemFreeAsync. This is a recent change, #778, motivated by the fact that our finalizers (destructors called by the garbage collector) don't know which stream an allocation was (asynchronously) created on, so we can't call cuMemFreeAsync there. I figured that cuMemFree would be synchronizing, but that's not entirely the case apparently. Synchronizing the entire context before every free is probably too expensive, so I'll have to think some more to figure out how to deal with this...

[aladram]

I meant to say that on CI, when running under compute-sanitizer, we also don't use cuMemAllocAsync, so I don't run into that problem there.

Oh my bad, makes sense!

Synchronizing the entire context before every free is probably too expensive, so I'll have to think some more to figure out how to deal with this...

I see. You could call cuMemFreeAsync on the NULL stream (which is what I did in the patch, I think, minus the synchronization), but it might incur unwanted synchronizations (queuing a task on the NULL stream will force all blocking streams to synchronize before that task is executed, c.f. https://developer.nvidia.com/blog/gpu-pro-tip-cuda-7-streams-simplify-concurrency/, "A Multi-threading Example"). However, if you're 100% sure that your allocation is not used anymore (meaning all tasks using it are completed), calling cuMemFree is totally fine, and will not incur any synchronization if the allocation is stream-ordered.

Let's say you reference-count your allocations, and there is no way for the device to access your allocation if your refcount is null. Then when you create your CuArray, you create a mem object with a refcount of 1, increment this refcount each time you launch a kernel using this allocation, and add a callback after any kernel launch using that allocation is completed (using cuLaunchHostFunc on the same stream you launched your kernel on) that decrement this refcount. On CuArray's destruction, you would just decrease the refcount on the mem object. If that refcount reaches 0, then you know you can safely call cuMemFree. Of coure the idea would be to do this using objects so you can take advantage of the garbage collection. This is not possible in CUDA since our users can pass data to their kernels that might contain a reference to anything, but maybe if you have more control it could be possible in your case. Unfortunately I'm not familiar enough with Julia and/or your project to make any assertion on the feasibility of such a thing. Let me know what you think about this!

[maleadt]

You could call cuMemFreeAsync on the NULL stream (which is what I did in the patch, I think, minus the synchronization), but it might incur unwanted synchronizations (queuing a task on the NULL stream will force all blocking streams to synchronize before that task is executed, c.f. https://developer.nvidia.com/blog/gpu-pro-tip-cuda-7-streams-simplify-concurrency/, "A Multi-threading Example").

Ah yes, so I don't even need to call device_synchronize/cuCtxSynchronize in #783, just using the default stream should be fine (our task-local streams are not explicitly non-blocking, and we don't use the ptds/ptsz API variants so we're still using the legacy stream behavior as far as I understand).

Then when you create your CuArray, you create a mem object with a refcount of 1, increment this refcount each time you launch a kernel using this allocation, and add a callback after any kernel launch using that allocation is completed (using cuLaunchHostFunc on the same stream you launched your kernel on) that decrement this refcount.

That's a neat idea (cc @vchuravy, you suggested something similar). I was afraid it would quickly become complicated though, because IIUC we'd need this for every asynchronous operation, i.e., not only kernel launches, but also memory operations, library calls (CUBLAS, CUDNN, ...), etc.

EDIT: actually, we do create our streams with STREAM_NON_BLOCKING, I'll have to change that.

[aladram]

I was afraid it would quickly become complicated though

It might. Just so I know, why can't you keep track of the stream in your CuArray object? And maybe if the stream doesn't exist anymore use the NULL stream them. Because what I said is more or less recreating a stream-ordered allocator... (EDIT: actually not really because with the method I described it doesn't matter anymore which stream you use). If you want ideas on how to create a neat memory management system, you could take a look at RAPIDS Memory Manager (read the "Using RMM in C++" section of the README).

[maleadt]

Just so I know, why can't you keep track of the stream in your CuArray object?

Because I've made the stream task-local (i.e. each task gets its own stream so that we can do task-based concurrency while GPU operations can overlap), yet it is possible to create an array on one task and pass it to another. For example:

# global code executes in the 'main task'

# create data
a = CuArray{Int}(undef, 1)

# make sure the data is actually available
synchronize()

task = @async do
  # this will execute on a different stream
  b = a .+ 1
end

# at this point, the GC might notice there aren't any uses of `a` anymore, so it's possible to garbage collect it.
# however, there's uses on another task's stream, so we can't free on the stream it was allocated on.

wait(task)

This task-based GPU programming model is a fairly recent development though, and I'm still figuring out what the best way is to map it into CUDA APIs (providing an easy/safe programming model without requiring lots of synchronization behind the scenes).

[aladram]

I see. Maybe CUDA graphs will help you then (sorry I'm sending a lot of links haha): to the best of my knowledge, it is the closest functionality we have that resemble task-based programming (you setup a dependency graph and run it, and I think it can include GPU or CPU tasks).

[maleadt]

No worries, thanks for the links. I'm thinking of first going with a maximally-correct version that synchronizes when freeing from finalizers, and improve it as we go / if it turns out problematic. Limited manpower, and all that.

While verifying #783, I'm confused about the behavior of cuMemFreeAsync though; it doesn't seem to respect legacy stream semantics. Given the following example:

println("start")
a = CuArray{Int}(undef, 1024, 1024, 1024)
device_synchronize()

# submit work from a different task, so using a different stream
t = @async begin
    println("start submit")
    a .= sin.(a)
    println("end submit")
end
wait(t)

# now force a collection, which will free `a` since the task has been completed.
# but since execution on the GPU hasn't, we need to block, so we'll use the default stream.
println("free")
a = nothing
GC.gc(true)

println("sync")
device_synchronize()

println("done")

Using #783, the free that gets triggered by the GC ends up using the default stream:

start
cuStreamCreate(Base.RefValue{Ptr{Nothing}}, CU_STREAM_DEFAULT) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x0000000003cc4df0
cuMemAllocAsync(Base.RefValue{CuPtr{Nothing}}, 8589934592, CuStream(0x0000000003cc4df0, CuContext(0x0000000002a447d0, instance a147b90c9f6c2143))) = CUDA_SUCCESS
 1: CuPtr{Nothing}(0x0000000302000000)
cuCtxSynchronize() = CUDA_SUCCESS


start submit
cuStreamCreate(Base.RefValue{Ptr{Nothing}}, CU_STREAM_DEFAULT) = CUDA_SUCCESS
 1: Ptr{Nothing} @0x00000000053af400
cuLaunchKernel(CuFunction(Ptr{Nothing} @0x00000000068c2e60, CuModule(Ptr{Nothing} @0x00000000052a98a0, CuContext(0x0000000002a447d0, instance a147b90c9f6c2143))), 4793491, 1, 1, 224, 1, 1, 0, CuStream(0x00000000053af400, CuContext(0x0000000002a447d0, instance a147b90c9f6c2143)), 3-element Vector{Ptr{Nothing}}, Ptr{Nothing} @0x0000000000000000) = CUDA_SUCCESS
end submit


free
cuMemFreeAsync(DeviceBuffer(8.000 GiB at 0x0000000302000000), CuStream(0x0000000000000000, nothing)) = CUDA_SUCCESS


sync
cuCtxSynchronize() = CUDA_SUCCESS

done

Checking with NSight Systems:

The cuMemFreeAsync returns immediately, even though we're using the default stream. I've used CU_STREAM_DEFAULT, so the streams aren't explicitly non-blocking. And explicitly using the legacy stream object 0x1 doesn't help here either. Am I misinterpreting the legacy stream behavior, and should I just explicitly call device_synchronize()? Or will this cuMemFreeAsync call only effectively free after all other streams have finished?

[aladram]

I'm thinking of first going with a maximally-correct version that synchronizes when freeing from finalizers, and improve it as we go / if it turns out problematic.

Makes sense.

The cuMemFreeAsync returns immediately, even though we're using the default stream.

cuMemFreeAsync called on the NULL stream always returns immediately, but will queue a barrier between all blocking streams (c.f. documentation for cuMemFreeAsync). There is a synchronization, but it happens device-side, not on your CPU thread. cuMemAllocAsync will do the same, and it will even return immediately a pointer, but the memory at the given address will not be available for use before the queued operation is executed (because the driver has knowledge of all previous asynchronous operations, including the ones that are queued but not executed yet, it can compute which part of the memory will be available before any of it actually happened). To the best of my knowledge, all cu*Async operations are asynchronous and should return immediately.

I apologize if I created any confusion with this synchronization term. I hope these explanations are clearer. If you have more questions I'll be more than happy to answer them!