@@ -0,0 +1,35 @@
+On-disk Kernel Caching
+======================
+
+When the ``cache`` keyword argument of the :func:`@cuda.jit <numba.cuda.jit>`
+decorator is ``True``, a file-based cache is enabled. This shortens compilation
+times when the function was already compiled in a previous invocation.
+
+The cache is maintained in the ``__pycache__`` subdirectory of the directory
+containing the source file; if the current user is not allowed to write to it,
+the cache implementation falls back to a platform-specific user-wide cache
+directory (such as ``$HOME/.cache/numba`` on Unix platforms).
+
+
+Compute capability considerations
+---------------------------------
+
+Separate cache files are maintained for each compute capability. When a cached
+kernel is loaded, the compute capability of the device the kernel is first
+launched on in the current run is used to determine which version to load.
+Therefore, on systems that have multiple GPUs with differing compute
+capabilities, the cached versions of kernels are only used for one compute
+capability, and recompilation will occur for other compute capabilities.
+
+For example: if a system has two GPUs, one of compute capability 7.5 and one of
+8.0, then:
+
+* If a cached kernel is first launched on the CC 7.5 device, then the cached
+  version for CC 7.5 is used. If it is subsequently launched on the CC 8.0
+  device, a recompilation will occur.
+* If in a subsequent run the cached kernel is first launched on the CC 8.0
+  device, then the cached version for CC 8.0 is used. A subsequent launch on
+  the CC 7.5 device will require a recompilation.
+
+This limitation is not expected to present issues in most practical scenarios,
+as multi-GPU production systems tend to have identical GPUs within each node.
@@ -25,4 +25,5 @@ Numba for CUDA GPUs
    external-memory.rst
    bindings.rst
    cuda_ffi.rst
+   caching.rst
    faq.rst
@@ -5,7 +5,7 @@
 from numba.core import config, serialize
 from numba.core.codegen import Codegen, CodeLibrary
 from numba.core.errors import NumbaInvalidConfigWarning
-from .cudadrv import devices, driver, nvvm
+from .cudadrv import devices, driver, nvvm, runtime
 
 import ctypes
 import numpy as np
