A program CUDA-t használó része hasonlítani fog a korábban látotthoz.

A belépési pontunk itt is az alábbi függvény:

void mexFunction(int nlhs, mxArray *plhs[],
                 int nrhs, mxArray const *prhs[])

A szerkezet nem sokat változott.

Különbség talán az, hogy inicializálni kell a GPU-t.

    mxInitGPU();

Ezek után megint csak az input-ok megfelelőségét ellenőrizzük.

    //  Check if parameters are 2 GPU Arrays.
    if ((nrhs!=2) || !(mxIsGPUArray(prhs[0])) || !(mxIsGPUArray(prhs[1]))) {
        mexErrMsgIdAndTxt(errId, errMsg);
    }

Ezek után viszont már más a GPU kezelése. A számítás előtt először a bemeneteket át kell konvertálni, hogy mxGPUArray típusúak legyenek.

    mxGPUArray const *A;
    mxGPUArray const *B;
    mxGPUArray *C;

    A = mxGPUCreateFromMxArray(prhs[0]);
    B = mxGPUCreateFromMxArray(prhs[1]);

Majd jöhet még egy kis típus ellenzőrzés. (Ezt megoldhatnánk típus konverzióval is, hogy a programunk minden formátumú adatra fusson, de ez itt most egy oktató program ami csak a lényeget mutatja.)

    //  Check if input are of type double
    if ((mxGPUGetClassID(A) != mxSINGLE_CLASS) || (mxGPUGetClassID(B) != mxSINGLE_CLASS)) {
        mexErrMsgIdAndTxt(errId, errMsg);
    }

Végül pedig a számítás előtt el kell kérnünk az adatokra mutató pointereket.

    d_A = (float const *)(mxGPUGetDataReadOnly(A));
    d_B = (float const *)(mxGPUGetDataReadOnly(B));

Ezen a ponton van két grafikus memóriában tárolt tömbre mutató pointerünk. Innetnől gyakorlatilag csak c++ kódot írunk. Vagyis majdnem, mert még először elő kell készíteni a kimenet tömbjét is.

    C = mxGPUCreateGPUArray(mxGPUGetNumberOfDimensions(A),
                            mxGPUGetDimensions(A),
                            mxGPUGetClassID(A),
                            mxGPUGetComplexity(A),
                            MX_GPU_DO_NOT_INITIALIZE);
    d_C = (float *)(mxGPUGetData(C));

De ettől a ponttól tényleg standard CUDA programozás jön. Előkészítünk a blokk struktúrát, és elindítjuk as zámító kernelt.

    N = (int)(mxGPUGetNumberOfElements(A));
    gridDim = (N + BLOCKDIM - 1) / BLOCKDIM;
    TimesTwo<<<gridDim, BLOCKDIM>>>(d_A, d_B, d_C, N);

A kernelt megnézve, itt már nem találun ksemmi meglepőt.

void __global__ TimesTwo(float const * const A,
                         float const * const B,
                         float * const C,
                         int const N)
{
    int const tid = blockDim.x * blockIdx.x + threadIdx.x;
    if (tid < N) {
        C[tid] = A[tid] + B[tid];
    }
}

Végül pedig az eredményt visszatesszük a kimeneti változóba.

    plhs[0] = mxGPUCreateMxArrayOnGPU(C);

És jó programozó módjára kitakarítunk magunk után.

    mxGPUDestroyGPUArray(A);
    mxGPUDestroyGPUArray(B);
    mxGPUDestroyGPUArray(C);

A fordítás menete nagyon is hansolóan működik, mint a CPU-n ment. Egy függvényhívás és kész is. Csak arra kell fögyelnünk, hogy a mexCuda parancsot használjuk, ami kezeli a GPU kódot.

mexcuda mexCudaAddSingle.cu;

Végül pedig meghívjuk a függvényünket ügyelve arra, hogy olyan adattípust adjunk neki, amit kezel.

gpu_A = gpuArray(single(A));
gpu_B = gpuArray(single(B));

gpu_C = mexCudaAddSingle(gpu_A, gpu_B)

#include "mex.h"
#include "gpu/mxGPUArray.h"

#define BLOCKDIM    256

void __global__ TimesTwo(float const * const A,
                         float const * const B,
                         float * const C,
                         int const N)
{
    int const tid = blockDim.x * blockIdx.x + threadIdx.x;
    if (tid < N) {
        C[tid] = A[tid] + B[tid];
    }
}

void mexFunction(int nlhs, mxArray *plhs[],
                 int nrhs, mxArray const *prhs[])
{
    mxGPUArray const *A;
    mxGPUArray const *B;
    mxGPUArray *C;
    float const *d_A;
    float const *d_B;
    float *d_C;
    int N;
    int gridDim;
    char const * const errId = "parallel:gpu:mexGPUExample:InvalidInput";
    char const * const errMsg = "Invalid input to MEX file.";

    mxInitGPU();

    //  Check if parameters are 2 GPU Arrays.
    if ((nrhs!=2) || !(mxIsGPUArray(prhs[0])) || !(mxIsGPUArray(prhs[1]))) {
        mexErrMsgIdAndTxt(errId, errMsg);
    }

    A = mxGPUCreateFromMxArray(prhs[0]);
    B = mxGPUCreateFromMxArray(prhs[1]);

    //  Check if input are of type double
    if ((mxGPUGetClassID(A) != mxSINGLE_CLASS) || (mxGPUGetClassID(B) != mxSINGLE_CLASS)) {
        mexErrMsgIdAndTxt(errId, errMsg);
    }

    d_A = (float const *)(mxGPUGetDataReadOnly(A));
    d_B = (float const *)(mxGPUGetDataReadOnly(B));

    C = mxGPUCreateGPUArray(mxGPUGetNumberOfDimensions(A),
                            mxGPUGetDimensions(A),
                            mxGPUGetClassID(A),
                            mxGPUGetComplexity(A),
                            MX_GPU_DO_NOT_INITIALIZE);
    d_C = (float *)(mxGPUGetData(C));

    N = (int)(mxGPUGetNumberOfElements(A));
    gridDim = (N + BLOCKDIM - 1) / BLOCKDIM;
    TimesTwo<<<gridDim, BLOCKDIM>>>(d_A, d_B, d_C, N);

    plhs[0] = mxGPUCreateMxArrayOnGPU(C);

    mxGPUDestroyGPUArray(A);
    mxGPUDestroyGPUArray(B);
    mxGPUDestroyGPUArray(C);
}