This file is indexed.

/usr/share/psychtoolbox-3/PsychGPGPU/memcpyCudaOpenGL.cpp is in psychtoolbox-3-common 3.0.14.20170103+git6-g605ff5c.dfsg1-1build1.

This file is owned by root:root, with mode 0o644.

The actual contents of the file can be viewed below.

  1
  2
  3
  4
  5
  6
  7
  8
  9
 10
 11
 12
 13
 14
 15
 16
 17
 18
 19
 20
 21
 22
 23
 24
 25
 26
 27
 28
 29
 30
 31
 32
 33
 34
 35
 36
 37
 38
 39
 40
 41
 42
 43
 44
 45
 46
 47
 48
 49
 50
 51
 52
 53
 54
 55
 56
 57
 58
 59
 60
 61
 62
 63
 64
 65
 66
 67
 68
 69
 70
 71
 72
 73
 74
 75
 76
 77
 78
 79
 80
 81
 82
 83
 84
 85
 86
 87
 88
 89
 90
 91
 92
 93
 94
 95
 96
 97
 98
 99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
/*
 * memcpyCudaOpenGL.cpp -- Fast memory copy between CUDA and OpenGL resources.
 *
 * This function, compiled into a MEX file for Mathworks Matlab or GNU/Octave
 * allows to exchange data between NVidia CUDA memory buffers and other resources
 * and OpenGL resources, such as textures, renderbuffers, framebuffer objects,
 * and pixelbuffer objects. Data exchange can happen via copying data between
 * OpenGL memory and CUDA memory, e.g., for use with high-level GPGPU toolkits,
 * or via shared access to the same underlying OpenGL resource, e.g., for even
 * more efficient use with low-level GPGPU toolkits.
 *
 * To avoid redundant calls to expensive high-overhead CUDA functions like
 * CUDA-OpenGL resource (un)registration and resource (un)mapping, it uses a
 * LRU cache with a default size of 8 slots, configurable to up to 1024 slots.
 * The cache stores registration and mapping info for all OpenGL resources which
 * are shared between CUDA and OpenGL and tries to keep resources registered as
 * long as possible and mapped as long as possible. If the cache reaches its
 * capacity, it uses a LRU strategy to decide which items to evict and thereby
 * unmap and unregister.
 *
 * Efficient use of the cache to minimize such operations is important, as they
 * can be very expensive. Example: CUDA-5.0 + GPUmat toolkit on 64-Bit Matlab under
 * 64-Bit Ubuntu Linux 12.04.1 LTS on a 8-core MacPro with a NVidia GeForce-8800:
 *
 * Time for 1024x768 pixel RGBA 32 bpc floating point data copies between CUDA
 * (GPUmat GPUsingle() matrix) and OpenGL texture (PTB texture of equivalent size):
 * Pure copy: < 0.5 msecs. Corresponds to about 97% of GPU peak bandwidth.
 * With mapping and unmapping: ~ 1.0 msecs.
 * With registering and unregistering of resources: ~ 4.9 msecs!
 * Non-Cached OpenGL interop can cause a 10x performance hit just for communication,
 * so effective caching is important.
 *
 *
 * See accompanying M-File for building the mex file.
 * See M-Wrapper file for actually using this low-level function.
 *
 * This file was originally developed for Psychtoolbox-3 (www.psychtoolbox.org)
 * but can be used for interop between any Matlab or Octave based CUDA and
 * OpenGL toolkit.
 *
 *
 * Copyright (c) 2013 Mario Kleiner. Licensed to you under the MIT license:
 *
 * Permission is hereby granted, free of charge, to any person obtaining a
 * copy of this software and associated documentation files (the
 * "Software"), to deal in the Software without restriction, including
 * without limitation the rights to use, copy, modify, merge, publish,
 * distribute, sublicense, and/or sell copies of the Software, and to permit
 * persons to whom the Software is furnished to do so, subject to the
 * following conditions:
 * 
 * The above copyright notice and this permission notice shall be included
 * in all copies or substantial portions of the Software.
 * 
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
 * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
 * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN
 * NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
 * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
 * OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
 * USE OR OTHER DEALINGS IN THE SOFTWARE.
 */

/* MATLAB / OCTAVE includes: */
#include "mex.h"

#include "string.h"

#ifdef _WIN32
#include <windows.h>
#endif

