coder.gpu.persistentMemory

Pragma to allocate a variable as persistent memory on the GPU

Since R2020b

Syntax

coder.gpu.persistentMemory(pvar)

Description

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.

Note

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.

Examples

Using coder.gpu.persistentMemory to allocate a variable on the GPU memory as persistent

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));
    end
    
    coder.gpu.persistentMemory(pvar);
    
    for i = 1:numel(input)
        pvar(i) = pvar(i) + input(i);
    end
    
    output = coder.nullcopy(input);
    for i = 1:numel(input)
        output(i) = pvar(i) * input(i);
    end
end

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 myPersistent.cu file is shown.

//
// File: myPersistent.cu
//
// 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) {
    myPersistent_initialize();
  }

  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);
  cudaFree(*gpu_input);
  cudaFree(*gpu_output);
}

//
// 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

pvar — Name of the persistent variable

name

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

Limitations

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

• 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

R2024b: Use GPU persistent memory in Simulink models

The coder.gpu.persistentMemory pragma is supported for Simulink^®. You can observe improved performance in Simulink simulations and generate optimized code for Simulink models.