# 19. CUDA model evaluator

## `sasmodels.kernelcuda`

GPU driver for C kernels (with CUDA)

To select cuda, use SAS_OPENCL=cuda, or SAS_OPENCL=cuda:n for a particular device number. If no device number is specified, then look for CUDA_DEVICE=n or a file ~/.cuda-device containing n for the device number. Otherwise, try all available device numbers.

TODO: docs are out of date

There should be a single GPU environment running on the system. This environment is constructed on the first call to `environment()`, and the same environment is returned on each call.

After retrieving the environment, the next step is to create the kernel. This is done with a call to `GpuEnvironment.compile_program()`, which returns the type of data used by the kernel.

Next a `GpuInput` object should be created with the correct kind of data. This data object can be used by multiple kernels, for example, if the target model is a weighted sum of multiple kernels. The data should include any extra evaluation points required to compute the proper data smearing. This need not match the square grid for 2D data if there is an index saying which q points are active.

Together the GpuInput, the program, and a device form a `GpuKernel`. This kernel is used during fitting, receiving new sets of parameters and evaluating them. The output value is stored in an output buffer on the devices, where it can be combined with other structure factors and form factors and have instrumental resolution effects applied.

In order to use OpenCL for your models, you will need OpenCL drivers for your machine. These should be available from your graphics card vendor. Intel provides OpenCL drivers for CPUs as well as their integrated HD graphics chipsets. AMD also provides drivers for Intel CPUs, but as of this writing the performance is lacking compared to the Intel drivers. NVidia combines drivers for CUDA and OpenCL in one package. The result is a bit messy if you have multiple drivers installed. You can see which drivers are available by starting python and running:

> import pyopencl as cl cl.create_some_context(interactive=True)

Once you have done that, it will show the available drivers which you can select. It will then tell you that you can use these drivers automatically by setting the SAS_OPENCL environment variable, which is PYOPENCL_CTX equivalent but not conflicting with other pyopnecl programs.

Some graphics cards have multiple devices on the same card. You cannot yet use both of them concurrently to evaluate models, but you can run the program twice using a different device for each session.

OpenCL kernels are compiled when needed by the device driver. Some drivers produce compiler output even when there is no error. You can see the output by setting PYOPENCL_COMPILER_OUTPUT=1. It should be harmless, albeit annoying.

class sasmodels.kernelcuda.GpuEnvironment(devnum: int = None)

  Bases: `object`

  GPU context for CUDA.

  compile_program(name: str, source: str, dtype: np.dtype, fast: bool, timestamp: float) → SourceModule

    Compile the program for the device in the given context.

  context: cuda.Context = None
  has_type(dtype: dtype) → bool

    Return True if all devices support a given type.

  release()

    Free the CUDA device associated with this context.

class sasmodels.kernelcuda.GpuInput(q_vectors: list[ndarray], dtype: dtype = dtype('float32'))

  Bases: `object`

  Make q data available to the gpu.

  q_vectors is a list of q vectors, which will be [q] for 1-D data, and [qx, qy] for 2-D data. Internally, the vectors will be reallocated to get the best performance on OpenCL, which may involve shifting and stretching the array to better match the memory architecture. Additional points will be evaluated with q=1e-3.

  dtype is the data type for the q vectors. The data type should be set to match that of the kernel, which is an attribute of `GpuModel`. Note that not all kernels support double precision, so even if the program was created for double precision, the GpuModel.dtype may be single precision.

  Call `release()` when complete. Even if not called directly, the buffer will be released when the data object is freed.

  release() → None

    Free the buffer associated with the q value.