/* CUDA includes: */
#include "cuda_runtime_api.h"
#include "cuda_gl_interop.h"

/* Global verbosity level and firsttime init flag: */
int verbose = 0;
static int firsttime = 1;

/* Global CUDA error state: */
cudaError_t cudastatus;

/* LRU resource registration/mapping cache: */
typedef struct cudares {
  unsigned int glhandle;    /* OpenGL object handle: Used as search key in the cache. 0 = Free slot. */
  unsigned int gltarget;    /* OpenGL object type: Texture, renderbuffer, pixelbuffer etc. */
  cudaGraphicsResource_t resource;  /* CUDA resource descriptor if OpenGL object registered with CUDA, NULL otherwise. */
  unsigned int ismapped;    /* 0 = Resource not mapped, 1 = Resource mapped. */
  unsigned int mapflags;    /* Resource mapping flags used. */
  unsigned long long lastaccess;  /* Time in "cache-ticks" (0 = Never) when this resource was last accessed: For LRU replacement. */
} cudares;

/* The resource mapping cache: Has MAX_CACHE_SLOTS slots maximum capacity. */
#define MAX_CACHE_SLOTS 1024
struct cudares resourceCache[MAX_CACHE_SLOTS];

/* Effective cache capacity: */
int cachesize;

/* First time cache full warning already emitted? */
int firstLRUCycle;

/* Current cache reference time for LRU replacement: */
unsigned long long cacheclock;

void doCudaUnmap(int i)
{
  /* Unmap if it is actually mapped */
  if (resourceCache[i].resource && resourceCache[i].ismapped) {
    /* Mark as unmapped */
    resourceCache[i].ismapped = 0;
    
	  cudastatus = cudaGraphicsUnmapResources(1, &(resourceCache[i].resource), 0);
	  if (cudastatus != cudaSuccess) {
	    mexPrintf("\nmemcpyCudaOpenGL: ERROR in cudaGraphicsUnmapResources(): %s\n", cudaGetErrorString(cudastatus));
	  }
	  
    if (verbose) mexPrintf("\n%s: doCudaUnmap(%i).\n", mexFunctionName(), i);	  
  }
}

void doCudaUnregister(int i)
{
  /* Unmap if it is mappped: */
  doCudaUnmap(i);
  
  /* Unregister if it is registered: */
  if (resourceCache[i].resource) {
    cudastatus = cudaGraphicsUnregisterResource(resourceCache[i].resource);
    if (cudastatus != cudaSuccess) {
      mexPrintf("\nmemcpyCudaOpenGL: ERROR in cudaGraphicsUnregisterResource(): %s\n", cudaGetErrorString(cudastatus));
    }
    
    if (verbose) mexPrintf("\n%s: doCudaUnregister(%i).\n", mexFunctionName(), i);    
  }
  
  /* Clear out and thereby free the cache slot: */
  memset(&(resourceCache[i]), 0, sizeof(resourceCache[i]));
}

/* Search for an item in the cache: */
int cacheLookup(unsigned int glhhandle, unsigned int gltarget)
{
  int i;

  /* Iterate over all cache slots and search for (glhandle,gltarget) key: */
  for (i = 1; i < cachesize; i++) {
    /* In use? */
    if ((resourceCache[i].glhandle == glhhandle) && (resourceCache[i].gltarget == gltarget)) {
      /* Cache hit: Update access timestamp and return. */
      resourceCache[i].lastaccess = cacheclock;
      return(i);
    }
  }

  /* Cache miss: */
  return(-1);
}

