Index: sasmodels/kernelcuda.py
===================================================================
--- sasmodels/kernelcuda.py	(revision 00afc158f4641f755fc8518e7528d23f8655d30b)
+++ sasmodels/kernelcuda.py	(revision 3199b177cb91c6bd74017b7f942ddbbbc9bb46c0)
@@ -68,5 +68,5 @@
 
 
-# Attempt to setup cuda. This may fail if the pycuda package is not
+# Attempt to setup CUDA. This may fail if the pycuda package is not
 # installed or if it is installed but there are no devices available.
 try:
@@ -108,7 +108,9 @@
 MAX_LOOPS = 2048
 
+
 def use_cuda():
-    env = os.environ.get("SAS_OPENCL", "").lower()
-    return HAVE_CUDA and (env == "" or env.startswith("cuda"))
+    sas_opencl = os.environ.get("SAS_OPENCL", "CUDA").lower()
+    return HAVE_CUDA and sas_opencl.startswith("cuda")
+
 
 ENV = None
@@ -122,4 +124,5 @@
         ENV.release()
     ENV = GpuEnvironment() if use_cuda() else None
+
 
 def environment():
@@ -139,11 +142,4 @@
     return ENV
 
-def free_context():
-    global ENV
-    if ENV is not None:
-        ENV.release()
-        ENV = None
-
-atexit.register(free_context)
 
 def has_type(dtype):
@@ -152,13 +148,13 @@
     Return true if device supports the requested precision.
     """
-    # Assume the nvidia card supports 32-bit and 64-bit floats.
-    # TODO: check if pycuda support F16
+    # Assume the NVIDIA card supports 32-bit and 64-bit floats.
+    # TODO: Check if pycuda support F16.
     return dtype in (generate.F32, generate.F64)
 
 
 FUNCTION_PATTERN = re.compile(r"""^