class sasmodels.kernelcuda.GpuKernel(model: GpuModel, q_vectors: list[ndarray])

  Bases: `Kernel`

  Callable SAS kernel.

  model is the GpuModel object to call

  The kernel is derived from `kernel.Kernel`, providing the _call_kernel() method to evaluate the kernel for a given set of parameters. Because of the need to move the q values to the GPU before evaluation, the kernel is instantiated for a particular set of q vectors, and can be called many times without transfering q each time.

  Call `release()` when done with the kernel instance.

  Fq(call_details: CallDetails, values: np.ndarray, cutoff: np.ndarray, magnetic: float, radius_effective_mode: bool = 0) → np.ndarray

    Returns <F(q)>, <F(q)^2>, effective radius, shell volume and form:shell volume ratio. The <F(q)> term may be None if the form factor does not support direct computation of \(F(q)\)

    \(P(q) = <F^2(q)>/<V>\) is used for structure factor calculations,

    \[I(q) = \text{scale} \cdot P(q) \cdot S(q) + \text{background}\]

    For the beta approximation, this becomes

    \[I(q) = \text{scale} P (1 + <F>^2/<F^2> (S - 1)) + \text{background} = \text{scale}/<V> (<F^2> + <F>^2 (S - 1)) + \text{background}\]

    \(<F(q)>\) and \(<F^2(q)>\) are averaged by polydispersity in shape and orientation, with each configuration \(x_k\) having form factor \(F(q, x_k)\), weight \(w_k\) and volume \(V_k\). The result is:

    \[P(q)=\frac{\sum w_k F^2(q, x_k) / \sum w_k}{\sum w_k V_k / \sum w_k}\]

    The form factor itself is scaled by volume and contrast to compute the total scattering. This is then squared, and the volume weighted F^2 is then normalized by volume F. For a given density, the number of scattering centers is assumed to scale linearly with volume. Later scaling the resulting \(P(q)\) by the volume fraction of particles gives the total scattering on an absolute scale. Most models incorporate the volume fraction into the overall scale parameter. An exception is vesicle, which includes the volume fraction parameter in the model itself, scaling \(F\) by \(\surd V_f\) so that the math for the beta approximation works out.

    By scaling \(P(q)\) by total weight \(\sum w_k\), there is no need to make sure that the polydisperisity distributions normalize to one. In particular, any distibution values \(x_k\) outside the valid domain of \(F\) will not be included, and the distribution will be implicitly truncated. This is controlled by the parameter limits defined in the model (which truncate the distribution before calling the kernel) as well as any region excluded using the INVALID macro defined within the model itself.

    The volume used in the polydispersity calculation is the form volume for solid objects or the shell volume for hollow objects. Shell volume should be used within \(F\) so that the normalizing scale represents the volume fraction of the shell rather than the entire form. This corresponds to the volume fraction of shell-forming material added to the solvent.

    The calculation of \(S\) requires the effective radius and the volume fraction of the particles. The model can have several different ways to compute effective radius, with the radius_effective_mode parameter used to select amongst them. The volume fraction of particles should be determined from the total volume fraction of the form, not just the shell volume fraction. This makes a difference for hollow shapes, which need to scale the volume fraction by the returned volume ratio when computing \(S\). For solid objects, the shell volume is set to the form volume so this scale factor evaluates to one and so can be used for both hollow and solid shapes.

  Iq(call_details: CallDetails, values: np.ndarray, cutoff: np.ndarray, magnetic: float) → np.ndarray

    Returns I(q) from the polydisperse average scattering.

    \[I(q) = \text{scale} \cdot P(q) + \text{background}\]

    With the correct choice of model and contrast, setting scale to the volume fraction \(V_f\) of particles should match the measured absolute scattering. Some models (e.g., vesicle) have volume fraction built into the model, and do not need an additional scale.

  dim: str = ''

    Kernel dimensions (1d or 2d).

  dtype: np.dtype = None

    Kernel precision.

  info: ModelInfo = None

    SAS model information structure.

  q_input: Any = None

    Q values at which the kernel is to be evaluated.

  release() → None

    Release resources associated with the kernel.

  result: np.ndarray = None

    Calculation results, updated after each call to _call_kernel().

class sasmodels.kernelcuda.GpuModel(source: dict[str, str], model_info: ModelInfo, dtype: dtype = dtype('float32'), fast: bool = False)

  Bases: `KernelModel`

  GPU wrapper for a single model.

  source and model_info are the model source and interface as returned from `generate.make_source()` and `modelinfo.make_model_info()`.

  dtype is the desired model precision. Any numpy dtype for single or double precision floats will do, such as ‘f’, ‘float32’ or ‘single’ for single and ‘d’, ‘float64’ or ‘double’ for double. Double precision is an optional extension which may not be available on all devices. Half precision (‘float16’,’half’) may be available on some devices. Fast precision (‘fast’) is a loose version of single precision, indicating that the compiler is allowed to take shortcuts.

  dtype: np.dtype = None
  fast: bool = False
  get_function(name: str) → cuda.Function

    Fetch the kernel from the environment by name, compiling it if it does not already exist.

  info: ModelInfo = None
  make_kernel(q_vectors: list[ndarray]) → GpuKernel

    Instantiate a kernel for evaluating the model at q_vectors.

  release() → None

    Free resources associated with the kernel.

  source: str = ''

sasmodels.kernelcuda.compile_model(source: str, dtype: np.dtype, fast: bool = False) → SourceModule

  Build a model to run on the gpu.

  Returns the compiled program and its type. The returned type will be float32 even if the desired type is float64 if any of the devices in the context do not support the cl_khr_fp64 extension.

sasmodels.kernelcuda.environment() → GpuEnvironment

  Returns a singleton `GpuEnvironment`.

  This provides an OpenCL context and one queue per device.

sasmodels.kernelcuda.has_type(dtype: dtype) → bool

  Return true if device supports the requested precision.

sasmodels.kernelcuda.mark_device_functions(source: str) → str

  Mark all function declarations as __device__ functions (except kernel).

sasmodels.kernelcuda.partition(n)

  Constructs block and grid arguments for n elements.

sasmodels.kernelcuda.reset_environment() → None

  Call to create a new OpenCL context, such as after a change to SAS_OPENCL.

sasmodels.kernelcuda.show_device_functions(source: str) → str

  Show all discovered function declarations, but don’t change any.

sasmodels.kernelcuda.sync()
  Overview:

    Waits for operation in the current context to complete.

  Note: Maybe context.synchronize() is sufficient.

sasmodels.kernelcuda.use_cuda() → bool

  Returns True if CUDA is the default compute engine.