/* Enter an item into the cache, unless it is already there, in which case its index is returned: */
int cacheInsert(unsigned int glhhandle, unsigned int gltarget, unsigned int mapflags)
{
  int i, lruvictim;
  unsigned long long curage;
  unsigned long long maxage = 0;

  /* Iterate over all cache slots and search for (glhandle,gltarget) key:
     We omit slot 0 as it is reserved. Otherwise we couldn't signal caching state
     to caller via the sign of the return value, as zero is ambiguous.
   */
  for (i = 1; i < cachesize; i++) {
    /* Found? */
    if ((resourceCache[i].glhandle == glhhandle) && (resourceCache[i].gltarget == gltarget)) {
      /* Cache hit: Update access timestamp and return slot. */
      resourceCache[i].lastaccess = cacheclock;
      
      /* Well, almost. We need to make sure that this resource is also registered
         with compatible 'mapflags', or bad things would happen.
       */
      if (mapflags != resourceCache[i].mapflags) {
        /* Incompatible flags. Need to unregister the resource and then return
         * its former host slot 'i' as "free" slot, so calling code thinks the
         * resource is not in cache and reregisters it with proper mapflags.
         */
         if (verbose) mexPrintf("\n%s: cacheInsert(%i): Incompatible cached map flags [%i vs. %i]. Triggering reregister.\n", mexFunctionName(), i, mapflags, resourceCache[i].mapflags);
         doCudaUnregister(i);
         return(-i);
      }
     
      /* Registered in a directly useable way: */ 
      return(i);
    }

    /* Nope. Keep track of first free slot, or slot which hasn't been accessed for
       the longest time, ie., the victim for lru replacement if neccessary. */
    curage = cacheclock - resourceCache[i].lastaccess;
    if (curage > maxage) {
      maxage = curage;
      lruvictim = i;
    }
  }

  /* If we reach this point, then the object isn't in the cache yet and needs
     to be entered into a free cache slot, if any. lruvictim is our candidate
     slot. If its lastaccess timestamp is zero, then we got a free slot for use.
     Otherwise our cache is full and lruvictim is the slot that wasn't used for
     the longest time which makes it our victim for expelling it from the cache.
   */
  if (resourceCache[lruvictim].lastaccess > 0) {
    /* Occupied slot. Free it up by unmapping and unregistering the cached
       OpenGL resource.
     */
    doCudaUnregister(lruvictim);
    if (firstLRUCycle) {
      firstLRUCycle = 0;
      mexPrintf("%s: Warning: CUDA<->OpenGL interop cache soft-limit of %i slots reached. Cache full!\n", mexFunctionName(), cachesize - 1);
      mexPrintf("%s: Warning: Will switch to LRU cache replacement mode now. Performance may be degraded!\n", mexFunctionName());
      mexPrintf("%s: Warning: Increase the cache soft-limit, or check your code for unintended resource leaks.\n", mexFunctionName());
    }

    if (verbose) mexPrintf("\n%s: cacheInsert(%i): LRU-Replacement done (Victims age %i).\n", mexFunctionName(), lruvictim, maxage);
  }
  else if (verbose) mexPrintf("\n%s: cacheInsert(%i): Inserted into free slot.\n", mexFunctionName(), lruvictim);

  /* Cache miss: Return the (now) free cache slot 'lruvictim' to be filled. Return
     a negative index to indicate this was a cache-miss and the cache needs to be filled.
   */
  return(-lruvictim);
}

/* Flush the resourceCache - Unmap and unregister all cached resources: */
void cacheFlush(void)
{
  int i;

  if (verbose) mexPrintf("\n%s: cacheFlush.\n", mexFunctionName());

  /* Iterate over all cache slots and unmap/unregister their resources: */
  for (i = 1; i < cachesize; i++) {
    /* In use? */
    if (resourceCache[i].glhandle) {
      doCudaUnregister(i);
    }
  }

  /* Nothing valid in cache anymore. Clear and reset it. */
  cacheclock = 0;
  memset(resourceCache, 0, sizeof(resourceCache[0]) * MAX_CACHE_SLOTS);

  /* Cache is in a pristine "as good as new" condition. */
  return;
}

/* Increment cache time and thereby age of all cached items for LRU ops: */
void ageCache(void)
{
  int i;
  unsigned long long lastaccess, mint = -1, maxt = 0;

  /* Increase age of items in cache by progressing time: */
  cacheclock++;

  /* Wraparound? */
  if (cacheclock == 0) {
    /* Clock wrapped around to zero. We need to recalculate a proper
       age and cacheclock value to keep LRU statistics intact:
       Note: This is mostly an exercise in correctness. In typical
       usage scenarios for this code, e.g., for visual stimulus presentation
       with Psychtoolbox or other OpenGL rendering toolkits, it would take
       anywhere between days and months of uninterruped operation to hit
       the wraparound. Few vision science sessions last more than a few hours.
     */
    for (i = 1; i < cachesize; i++) {
      lastaccess = resourceCache[i].lastaccess;
      if (lastaccess > 0) {
        mint = (lastaccess < mint) ? lastaccess : mint;
        maxt = (lastaccess > maxt) ? lastaccess : maxt;
      }
    }

    if (maxt > 0) {
      /* Shift access times of all valid items by mint - 1, so that
       * the oldest item has access time 1, and the newest item has
       * a much lower access time:
       */
      mint = mint - 1;
      maxt -= mint;
      for (i = 1; i < cachesize; i++) {
        if (resourceCache[i].lastaccess > 0) {
      	  resourceCache[i].lastaccess -= mint;
        }
      }

      /* Adjust cacheclock to maxt + 1. This way, the scale is compressed,
         but the relative order is preserved, ie., distances between cacheclock
         and lastaccess times of all active items in the cache are as before the
         wraparound and LRU semantics is preserved across the wraparound.
       */
      cacheclock = maxt + 1;
    }
    else {
      /* All items were inactive the whole time. Just set clock to 1
         and all is good, as all items have lastaccess time zero.
       */
      cacheclock = 1;
    }
  }
}