-  (?P<space>\s*)                   # initial space
-  (?P<qualifiers>^(?:\s*\b\w+\b\s*)+) # one or more qualifiers before function
-  (?P<function>\s*\b\w+\b\s*[(])      # function name plus open parens
+  (?P<space>\s*)                       # Initial space.
+  (?P<qualifiers>^(?:\s*\b\w+\b\s*)+)  # One or more qualifiers before function.
+  (?P<function>\s*\b\w+\b\s*[(])       # Function name plus open parens.
   """, re.VERBOSE|re.MULTILINE)
 
@@ -167,7 +163,8 @@
   """, re.VERBOSE|re.MULTILINE)
 
+
 def _add_device_tag(match):
     # type: (None) -> str
-    # Note: should be re.Match, but that isn't a simple type
+    # Note: Should be re.Match, but that isn't a simple type.
     """
     replace qualifiers with __device__ qualifiers if needed
@@ -182,4 +179,5 @@
         return "".join((space, "__device__ ", qualifiers, function))
 
+
 def mark_device_functions(source):
     # type: (str) -> str
@@ -188,4 +186,5 @@
     """
     return FUNCTION_PATTERN.sub(_add_device_tag, source)
+
 
 def show_device_functions(source):
@@ -197,4 +196,5 @@
         print(match.group('qualifiers').replace('\n',r'\n'), match.group('function'), '(')
     return source
+
 
 def compile_model(source, dtype, fast=False):
@@ -221,5 +221,5 @@
     #options = ['--verbose', '-E']
     options = ['--use_fast_math'] if fast else None
-    program = SourceModule(source, no_extern_c=True, options=options) # include_dirs=[...]
+    program = SourceModule(source, no_extern_c=True, options=options) #, include_dirs=[...])
 
     #print("done with "+program)
@@ -227,20 +227,17 @@
 
 
-# for now, this returns one device in the context
-# TODO: create a context that contains all devices on all platforms
+# For now, this returns one device in the context.
+# TODO: Create a context that contains all devices on all platforms.
 class GpuEnvironment(object):
     """
-    GPU context, with possibly many devices, and one queue per device.
+    GPU context for CUDA.
     """
     context = None # type: cuda.Context
     def __init__(self, devnum=None):
         # type: (int) -> None
-        # Byte boundary for data alignment
-        #self.data_boundary = max(d.min_data_type_align_size
-        #                         for d in self.context.devices)
-        self.compiled = {}
         env = os.environ.get("SAS_OPENCL", "").lower()
         if devnum is None and env.startswith("cuda:"):
             devnum = int(env[5:])
+
         # Set the global context to the particular device number if one is
         # given, otherwise use the default context.  Perhaps this will be set
@@ -251,4 +248,11 @@
             self.context = make_default_context()
 
+        ## Byte boundary for data alignment.
+        #self.data_boundary = max(d.min_data_type_align_size
+        #                         for d in self.context.devices)
+
+        # Cache for compiled programs, and for items in context.
+        self.compiled = {}
+
     def release(self):
         if self.context is not None:
@@ -271,10 +275,9 @@
         Compile the program for the device in the given context.
         """
-        # Note: PyOpenCL caches based on md5 hash of source, options and device
-        # so we don't really need to cache things for ourselves.  I'll do so
-        # anyway just to save some data munging time.
+        # Note: PyCuda (probably) caches but I'll do so as well just to
+        # save some data munging time.
         tag = generate.tag_source(source)
         key = "%s-%s-%s%s"%(name, dtype, tag, ("-fast" if fast else ""))
-        # Check timestamp on program
+        # Check timestamp on program.
         program, program_timestamp = self.compiled.get(key, (None, np.inf))
         if program_timestamp < timestamp:
@@ -286,4 +289,5 @@
         return program
 
+
 class GpuModel(KernelModel):
     """
@@ -301,10 +305,10 @@
     that the compiler is allowed to take shortcuts.
     """
-    info = None # type: ModelInfo
-    source = "" # type: str
-    dtype = None # type: np.dtype
-    fast = False # type: bool
-    _program = None # type: SourceModule
-    _kernels = None # type: Dict[str, cuda.Function]
+    info = None  # type: ModelInfo
+    source = ""  # type: str
+    dtype = None  # type: np.dtype
+    fast = False  # type: bool
+    _program = None  # type: SourceModule
+    _kernels = None  # type: Dict[str, cuda.Function]
 
     def __init__(self, source, model_info, dtype=generate.F32, fast=False):
@@ -352,8 +356,9 @@
         functions = [program.get_function(k) for k in names]
         self._kernels = {k: v for k, v in zip(variants, functions)}
-        # keep a handle to program so GC doesn't collect
+        # Keep a handle to program so GC doesn't collect.
         self._program = program
 
-# TODO: check that we don't need a destructor for buffers which go out of scope
+
+# TODO: Check that we don't need a destructor for buffers which go out of scope.
 class GpuInput(object):
     """
@@ -377,21 +382,19 @@
     def __init__(self, q_vectors, dtype=generate.F32):
         # type: (List[np.ndarray], np.dtype) -> None
-        # TODO: do we ever need double precision q?
+        # TODO: Do we ever need double precision q?
         self.nq = q_vectors[0].size
         self.dtype = np.dtype(dtype)
         self.is_2d = (len(q_vectors) == 2)
-        # TODO: stretch input based on get_warp()
-        # not doing it now since warp depends on kernel, which is not known
+        # TODO: Stretch input based on get_warp().
+        # Not doing it now since warp depends on kernel, which is not known
         # at this point, so instead using 32, which is good on the set of
         # architectures tested so far.
         if self.is_2d:
-            # Note: 16 rather than 15 because result is 1 longer than input.
-            width = ((self.nq+16)//16)*16
+            width = ((self.nq+15)//16)*16
             self.q = np.empty((width, 2), dtype=dtype)
             self.q[:self.nq, 0] = q_vectors[0]
             self.q[:self.nq, 1] = q_vectors[1]
         else:
-            # Note: 32 rather than 31 because result is 1 longer than input.
-            width = ((self.nq+32)//32)*32
+            width = ((self.nq+31)//32)*32
             self.q = np.empty(width, dtype=dtype)
             self.q[:self.nq] = q_vectors[0]
@@ -399,5 +402,5 @@
         #print("creating inputs of size", self.global_size)
 
-        # transfer input value to gpu
+        # Transfer input value to GPU.
         self.q_b = cuda.to_device(self.q)
 
@@ -405,5 +408,5 @@
         # type: () -> None
         """
-        Free the memory.
+        Free the buffer associated with the q value.
         """
         if self.q_b is not None:
@@ -414,4 +417,5 @@
         # type: () -> None
         self.release()
+
 
 class GpuKernel(Kernel):
@@ -429,12 +433,12 @@
     Call :meth:`release` when done with the kernel instance.
     """
-    #: SAS model information structure
-    info = None # type: ModelInfo
-    #: kernel precision
-    dtype = None # type: np.dtype
-    #: kernel dimensions (1d or 2d)
-    dim = "" # type: str
-    #: calculation results, updated after each call to :meth:`_call_kernel`
-    result = None # type: np.ndarray
+    #: SAS model information structure.
+    info = None  # type: ModelInfo
+    #: Kernel precision.
+    dtype = None  # type: np.dtype
+    #: Kernel dimensions (1d or 2d).
+    dim = ""  # type: str
+    #: Calculation results, updated after each call to :meth:`_call_kernel`.
+    result = None  # type: np.ndarray
 
     def __init__(self, model, q_vectors):
@@ -443,39 +447,49 @@
         self.q_input = GpuInput(q_vectors, dtype)
         self._model = model
-        # F16 isn't sufficient, so don't support it
-        self._as_dtype = np.float64 if dtype == generate.F64 else np.float32
-
-        # attributes accessed from the outside
+
+        # Attributes accessed from the outside.
         self.dim = '2d' if self.q_input.is_2d else '1d'
         self.info = model.info
-        self.dtype = model.dtype
-
-        # holding place for the returned value
+        self.dtype = dtype
+
+        # Converter to translate input to target type.
+        self._as_dtype = np.float64 if dtype == generate.F64 else np.float32
+
+        # Holding place for the returned value.
         nout = 2 if self.info.have_Fq and self.dim == '1d' else 1
-        extra_q = 4  # total weight, form volume, shell volume and R_eff
-        self.result = np.empty(self.q_input.nq*nout+extra_q, dtype)
-
-        # allocate result value on gpu
+        extra_q = 4  # Total weight, form volume, shell volume and R_eff.
+        self.result = np.empty(self.q_input.nq*nout + extra_q, dtype)
+
+        # Allocate result value on GPU.
         width = ((self.result.size+31)//32)*32 * self.dtype.itemsize
         self._result_b = cuda.mem_alloc(width)
 
-    def _call_kernel(self, call_details, values, cutoff, magnetic, effective_radius_type):
-        # type: (CallDetails, np.ndarray, np.ndarray, float, bool) -> np.ndarray
-        # Arrange data transfer to card
+    def _call_kernel(self, call_details, values, cutoff, magnetic,
+                     effective_radius_type):
+        # type: (CallDetails, np.ndarray, float, bool, int) -> np.ndarray
+
+        # Arrange data transfer to card.
         details_b = cuda.to_device(call_details.buffer)
         values_b = cuda.to_device(values)
 
+        # Setup kernel function and arguments.
         name = 'Iq' if self.dim == '1d' else 'Imagnetic' if magnetic else 'Iqxy'
         kernel = self._model.get_function(name)
         kernel_args = [
-            np.uint32(self.q_input.nq), None, None,
-            details_b, values_b, self.q_input.q_b, self._result_b,
-            self._as_dtype(cutoff),
-            np.uint32(effective_radius_type),
+            np.uint32(self.q_input.nq),  # Number of inputs.
+            None,  # Placeholder for pd_start.
+            None,  # Placeholder for pd_stop.
+            details_b,  # Problem definition.
+            values_b,  # Parameter values.
+            self.q_input.q_b,  # Q values.
+            self._result_b,   # Result storage.
+            self._as_dtype(cutoff),  # Probability cutoff.
+            np.uint32(effective_radius_type),  # R_eff mode.
         ]
         grid = partition(self.q_input.nq)
-        #print("Calling OpenCL")
+
+        # Call kernel and retrieve results.
+        #print("Calling CUDA")
         #call_details.show(values)
-        # Call kernel and retrieve results
         last_nap = time.clock()
         step = 100000000//self.q_input.nq + 1
@@ -488,5 +502,5 @@
             if stop < call_details.num_eval:
                 sync()
-                # Allow other processes to run
+                # Allow other processes to run.
                 current_time = time.clock()
                 if current_time - last_nap > 0.5:
@@ -522,6 +536,5 @@
     Note: Maybe context.synchronize() is sufficient.
     """
-    #return # The following works in C++; don't know what pycuda is doing
-    # Create an event with which to synchronize
+    # Create an event with which to synchronize.
     done = cuda.Event()
 
@@ -529,5 +542,5 @@
     done.record()
 
-    #line added to not hog resources
+    # Make sure we don't hog resource while waiting to sync.
     while not done.query():
         time.sleep(0.01)
@@ -535,4 +548,5 @@
     # Block until the GPU executes the kernel.
     done.synchronize()
+
     # Clean up the event; I don't think they can be reused.
     del done
