Documentation

Compile a PTX File from a CU File

If you have a CU file you want to execute on the GPU, you must first compile it to create a PTX file. One way to do this is with the nvcc compiler in the NVIDIA^® CUDA^® Toolkit. For example, if your CU file is called myfun.cu, you can create a compiled PTX file with the shell command:

nvcc -ptx myfun.cu

This generates the file named myfun.ptx.

Construct CUDAKernel Object with CU File Input

With a .cu file and a .ptx file you can create a CUDAKernel object in MATLAB that you can then use to evaluate the kernel:

k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu');

Construct CUDAKernel Object with C Prototype Input

If you do not have the CU file corresponding to your PTX file, you can specify the C prototype for your C kernel instead of the CU file. For example:

k = parallel.gpu.CUDAKernel('myfun.ptx','float *, const float *, float');

Another use for C prototype input is when your source code uses an unrecognized renaming of a supported data type. (See the supported types below.) Suppose your kernel comprises the following code.

typedef float ArgType;
__global__ void add3( ArgType * v1, const ArgType * v2 )
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];
}

ArgType itself is not recognized as a supported data type, so the CU file that includes it cannot be directly used as input when creating the CUDAKernel object in MATLAB. However, the supported input types to the add3 kernel can be specified as C prototype input to the CUDAKernel constructor. For example:

k = parallel.gpu.CUDAKernel('test.ptx','float *, const float *','add3');

Supported Data Types

The supported C/C++ standard data types are listed in the following table.

Also, the following integer types are supported when you include the tmwtypes.h header file in your program.

The header file is shipped as matlabroot/extern/include/tmwtypes.h. You include the file in your program with the line:

#include "tmwtypes.h"

Argument Restrictions

All inputs can be scalars or pointers, and can be labeled const.

The C declaration of a kernel is always of the form:

__global__ void aKernel(inputs ...)

When translating the definition of a kernel in C into MATLAB:

These rules have some implications. The most notable is that every output from a kernel must necessarily also be an input to the kernel, since the input allows the user to define the size of the output (which follows from being unable to allocate memory on the GPU).

CUDAKernel Object Properties

When you create a kernel object without a terminating semicolon, or when you type the object variable at the command line, MATLAB displays the kernel object properties. For example:

k = parallel.gpu.CUDAKernel('conv.ptx','conv.cu')

k = 
  parallel.gpu.CUDAKernel handle
  Package: parallel.gpu

  Properties:
     ThreadBlockSize: [1 1 1]
  MaxThreadsPerBlock: 512
            GridSize: [1 1 1]
    SharedMemorySize: 0
          EntryPoint: '_Z8theEntryPf'
  MaxNumLHSArguments: 1
     NumRHSArguments: 2
       ArgumentTypes: {'in single vector'  'inout single vector'}

The properties of a kernel object control some of its execution behavior. Use dot notation to alter those properties that can be changed.

For a descriptions of the object properties, see the CUDAKernel object reference page. A typical reason for modifying the settable properties is to specify the number of threads, as described below.

Specify Entry Points

If your PTX file contains multiple entry points, you can identify the particular kernel in myfun.ptx that you want the kernel object k to refer to:

k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu','myKernel1');

A single PTX file can contain multiple entry points to different kernels. Each of these entry points has a unique name. These names are generally mangled (as in C++ mangling). However, when generated by nvcc the PTX name always contains the original function name from the CU file. For example, if the CU file defines the kernel function as

__global__ void simplestKernelEver( float * x, float val )

then the PTX code contains an entry that might be called _Z18simplestKernelEverPff.

When you have multiple entry points, specify the entry name for the particular kernel when calling CUDAKernel to generate your kernel.

You might not have control over the original entry names, in which case you must be aware of the unique mangled derived for each. For example, consider the following function template.

template <typename T>
__global__ void add4( T * v1, const T * v2 )
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];
}

When the template is expanded out for float and double, it results in two entry points, both of which contain the substring add4.

template __global__ void add4<float>(float *, const float *);
template __global__ void add4<double>(double *, const double *);