void unregisterResource(unsigned int glhandle, unsigned int gltarget)
{
  /* Look it up in cache. */
  int i = cacheLookup(glhandle, gltarget);
  if (i > 0) {
    /* There it is! Unmap and unregister it, release its cache slot: */
    doCudaUnregister(i);
  }
}

void unmapResource(unsigned int glhandle, unsigned int gltarget)
{
  /* Look it up in cache. */
  int i = cacheLookup(glhandle, gltarget);
  if (i > 0) {
    /* There it is! Unmap it: */
    doCudaUnmap(i);
  }
}

void usageExit(int errorAbort)
{
  mexPrintf("\nUsage:\n");
  mexPrintf("======\n\n");
  mexPrintf("%s(0); - Unmap and unregister all resources. Automatically called at shutdown.\n\n", mexFunctionName());
  mexPrintf("%s(1, glObject, glTarget); - Unmap resource glObject of type glTarget\n\n", mexFunctionName());
  mexPrintf("%s(2, glObject, glTarget); - Unregister (and unmap if mapped) resource glObject of type glTarget\n\n", mexFunctionName());
  mexPrintf("%s(3, glObject, glTarget, gpuptr, nrbytes, direction, keepmapped, mapflags);\n", mexFunctionName());
  mexPrintf("- Copy data in VRAM between a CUDA memory buffer and a OpenGL resource. The resource gets\n");
  mexPrintf("  registered for use with CUDA if it isn't registered already, and mapped for use if it\n");
  mexPrintf("  isn't mapped already. 'glObject' and 'glTarget' are OpenGL object and target type handles\n");
  mexPrintf("  of the OpenGL resource. 'gpuptr' is a CUDA memory pointer into the CUDA memory buffer.\n");
  mexPrintf("  'nrbytes' = Number of bytes to copy. 'direction' 0 = OpenGL->CUDA, 1 = CUDA->OpenGL.\n");
  mexPrintf("  'keepmapped' 0 = Unmap immediately after copy [default], 1 = Keep mapped for future use.\n");
  mexPrintf("  Mapping and unmapping comes at some overhead, so should be avoided if possible. However,\n");
  mexPrintf("  OpenGL can't operate on mapped buffers, so one must unmap before OpenGL object is used again\n");
  mexPrintf("  by OpenGL. For a processing flow where a specific object is used as input and output for a\n");
  mexPrintf("  CUDA computation and not used by OpenGL during the CUDA computation, it can make sense to\n");
  mexPrintf("  keep the object mapped inbetween, avoiding one extra map/unmap cycle.\n");
  mexPrintf("  'mapflags': Optional, auto-selected. If set to a >= 0 value, use them as flags for CUDA resource registration.\n\n");
  mexPrintf("%s(4, glObject, glTarget, gpuptr, nrbytes, direction, keepmapped, mapflags);\n\n", mexFunctionName());
  mexPrintf("- Like 3, but avoid data copy. Instead return a uint64 64-Bit memory pointer to the actual mapped resource.\n");
  mexPrintf("  This implies 'keepmapped' == 1, regardless of what you specify.\n");
  mexPrintf("  This allows C-Level plugins to access the mapped resource directly via CUDA pointers.\n\n");
  mexPrintf("%s(5, newlimit); - Increase soft-limit of LRU cache to 'newlimit'.\n\n", mexFunctionName());
  mexPrintf("%s(6, verbose); - Set 'verbose' flag. Defaults to zero. 1 = Verbose status output.\n\n", mexFunctionName());

  /* mexPrintf("%s(5) \n", mexFunctionName()); */

  if (errorAbort) mexErrMsgTxt("Wrong number of arguments. See usage help above.");

  return;
}