@@ -302,7 +302,7 @@ def _reduce_states(self):
                    'libraries to link against')
             raise RuntimeError(msg)
         return dict(
-            codegen=self._codegen,
+            codegen=None,
             name=self.name,
             entry_name=self._entry_name,
             module=self._module,
@@ -337,6 +337,8 @@ def _rebuild(cls, codegen, name, entry_name, module, linking_libraries,
         instance._max_registers = max_registers
         instance._nvvm_options = nvvm_options
 
+        return instance
+
 
 class JITCUDACodegen(Codegen):
     """
@@ -360,3 +362,11 @@ def _create_empty_module(self, name):
 
     def _add_module(self, module):
         pass
+
+    def magic_tuple(self):
+        """
+        Return a tuple unambiguously describing the codegen behaviour.
+        """
+        ctx = devices.get_context()
+        cc = ctx.device.compute_capability
+        return (runtime.runtime.get_version(), cc)
@@ -12,7 +12,7 @@
 
 
 def jit(func_or_sig=None, device=False, inline=False, link=[], debug=None,
-        opt=True, **kws):
+        opt=True, cache=False, **kws):
     """
     JIT compile a python function conforming to the CUDA Python specification.
     If a signature is supplied, then a function is returned that takes a
@@ -49,6 +49,8 @@ def jit(func_or_sig=None, device=False, inline=False, link=[], debug=None,
        assembly code. This enables inspection of the source code in NVIDIA
        profiling tools and correlation with program counter sampling.
     :type lineinfo: bool
+    :param cache: If True, enables the file-based cache for this function.
+    :type cache: bool
     """
 
     if link and config.ENABLE_CUDASIM:
@@ -102,6 +104,9 @@ def _jit(func):
 
             disp = CUDADispatcher(func, targetoptions=targetoptions)
 
+            if cache:
+                disp.enable_caching()
+
             if device:
                 from numba.core import typeinfer
                 with typeinfer.register_dispatcher(disp):
@@ -124,7 +129,7 @@ def autojitwrapper(func):
             else:
                 def autojitwrapper(func):
                     return jit(func, device=device, debug=debug, opt=opt,
-                               link=link, **kws)
+                               link=link, cache=cache, **kws)
 
             return autojitwrapper
         # func_or_sig is a function
@@ -140,7 +145,12 @@ def autojitwrapper(func):
                 targetoptions['fastmath'] = fastmath
                 targetoptions['device'] = device
                 targetoptions['extensions'] = extensions
-                return CUDADispatcher(func_or_sig, targetoptions=targetoptions)
+                disp = CUDADispatcher(func_or_sig, targetoptions=targetoptions)
+
+                if cache:
+                    disp.enable_caching()
+
+                return disp
 
 
 def declare_device(name, sig):
 
@@ -6,6 +6,7 @@
 import warnings
 
 from numba.core import config, serialize, sigutils, types, typing, utils
+from numba.core.caching import Cache, CacheImpl
 from numba.core.compiler_lock import global_compiler_lock
 from numba.core.dispatcher import Dispatcher
 from numba.core.errors import NumbaPerformanceWarning, NumbaDeprecationWarning
@@ -106,23 +107,50 @@ def __init__(self, py_func, argtypes, link=None, debug=False,
         self._codelibrary = lib
         self.call_helper = cres.call_helper
 
+        # The following are referred to by the cache implementation. Note:
+        # - There are no referenced environments in CUDA.
+        # - Kernels don't have lifted code.
+        # - reload_init is only for parfors.
+        self.target_context = tgt_ctx
+        self.fndesc = cres.fndesc
+        self.environment = cres.environment
+        self._referenced_environments = []
+        self.lifted = []
+        self.reload_init = []
+
+    @property
+    def library(self):
+        return self._codelibrary
+
+    @property
+    def type_annotation(self):
+        return self._type_annotation
+
+    def _find_referenced_environments(self):
+        return self._referenced_environments
+
+    @property
+    def codegen(self):
+        return self.target_context.codegen()
+
     @property
     def argument_types(self):
         return tuple(self.signature.args)
 
     @classmethod
-    def _rebuild(cls, cooperative, name, argtypes, codelibrary, link, debug,
-                 lineinfo, call_helper, extensions):
+    def _rebuild(cls, cooperative, name, signature, codelibrary,
+                 debug, lineinfo, call_helper, extensions):
         """
         Rebuild an instance.
         """
         instance = cls.__new__(cls)
         # invoke parent constructor
         super(cls, instance).__init__()
         # populate members
+        instance.entry_point = None
         instance.cooperative = cooperative
         instance.entry_name = name
-        instance.argument_types = tuple(argtypes)
+        instance.signature = signature
         instance._type_annotation = None
         instance._codelibrary = codelibrary
         instance.debug = debug
@@ -140,7 +168,7 @@ def _reduce_states(self):
         Stream information is discarded.
         """
         return dict(cooperative=self.cooperative, name=self.entry_name,
-                    argtypes=self.argtypes, codelibrary=self.codelibrary,
+                    signature=self.signature, codelibrary=self._codelibrary,
                     debug=self.debug, lineinfo=self.lineinfo,
                     call_helper=self.call_helper, extensions=self.extensions)
 
@@ -479,6 +507,31 @@ def __call__(self, *args):
                                     self.stream, self.sharedmem)
 
 
+class CUDACacheImpl(CacheImpl):
+    def reduce(self, kernel):
+        return kernel._reduce_states()
+
+    def rebuild(self, target_context, payload):
+        return _Kernel._rebuild(**payload)
+
+    def check_cachable(self, cres):
+        # CUDA Kernels are always cachable - the reasons for an entity not to
+        # be cachable are:
+        #
+        # - The presence of lifted loops, or
+        # - The presence of dynamic globals.
+        #
+        # neither of which apply to CUDA kernels.
+        return True
+
+
+class CUDACache(Cache):
+    """
+    Implements a cache that saves and loads CUDA kernels and compile results.
+    """
+    _impl_class = CUDACacheImpl
+
+
 class CUDADispatcher(Dispatcher, serialize.ReduceMixin):
     '''
     CUDA Dispatcher object. When configured and called, the dispatcher will
@@ -517,6 +570,9 @@ def __init__(self, py_func, targetoptions, pipeline_class=CUDACompiler):
     def _numba_type_(self):
         return cuda_types.CUDADispatcher(self)
 
+    def enable_caching(self):
+        self._cache = CUDACache(self.py_func)
+
     @functools.lru_cache(maxsize=128)
     def configure(self, griddim, blockdim, stream=0, sharedmem=0):
         griddim, blockdim = normalize_kernel_dimensions(griddim, blockdim)
@@ -718,30 +774,45 @@ def compile_device(self, args):
 
         return cres
 
+    def add_overload(self, kernel, argtypes):
+        c_sig = [a._code for a in argtypes]
+        self._insert(c_sig, kernel, cuda=True)
+        self.overloads[argtypes] = kernel
+
     def compile(self, sig):
         '''
         Compile and bind to the current context a version of this kernel
         specialized for the given signature.
         '''
         argtypes, return_type = sigutils.normalize_signature(sig)
         assert return_type is None or return_type == types.none
+
+        # Do we already have an in-memory compiled kernel?
         if self.specialized:
             return next(iter(self.overloads.values()))
         else:
             kernel = self.overloads.get(argtypes)
-        if kernel is None:
+            if kernel is not None:
+                return kernel
+
+        # Can we load from the disk cache?
+        kernel = self._cache.load_overload(sig, self.targetctx)
+
+        if kernel is not None:
+            self._cache_hits[sig] += 1
+        else:
+            # We need to compile a new kernel
+            self._cache_misses[sig] += 1
             if not self._can_compile:
                 raise RuntimeError("Compilation disabled")
-            kernel = _Kernel(self.py_func, argtypes,
-                             **self.targetoptions)
-            # Inspired by _DispatcherBase.add_overload, but differs slightly
-            # because we're inserting a _Kernel object instead of a compiled
-            # function.
-            c_sig = [a._code for a in argtypes]
-            self._insert(c_sig, kernel, cuda=True)
-            self.overloads[argtypes] = kernel
 
+            kernel = _Kernel(self.py_func, argtypes, **self.targetoptions)
+            # We call bind to force codegen, so that there is a cubin to cache
             kernel.bind()
+            self._cache.save_overload(sig, kernel)
+
+        self.add_overload(kernel, argtypes)
+
         return kernel
 
     def inspect_llvm(self, signature=None):
 
@@ -77,7 +77,7 @@ def elapsed_time(self, event):
 
 def jit(func_or_sig=None, device=False, debug=False, argtypes=None,
         inline=False, restype=None, fastmath=False, link=None,
-        boundscheck=None, opt=True
+        boundscheck=None, opt=True, cache=None
         ):
     # Here for API compatibility
     if boundscheck: