Main Content


Pragma to allocate a variable as persistent memory on the GPU

Since R2020b



    coder.gpu.persistentMemory(pvar) maps the persistent MATLAB® variable pvar to the CUDA® enabled NVIDIA® GPU as variable with persistent memory. The variable must be fixed size and of a data type supported for GPU code generation.

    This is a code generation function. It has no effect in MATLAB.


    For CUDA MEX, the persistent memory on the GPU is retained during the entirety of the MATLAB session. For freeing the GPU memory, use the clear mex MATLAB command. For freeing the GPU memory in static libraries, dynamic libraries, or executable targets, call the generated <primary function name>_terminate() housekeeping function.


    collapse all

    This example shows how to map a persistent MATLAB variable to the persistent memory on the GPU by using the coder.gpu.persistentMemory pragma.

    Consider the following MATLAB entry-point function myPersistent containing a simple loop.

    function output = myPersistent(input)
        persistent pvar;
        if isempty(pvar)
            pvar = ones(size(input));
        for i = 1:numel(input)
            pvar(i) = pvar(i) + input(i);
        output = coder.nullcopy(input);
        for i = 1:numel(input)
            output(i) = pvar(i) * input(i);

    Create a code generation configuration object for a standalone CUDA static library.

    cfg = coder.gpuConfig('lib');

    Define a cell array input that declares the size and data type of the inputs to the function myPersistent.

    input = {rand(1,1024)}

    Generate a MEX function myPersistent_mex by using codegen function with the -config, -args, and -report options to specify configuration, provide input arguments, and generate a code generation report.

    codegen -config cfg -args input -report myPersistent

    A snippet of the generated file is shown.

    // File:
    // GPU Coder version                    : 2.0
    // CUDA/C/C++ source code generated on  : 16-Jul-2020 20:08:46
    // Include Files
    #include "myPersistent.h"
    #include "myPersistent_data.h"
    #include "myPersistent_initialize.h"
    #include "MWCudaDimUtility.hpp"
    // Arguments    : const double input[1024]
    //                double output[1024]
    // Return Type  : void
    void myPersistent(const double input[1024], double output[1024])
      double (*gpu_input)[1024];
      double (*gpu_output)[1024];
      if (!isInitialized_myPersistent) {
      cudaMalloc(&gpu_output, 8192UL);
      cudaMalloc(&gpu_input, 8192UL);
      cudaMemcpy(gpu_input, (void *)&input[0], 8192UL, cudaMemcpyHostToDevice);
      myPersistent_kernel1<<<dim3(2U, 1U, 1U), dim3(512U, 1U, 1U)>>>(*gpu_input,
        *gpu_output, *pvar);
      cudaMemcpy(&output[0], gpu_output, 8192UL, cudaMemcpyDeviceToHost);
    // Arguments    : void
    // Return Type  : void
    void myPersistent_init()
      double b_pvar[1024];
      boolean_T pvar_dirtyOnCpu;
      boolean_T pvar_dirtyOnGpu;
      pvar_dirtyOnCpu = false;
      pvar_dirtyOnGpu = true;
      for (int i = 0; i < 1024; i++) {
        if (pvar_dirtyOnGpu) {
          cudaMemcpy(&b_pvar[0], pvar, 8192UL, cudaMemcpyDeviceToHost);
          pvar_dirtyOnGpu = false;
        b_pvar[i] = 1.0;
        pvar_dirtyOnCpu = true;
      if (pvar_dirtyOnCpu) {
        cudaMemcpy(pvar, &b_pvar[0], 8192UL, cudaMemcpyHostToDevice);

    The persistent variable pvar is retained on the GPU between calls to the myPersistent() function.

    Input Arguments

    collapse all

    The name of the variable that must be mapped to the GPU memory space as a persistent variable.


    • MATLAB classes are not supported by the coder.gpu.persistentMemory pragma.

    • The coder.gpu.persistentMemory pragma is not supported in Simulink®.

    • If the persistent variable in MATLAB is variable-sized, the code generator does not map the variable to a persistent memory on the GPU.

    Version History

    Introduced in R2020b