void mexExit(void)
{
  cacheFlush();
  firsttime = 1;
}

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

  unsigned int cmd, glhandle, gltarget, direction, flags, keepmapped, ispbo;
  unsigned long nrbytes;
  void *gpuptr;
  int slot = 0;
  cudaGraphicsResource_t resource = NULL;
  struct cudaArray *mappedArray = NULL;
  void* mappedPtr = NULL;
  size_t mappedSize = 0;

  /* Be optimistic, assume success unless told otherwise: */
  cudastatus = cudaSuccess;
  
  if (firsttime) {
      firsttime = 0;

      mexPrintf("\n%s: A simple CUDA <=> OpenGL interoperation interface.\n", mexFunctionName());
      mexPrintf("(c) 2013 by Mario Kleiner. Licensed to you under the MIT license.\n\n");

      /* Reset cache clock to zero and clear the cache: */
      cacheclock = 0;
      memset(resourceCache, 0, sizeof(resourceCache[0]) * MAX_CACHE_SLOTS);

      /* Start off with an effective cache capacity of 8 slots (1 slot is blocked from use): */
      cachesize = 8 + 1;
      firstLRUCycle = 1;

      /* Make sure the cache is flushed at mex file shutdown time: */
      mexAtExit(mexExit);
  }
  
  /* Retrieve command code: Give usage info if none given. */
  if (nrhs < 1) { usageExit(0); return; }

  cmd = (unsigned int) mxGetScalar(prhs[0]);

  /* Change of verbosity? */
  if (cmd == 6) {
    if (nrhs < 2) usageExit(1);  
    verbose = (unsigned int) mxGetScalar(prhs[1]);
    if (verbose) mexPrintf("\n%s: Verbose tracing of operations enabled.\n", mexFunctionName());
    return;
  }

  /* Resizing the LRU cache requested? */
  if (cmd == 5) {
    if (nrhs < 2) usageExit(1);

    /* Reset LRU cache full warning: */
    firstLRUCycle = 1;
    
    slot = (unsigned int) mxGetScalar(prhs[1]);
    
    /* Increment request by 1 to compensate for the "lost" slot 0: */
    slot = slot + 1;
    
    /* Child protections: */
    if (slot > MAX_CACHE_SLOTS) {
      mexPrintf("%s: Requested new softlimit %i for cache exceeds compiled in maximum %i. Will clamp to maximum.\n", mexFunctionName(), slot - 1, MAX_CACHE_SLOTS - 1);
      cachesize = MAX_CACHE_SLOTS;
      return;
    }
    
    if (slot < cachesize) {
      /* Shrinking the cache requested. This implies a full cache flush: */
      mexPrintf("%s: Requested new softlimit %i for cache is smaller than old softlimit %i. Will flush the cache before shrinking it.\n", mexFunctionName(), slot - 1, cachesize - 1);
      cacheFlush();
    }
    
    /* Set new softlimit: */
    cachesize = slot;
    mexPrintf("%s: New softlimit for LRU cache set to %i slots.\n", mexFunctionName(), cachesize - 1);
    
    return;
  }

  if (cmd == 0) {
    /* Cache flush requested: */
    cacheFlush();
    return;
  }

  /* Following ops require at least object handle and target type: */
  if (nrhs < 3) usageExit(1);

  /* Time to increment the age of our cached items by a clock tick: */
  ageCache();

  /* Retrieve OpenGL object handle to our image buffer: */
  glhandle = (unsigned int) mxGetScalar(prhs[1]);

  /* Get GLEnum target: */
  gltarget = (unsigned int) mxGetScalar(prhs[2]);
  
  if (cmd == 1) {
    /* Unmap resource if it is mapped: */
    unmapResource(glhandle, gltarget);
    return;
  }

  if (cmd == 2) {
    /* Unmap and unregister resource if it is mapped and/or registered: */
    unregisterResource(glhandle, gltarget);
    return;
  }

  if (nrhs < 6) usageExit(1);

  /* Retrieve CUDA memory pointer to source/destination CUDA memory buffer: */
  gpuptr = (void*) (unsigned long) mxGetScalar(prhs[3]);
  
  /* Retrieve number of bytes to copy: */
  nrbytes = (unsigned long) mxGetScalar(prhs[4]);
  
  /* Retrieve direction: 0 = OpenGL -> CUDA, 1 = CUDA -> OpenGL : */
  direction = (unsigned int) mxGetScalar(prhs[5]);

  /* Retrieve optional 'keepmapped' flag. */
  keepmapped = 0;
  if (nrhs >= 7) keepmapped = (unsigned int) mxGetScalar(prhs[6]);

  /* Define CUDA optimization flags, depending if this is a OpenGL->CUDA or
   * CUDA->OpenGL copy operation.
   */
  if ((nrhs >= 8) && (mxGetScalar(prhs[7]) >= 0)) {
    /* Override map flags provided. Use them: */
    flags = (unsigned int) mxGetScalar(prhs[7]);
  }
  else {
    /* Use auto-selected map flags: */
    flags = (direction) ? cudaGraphicsRegisterFlagsWriteDiscard : cudaGraphicsRegisterFlagsReadOnly;
  }

  /* Is gltarget a OpenGL pixelbuffer object? Check for gltarget == GL_PACK_BUFFER or GL_UNPACK_BUFFER. */
  ispbo = (gltarget == 35051 || gltarget == 35052) ? 1 : 0;

  /* Copy of data or mapped resource access pointer requested? */
  if (cmd == 3 || cmd == 4) {
    /* Register OpenGL object with CUDA as 'resource': */
    
    /* Already in cache? This would mean it is registered already with compatible mapping flags: */
    slot = cacheInsert(glhandle, gltarget, flags);
    if (slot < 0) {
      /* Not yet in cache. This means it is not registered at this time, either because it
         wasn't registered at all, or because it was registered with incompatible 'flags',
         so it just got unregistered and expelled from the cache. In any case, we need to
         insert it into the cache and register it. -slot is the free target slot for this
         purpose.
       */
       
      /* Turn slot into something useful: */
      slot = -slot;
    
      if (ispbo) {
        /* OpenGL Pixelbuffer object (GL_PACK_BUFFER or GL_UNPACK_BUFFER): */
        cudastatus = cudaGraphicsGLRegisterBuffer(&(resourceCache[slot].resource), glhandle, flags);
      }
      else {
        /* OpenGL texture or renderbuffer object: */
        cudastatus = cudaGraphicsGLRegisterImage(&(resourceCache[slot].resource), glhandle, gltarget, flags);
      }
      
      if (cudastatus != cudaSuccess) {
        mexPrintf("\nmemcpyCudaOpenGL: ERROR in %s(): %s\n", (ispbo) ? "cudaGraphicsGLRegisterBuffer" : "cudaGraphicsGLRegisterImage", cudaGetErrorString(cudastatus));
        resourceCache[slot].resource = NULL;
        goto err_final;
      }

      if (verbose) mexPrintf("\n%s: cacheInsert(%i): CUDA resource registered (globject %i, gltarget %i, flags %i).\n", mexFunctionName(), slot, glhandle, gltarget, flags);
      
      /* Fill cache slot: */
      resourceCache[slot].glhandle = glhandle;
      resourceCache[slot].gltarget = gltarget;
      resourceCache[slot].mapflags = flags;
      resourceCache[slot].lastaccess = cacheclock;
      resourceCache[slot].ismapped = 0;
    }
      
    /* At this point, the resource is stored in slot 'slot' of the cache and registered in a compatible way: */
    
    /* Map the 'resource', unless it is already mapped: */
    if (!resourceCache[slot].ismapped) {
      /* Map it: */
      cudastatus = cudaGraphicsMapResources(1, &(resourceCache[slot].resource), 0);
      if (cudastatus != cudaSuccess) {
        mexPrintf("\nmemcpyCudaOpenGL: ERROR in cudaGraphicsMapResources(): %s\n", cudaGetErrorString(cudastatus));
        goto err_unregister;
      }

      if (verbose) mexPrintf("\n%s: CUDA resource %i mapped (globject %i, gltarget %i, flags %i).\n", mexFunctionName(), slot, glhandle, gltarget, flags);
      
      /* Successfully mapped: */
      resourceCache[slot].ismapped = 1;
    }
    
    /* Get simpler handle: */
    resource = resourceCache[slot].resource;
    
    /* Get mapped resource image array handle or PBO pointer: */
    if (ispbo) {
      cudastatus = cudaGraphicsResourceGetMappedPointer(&mappedPtr, &mappedSize, resource);
    }
    else {
      cudastatus = cudaGraphicsSubResourceGetMappedArray(&mappedArray, resource, 0, 0);
    }
    
    if (cudastatus != cudaSuccess) {
      mexPrintf("\nmemcpyCudaOpenGL: ERROR in %s(): %s\n", (ispbo) ? "cudaGraphicsResourceGetMappedPointer" : "cudaGraphicsSubResourceGetMappedArray", cudaGetErrorString(cudastatus));
      goto err_unmap;
    }
  }
  
  /* Copy of PBO data between CUDA and OpenGL requested? */
  if (cmd == 3 && ispbo) {
    /* Copy from OpenGL PBO to CUDA buffer? */
    if (direction == 0) {
        /* OpenGL -> CUDA copy: */
        cudastatus = cudaMemcpyAsync(gpuptr, (const void*) mappedPtr, (size_t) nrbytes, cudaMemcpyDeviceToDevice, 0);
        if (cudastatus != cudaSuccess) {
            mexPrintf("\nmemcpyCudaOpenGL: ERROR in cudaMemcpyAsync(): %s\n", cudaGetErrorString(cudastatus));
            goto err_unmap;
        }
    }

    if (direction == 1) {
        /* CUDA -> OpenGL copy: */
        cudastatus = cudaMemcpyAsync(mappedPtr, (const void*) gpuptr, (size_t) nrbytes, cudaMemcpyDeviceToDevice, 0);
        if (cudastatus != cudaSuccess) {
            mexPrintf("\nmemcpyCudaOpenGL: ERROR in cudaMemcpyAsync(): %s\n", cudaGetErrorString(cudastatus));
            goto err_unmap;
        }
    }
  }

  /* Copy of texture or renderbuffer data between CUDA and OpenGL requested? */
  if (cmd == 3 && !ispbo) {
    /* Copy from OpenGL object to CUDA buffer? */
    if (direction == 0) {
        /* OpenGL -> CUDA copy: */
        cudastatus = cudaMemcpyFromArrayAsync(gpuptr, (const struct cudaArray*) mappedArray, 0, 0, (size_t) nrbytes, cudaMemcpyDeviceToDevice, 0);
        if (cudastatus != cudaSuccess) {
            mexPrintf("\nmemcpyCudaOpenGL: ERROR in cudaMemcpyFromArray(): %s\n", cudaGetErrorString(cudastatus));
            goto err_unmap;
        }
    }

    if (direction == 1) {
        /* CUDA -> OpenGL copy: */
        cudastatus = cudaMemcpyToArrayAsync((struct cudaArray*) mappedArray, 0, 0, (const void*) gpuptr, (size_t) nrbytes, cudaMemcpyDeviceToDevice, 0);
        if (cudastatus != cudaSuccess) {
            mexPrintf("\nmemcpyCudaOpenGL: ERROR in cudaMemcpyToArray(): %s\n", cudaGetErrorString(cudastatus));
            goto err_unmap;
        }
    }
  }
  
  /* Return of pointers to mapped resource requested? */
  if (cmd == 4) {
    /* Yes: This implies we must not unmap the resource now, as otherwise the
     * returned pointers would be dead on arrival.
     */
    keepmapped = 1;
    
    /* Cast pointer to void* then store it in a 64-Bit unsigned integer return value: */
    plhs[0] = mxCreateNumericMatrix(1, 1, mxUINT64_CLASS, mxREAL);
    *((unsigned long long*) mxGetData(plhs[0])) = (unsigned long long) (void*) ((ispbo) ? mappedPtr : mappedArray);
  }
  
  /* Keep resource mapped? */
  if (slot && !keepmapped) doCudaUnmap(slot);
  
  /* Successfully completed: */
  return;
  
  /* Error handling -- Unwind in reverse order: */
  
err_unmap:
  
  /* Unmap the 'resource': */
  unmapResource(glhandle, gltarget);

err_unregister:
  
  /* Unregister the 'resource': */
  unregisterResource(glhandle, gltarget);

err_final:      
      
  if (cudastatus != cudaSuccess) mexErrMsgTxt("Error in memcpyCudaOpenGL(), reason see above.");
}