The PTX has corresponding entries:

_Z4add4IfEvPT_PKS0_
_Z4add4IdEvPT_PKS0_

Use entry point add4If for the float version, and add4Id for the double version.

k = parallel.gpu.CUDAKernel('test.ptx','double *, const double *','add4Id');

Specify Number of Threads

You specify the number of computational threads for your CUDAKernel by setting two of its object properties:

The default value for both of these properties is [1 1 1], but suppose you want to use 500 threads to run element-wise operations on vectors of 500 elements in parallel. A simple way to achieve this is to create your CUDAKernel and set its properties accordingly:

k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu');
k.ThreadBlockSize = [500,1,1];

Generally, you set the grid and thread block sizes based on the sizes of your inputs. For information on thread hierarchy, and multiple-dimension grids and blocks, see the NVIDIA CUDA C Programming Guide.

Use Workspace Variables

Assume that you have already written some kernels in a native language and want to use them in MATLAB to execute on the GPU. You have a kernel that does a convolution on two vectors; load and run it with two random input vectors:

k = parallel.gpu.CUDAKernel('conv.ptx','conv.cu');

result = feval(k,rand(100,1),rand(100,1));

Even if the inputs are constants or variables for MATLAB workspace data, the output is gpuArray.

Use gpuArray Variables

It might be more efficient to use gpuArray objects as input when running a kernel:

k = parallel.gpu.CUDAKernel('conv.ptx','conv.cu');

i1 = gpuArray(rand(100,1,'single'));
i2 = gpuArray(rand(100,1,'single'));

result1 = feval(k,i1,i2);

Because the output is a gpuArray, you can now perform other operations using this input or output data without further transfers between the MATLAB workspace and the GPU. When all your GPU computations are complete, gather your final result data into the MATLAB workspace:

result2 = feval(k,i1,i2);

r1 = gather(result1);
r2 = gather(result2);

Determine Input and Output Correspondence

When calling [out1, out2] = feval(kernel, in1, in2, in3), the inputs in1, in2, and in3 correspond to each of the input arguments to the C function within your CU file. The outputs out1 and out2 store the values of the first and second non-const pointer input arguments to the C function after the C kernel has been executed.

For example, if the C kernel within a CU file has the following signature:

void reallySimple( float * pInOut, float c )

the corresponding kernel object (k) in MATLAB has the following properties:

MaxNumLHSArguments: 1
   NumRHSArguments: 2
     ArgumentTypes: {'inout single vector'  'in single scalar'}

Therefore, to use the kernel object from this code with feval, you need to provide feval two input arguments (in addition to the kernel object), and you can use one output argument:

y = feval(k,x1,x2)

The input values x1 and x2 correspond to pInOut and c in the C function prototype. The output argument y corresponds to the value of pInOut in the C function prototype after the C kernel has executed.

The following is a slightly more complicated example that shows a combination of const and non-const pointers:

void moreComplicated( const float * pIn, float * pInOut1, float * pInOut2 )

The corresponding kernel object in MATLAB then has the properties:

MaxNumLHSArguments: 2
   NumRHSArguments: 3
     ArgumentTypes: {'in single vector'  'inout single vector'  'inout single vector'}

You can use feval on this code’s kernel (k) with the syntax:

[y1,y2] = feval(k,x1,x2,x3)

The three input arguments x1, x2, and x3, correspond to the three arguments that are passed into the C function. The output arguments y1 and y2, correspond to the values of pInOut1 and pInOut2 after the C kernel has executed.

Add Two Numbers

This example adds two doubles together in the GPU. You should have the NVIDIA CUDA Toolkit installed, and have CUDA-capable drivers for your device.

Add Two Vectors

This example extends the previous one to add two vectors together. For simplicity, assume that there are exactly the same number of threads as elements in the vectors and that there is only one thread block.

Example with CU and PTX Files

For an example that shows how to work with CUDA, and provides CU and PTX files for you to experiment with, see Illustrating Three Approaches to GPU Computing: The Mandelbrot Set.