Add ability to link CUDA functions with in-memory PTX.

docs/source/cuda-reference/kernel.rst

@@ -10,6 +10,14 @@ be configured and launched:
 .. autofunction:: numba.cuda.jit
 
 
+.. _cuda-ptx-code:
+
+The ``cuda.PTXCode`` class is used to store in-memory PTX code to be linked
+with a CUDA dispatcher object
+
+.. autoclass:: numba.cuda.PTXCode
+
+
 Dispatcher objects
 ------------------
 

numba/cuda/codegen.py

@@ -13,6 +13,18 @@
 CUDA_TRIPLE = 'nvptx64-nvidia-cuda'
 
 
+class PTXCode :
+    """
+    Store in-memory PTX code to link with a compiled function.
+    """
+
+    def __init__(self, ptx) :
+        self._ptx = ptx
+
+    def __str__(self) :
+        return self._ptx
+
+
 def run_nvdisasm(cubin, flags):
     # nvdisasm only accepts input from a file, so we need to write out to a
     # temp file and clean up afterwards.
@@ -169,7 +181,10 @@ def get_cubin(self, cc=None):
         for ptx in ptxes:
             linker.add_ptx(ptx.encode())
         for path in self._linking_files:
-            linker.add_file_guess_ext(path)
+            if isinstance(path, PTXCode):
+                linker.add_ptx(path._ptx.encode(encoding='ascii'))
+            else :
+                linker.add_file_guess_ext(path)
         if self.needs_cudadevrt:
             linker.add_file_guess_ext(get_cudalib('cudadevrt', static=True))
 

numba/cuda/decorators.py

@@ -28,8 +28,8 @@ def jit(func_or_sig=None, device=False, inline=False, link=[], debug=None,
        .. note:: A kernel cannot have any return value.
     :param device: Indicates whether this is a device function.
     :type device: bool
-    :param link: A list of files containing PTX or CUDA C/C++ source to link
-       with the function
+    :param link: A list of files containing PTX or CUDA C/C++ source or
+      :class:`PTXCode <numba.cuda.PTXCode>` to link with the function.
     :type link: list
     :param debug: If True, check for exceptions thrown when executing the
        kernel. Since this degrades performance, this should only be used for

numba/cuda/device_init.py

@@ -22,6 +22,7 @@
 from .errors import KernelRuntimeError
 
 from .decorators import jit, declare_device
+from .codegen import PTXCode
 from .api import *
 from .api import _auto_device
 from .args import In, Out, InOut