Consider running GC when allocating and synchronizing

docs/src/lib/driver.md

@@ -138,7 +138,7 @@ methods then work with these raw pointers:
 
 ```@docs
 CUDA.memory_status
-CUDA.available_memory
+CUDA.free_memory
 CUDA.total_memory
 ```
 

lib/cudadrv/memory.jl

@@ -767,11 +767,11 @@ end
 end # module Mem
 
 """
-    available_memory()
+    free_memory()
 
-Returns the available amount of memory (in bytes), available for allocation by the CUDA context.
+Returns the free amount of memory (in bytes), available for allocation by the CUDA context.
 """
-available_memory() = Mem.info()[1]
+free_memory() = Mem.info()[1]
 
 """
     total_memory()

lib/cudadrv/synchronization.jl

@@ -182,9 +182,11 @@ function device_synchronize(; blocking::Bool=false, spin::Bool=true)
         if spin && spinning_synchronization(isdone, legacy_stream())
             cuCtxSynchronize()
         else
+            maybe_collect(true)
             nonblocking_synchronize(context())
         end
     else
+        maybe_collect(true)
         cuCtxSynchronize()
     end
 
@@ -196,9 +198,11 @@ function synchronize(stream::CuStream=stream(); blocking::Bool=false, spin::Bool
         if spin && spinning_synchronization(isdone, stream)
             cuStreamSynchronize(stream)
         else
+            maybe_collect(true)
             nonblocking_synchronize(stream)
         end
     else
+        maybe_collect(true)
         cuStreamSynchronize(stream)
     end
 
@@ -210,9 +214,11 @@ function synchronize(event::CuEvent; blocking::Bool=false, spin::Bool=true)
         if spin && spinning_synchronization(isdone, event)
             cuEventSynchronize(event)
         else
+            maybe_collect(true)
             nonblocking_synchronize(event)
         end
     else
+        maybe_collect(true)
         cuEventSynchronize(event)
     end
 end
@@ -269,6 +275,7 @@ function device_synchronize(; blocking::Bool=false, spin::Bool=true)
             nonblocking_synchronize(stream)
         end
     end
+    maybe_collect(true)
     cuCtxSynchronize()
 
     check_exceptions()
@@ -280,6 +287,7 @@ function synchronize(stream::CuStream=stream(); blocking::Bool=false, spin::Bool
             nonblocking_synchronize(stream)
         end
     end
+    maybe_collect(true)
     cuStreamSynchronize(stream)
 
     check_exceptions()
@@ -289,6 +297,7 @@ function synchronize(event::CuEvent; blocking::Bool=false, spin::Bool=true)
     if use_nonblocking_synchronization && !blocking
         spin && spinning_synchronization(isdone, event)
     end
+    maybe_collect(true)
     cuEventSynchronize(event)
 end
 

lib/cudnn/src/convolution.jl

@@ -331,6 +331,6 @@ function cudnnFindConvolutionAlgorithmWorkspaceSize(x)
     # Because algorithm discovery runs infrequently yet allocates more than conv functions,
     # This is a good place to synchronize and trim the memory pool to reduce fragmentation.
     CUDA.reclaim()
-    gpufree = CUDA.available_memory() + coalesce(CUDA.cached_memory(), 0)
+    gpufree = CUDA.free_memory() + coalesce(CUDA.cached_memory(), 0)
     min(gpufree ÷ 10, sizeof(x) * 100)
 end

src/CUDA.jl

@@ -125,6 +125,7 @@ export CUDABackend
 # StaticArrays is still a direct dependency, so directly include the extension
 include("../ext/StaticArraysExt.jl")
 
+include("deprecated.jl")
 include("precompile.jl")
 
 end

src/deprecated.jl

@@ -0,0 +1 @@
+@deprecate available_memory() free_memory()

src/initialization.jl

@@ -72,9 +72,9 @@ function __init__()
         return
     end
 
-    if driver < v"11.2"
+    if driver < v"11.3"
         @warn """The NVIDIA driver on this system only supports up to CUDA $driver.
-                 For performance reasons, it is recommended to upgrade to a driver that supports CUDA 11.2 or higher."""
+                 For performance reasons, it is recommended to upgrade to a driver that supports CUDA 11.3 or higher."""
     end
 
     # check that we have a runtime