Commented out some potentially problem-causing code.
[charm.git] / src / arch / cuda / hybridAPI / cuda-hybrid-api.cu
1 /* 
2  * cuda-hybrid-api.cu
3  *
4  * by Lukasz Wesolowski
5  * 04.01.2008
6  *
7  * an interface for execution on the GPU
8  *
9  * description: 
10  * -user enqueues one or more work requests to the work
11  * request queue (wrQueue) to be executed on the GPU
12  * - a converse function (gpuProgressFn) executes periodically to
13  * offload work requests to the GPU one at a time
14  *
15  */
16
17 #include "wrqueue.h"
18 #include "cuda-hybrid-api.h"
19 #include "stdio.h"
20 #include <cutil.h>
21
22 /* A function in ck.C which casts the void * to a CkCallback object
23  *  and executes the callback 
24  */ 
25 extern void CUDACallbackManager(void * fn); 
26
27 /* initial size of the user-addressed portion of host/device buffer
28  * arrays; the system-addressed portion of host/device buffer arrays
29  * (used when there is no need to share buffers between work requests)
30  * will be equivalant in size.  
31  */ 
32 #define NUM_BUFFERS 128
33 #define MAX_PINNED_REQ 64  
34
35 /* a flag which tells the system to record the time for invocation and
36  *  completion of GPU events: memory allocation, transfer and
37  *  kernel execution
38  */  
39 //#define GPU_PROFILE
40 //#define GPU_DEBUG
41 //#define GPU_TRACE
42 //#define _DEBUG
43
44 /* work request queue */
45 workRequestQueue *wrQueue = NULL; 
46
47 /* pending page-locked memory allocation requests */
48 unsigned int pinnedMemQueueIndex = 0; 
49 pinnedMemReq pinnedMemQueue[MAX_PINNED_REQ];
50
51
52 /* The runtime system keeps track of all allocated buffers on the GPU.
53  * The following arrays contain pointers to host (CPU) data and the
54  * corresponding data on the device (GPU). 
55  */ 
56
57 /* host buffers  */ 
58 void **hostBuffers = NULL; 
59
60 /* device buffers */
61 void **devBuffers = NULL; 
62
63 /* used to assign bufferIDs automatically by the system if the user 
64    specifies an invalid bufferID */
65 unsigned int nextBuffer; 
66
67 unsigned int timerHandle; 
68
69 #ifdef GPU_PROFILE
70
71 /* event types */
72 #define DATA_SETUP          1            
73 #define KERNEL_EXECUTION    2
74 #define DATA_CLEANUP        3
75
76 typedef struct gpuEventTimer {
77   float startTime; 
78   float endTime; 
79   int eventType;
80   int ID; 
81 #ifdef GPU_TRACE
82   int stage; 
83   double cmistartTime; 
84   double cmiendTime; 
85 #endif
86 } gpuEventTimer; 
87
88 gpuEventTimer gpuEvents[QUEUE_SIZE_INIT * 3]; 
89 unsigned int timeIndex = 0; 
90 unsigned int runningKernelIndex = 0; 
91 unsigned int dataSetupIndex = 0; 
92 unsigned int dataCleanupIndex = 0; 
93
94 #ifdef GPU_TRACE
95 extern "C" int traceRegisterUserEvent(const char*x, int e);
96 extern "C" void traceUserBracketEvent(int e, double beginT, double endT);
97 extern "C" double CmiWallTimer(); 
98
99 #define GPU_MEM_SETUP 8800
100 #define GPU_KERNEL_EXEC 8801
101 #define GPU_MEM_CLEANUP 8802
102
103 #endif
104
105 #endif
106
107 /* There are separate CUDA streams for kernel execution, data transfer
108  *  into the device, and data transfer out. This allows prefetching of
109  *  data for a subsequent kernel while the previous kernel is
110  *  executing and transferring data out of the device. 
111  */
112 cudaStream_t kernel_stream; 
113 cudaStream_t data_in_stream;
114 cudaStream_t data_out_stream; 
115
116 /* pinnedMallocHost
117  *
118  * schedules a pinned memory allocation so that it does not impede
119  * concurrent asynchronous execution 
120  *
121  */
122 void pinnedMallocHost(pinnedMemReq *reqs) {
123   /*
124   if ( (cudaStreamQuery(kernel_stream) == cudaSuccess) &&
125        (cudaStreamQuery(data_in_stream) == cudaSuccess) &&
126        (cudaStreamQuery(data_out_stream) == cudaSuccess) ) {    
127   */
128
129   /*
130     for (int i=0; i<reqs->nBuffers; i++) {
131       CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **) reqs->hostPtrs[i], 
132                                             reqs->sizes[i])); 
133     }
134
135     free(reqs->hostPtrs);
136     free(reqs->sizes);
137
138     CUDACallbackManager(reqs->callbackFn);
139
140   */
141
142     /*
143   }
144   else {
145     pinnedMemQueue[pinnedMemQueueIndex].hostPtrs = reqs->hostPtrs;
146     pinnedMemQueue[pinnedMemQueueIndex].sizes = reqs->sizes; 
147     pinnedMemQueue[pinnedMemQueueIndex].callbackFn = reqs->callbackFn;     
148     pinnedMemQueueIndex++;
149     if (pinnedMemQueueIndex == MAX_PINNED_REQ) {
150       printf("Error: pinned memory request buffer is overflowing\n"); 
151     }
152   }
153     */  
154 }
155
156 /* flushPinnedMemQueue
157  *
158  * executes pending pinned memory allocation requests
159  *
160  */
161 void flushPinnedMemQueue() {
162   /*
163   for (int i=0; i<pinnedMemQueueIndex; i++) {
164     pinnedMemReq *req = &pinnedMemQueue[i]; 
165     for (int j=0; j<req->nBuffers; j++) {
166       CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **) req->hostPtrs[j], 
167                                             req->sizes[j])); 
168     }
169     free(req->hostPtrs);
170     free(req->sizes);
171     CUDACallbackManager(pinnedMemQueue[i].callbackFn);    
172   }
173   pinnedMemQueueIndex = 0; 
174   */
175 }
176
177 /* allocateBuffers
178  *
179  * allocates a work request's data on the GPU
180  *
181  * used to allocate memory for work request data in advance in order
182  * to allow overlapping the work request's data transfer to the GPU
183  * with the execution of the previous kernel; the allocation needs to
184  * take place before the kernel starts executing in order to allow overlap
185  *
186  */
187
188 void allocateBuffers(workRequest *wr) {
189   dataInfo *bufferInfo = wr->bufferInfo; 
190
191   if (bufferInfo != NULL) {
192
193     for (int i=0; i<wr->nBuffers; i++) {
194       int index = bufferInfo[i].bufferID; 
195       int size = bufferInfo[i].size; 
196
197       if (bufferInfo[i].transferToDevice == 0) {
198         continue; 
199       }
200
201       // if index value is invalid, use an available ID  
202       if (index < 0 || index >= NUM_BUFFERS) {
203         int found = 0; 
204         for (int j=nextBuffer; j<NUM_BUFFERS*2; j++) {
205           if (devBuffers[j] == NULL) {
206             index = j;
207             found = 1; 
208             break;
209           }
210         }
211
212         /* if no index was found, try to search for a value at the
213          * beginning of the system addressed space 
214          */
215         
216         if (!found) {
217           for (int j=NUM_BUFFERS; j<nextBuffer; j++) {
218             if (devBuffers[j] == NULL) {        
219               index = j;
220               found = 1; 
221               break;
222             }
223           }
224         }
225
226         /* if no index was found, print an error */
227         if (!found) {
228           printf("Error: devBuffers is full \n");
229         }
230
231         nextBuffer = index+1; 
232         if (nextBuffer == NUM_BUFFERS * 2) {
233           nextBuffer = NUM_BUFFERS; 
234         }
235         
236         bufferInfo[i].bufferID = index; 
237
238       }      
239       
240       // allocate if the buffer for the corresponding index is NULL 
241       if (devBuffers[index] == NULL) {
242 #ifdef GPU_PRINT_BUFFER_ALLOCATE
243         double mil = 1e6;
244         printf("*** ALLOCATE buffer 0x%x size %f mb\n", devBuffers[index], 1.0*size/mil);
245 #endif
246
247         CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **) &devBuffers[index], size));
248 #ifdef GPU_DEBUG
249         printf("buffer %d allocated at time %.2f size: %d error string: %s\n", 
250                index, cutGetTimerValue(timerHandle), size, 
251                cudaGetErrorString( cudaGetLastError() ) );
252 #endif
253       }
254     }
255   }
256 }
257
258
259 /* setupData
260  *  copy data to the GPU before kernel execution 
261  */
262 void setupData(workRequest *wr) {
263   dataInfo *bufferInfo = wr->bufferInfo; 
264
265   if (bufferInfo != NULL) {
266     for (int i=0; i<wr->nBuffers; i++) {
267       int index = bufferInfo[i].bufferID; 
268       int size = bufferInfo[i].size; 
269       hostBuffers[index] = bufferInfo[i].hostBuffer; 
270       
271       /* allocate if the buffer for the corresponding index is NULL */
272       /*
273       if (devBuffers[index] == NULL) {
274         CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **) &devBuffers[index], size));
275 #ifdef GPU_DEBUG
276         printf("buffer %d allocated %.2f\n", index,
277                cutGetTimerValue(timerHandle)); 
278 #endif
279       }
280       */
281       
282       if (bufferInfo[i].transferToDevice) {
283         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpyAsync(devBuffers[index], 
284           hostBuffers[index], size, cudaMemcpyHostToDevice, data_in_stream));
285 #ifdef GPU_DEBUG
286         printf("transferToDevice bufId: %d at time %.2f size: %d " 
287                "error string: %s\n", index, cutGetTimerValue(timerHandle), 
288                size, cudaGetErrorString( cudaGetLastError() )); 
289 #endif  
290         /*
291         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(devBuffers[index], 
292           hostBuffers[index], size, cudaMemcpyHostToDevice));
293         */
294
295       }
296     }
297   }
298
299
300 /* copybackData
301  *  transfer data from the GPU to the CPU after a work request is done 
302  */ 
303 void copybackData(workRequest *wr) {
304   dataInfo *bufferInfo = wr->bufferInfo; 
305
306   if (bufferInfo != NULL) {
307     int nBuffers = wr->nBuffers; 
308     
309     for (int i=0; i<nBuffers; i++) {
310       int index = bufferInfo[i].bufferID; 
311       int size = bufferInfo[i].size; 
312       
313       if (bufferInfo[i].transferFromDevice) {
314 #ifdef GPU_DEBUG
315         printf("transferFromDevice: %d at time %.2f size: %d "
316                "error string: %s\n", index, cutGetTimerValue(timerHandle), 
317                size, cudaGetErrorString( cudaGetLastError() )); 
318 #endif
319         
320         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpyAsync(hostBuffers[index], 
321           devBuffers[index], size, cudaMemcpyDeviceToHost,
322           data_out_stream));
323         
324         /*
325         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(hostBuffers[index], 
326           devBuffers[index], size, cudaMemcpyDeviceToHost));
327         */
328       }
329     }     
330   }
331 }
332
333 /* frees GPU memory for buffers specified by the user; also frees the
334  *  work request's bufferInfo array
335  */
336 void freeMemory(workRequest *wr) {
337   dataInfo *bufferInfo = wr->bufferInfo;   
338   int nBuffers = wr->nBuffers; 
339   if (bufferInfo != NULL) {
340     for (int i=0; i<nBuffers; i++) {    
341       int index = bufferInfo[i].bufferID; 
342       if (bufferInfo[i].freeBuffer) {
343 #ifdef GPU_PRINT_BUFFER_ALLOCATE
344         printf("*** FREE buffer 0x%x\n", devBuffers[index]);
345 #endif
346
347 #ifdef GPU_DEBUG
348         printf("buffer %d freed at time %.2f error string: %s\n", 
349                index, cutGetTimerValue(timerHandle),  
350                cudaGetErrorString( cudaGetLastError() ));
351 #endif 
352         CUDA_SAFE_CALL_NO_SYNC(cudaFree(devBuffers[index])); 
353         devBuffers[index] = NULL; 
354       }
355     }
356     free(bufferInfo); 
357   }
358 }
359
360 /* kernelSelect
361  * a switch statement defined by the user to allow the library to execute
362  * the correct kernel 
363  */ 
364 void kernelSelect(workRequest *wr);
365
366 /* initHybridAPI
367  *   initializes the work request queue, host/device buffer pointer
368  *   arrays, and CUDA streams
369  */
370 void initHybridAPI(int myPe) {
371
372   int deviceCount;
373   cudaGetDeviceCount(&deviceCount);
374
375   cudaSetDevice(myPe % deviceCount); 
376
377   initWRqueue(&wrQueue);
378
379   /* allocate host/device buffers array (both user and
380      system-addressed) */
381   hostBuffers = (void **) malloc(NUM_BUFFERS * 2 * sizeof(void *)); 
382   devBuffers = (void **) malloc(NUM_BUFFERS * 2 * sizeof(void *)); 
383
384   /* initialize device array to NULL */ 
385   for (int i=0; i<NUM_BUFFERS*2; i++) {
386     devBuffers[i] = NULL; 
387   }
388   
389   CUDA_SAFE_CALL_NO_SYNC(cudaStreamCreate(&kernel_stream)); 
390   CUDA_SAFE_CALL_NO_SYNC(cudaStreamCreate(&data_in_stream)); 
391   CUDA_SAFE_CALL_NO_SYNC(cudaStreamCreate(&data_out_stream)); 
392
393 #ifdef GPU_PROFILE
394   CUT_SAFE_CALL(cutCreateTimer(&timerHandle));
395   CUT_SAFE_CALL(cutStartTimer(timerHandle));
396 #endif
397
398   nextBuffer = NUM_BUFFERS;  
399
400 #ifdef GPU_TRACE
401   traceRegisterUserEvent("GPU Memory Setup", GPU_MEM_SETUP);
402   traceRegisterUserEvent("GPU Kernel Execution", GPU_KERNEL_EXEC);
403   traceRegisterUserEvent("GPU Memory Cleanup", GPU_MEM_CLEANUP);
404 #endif
405 }
406
407 /* gpuProgressFn
408  *  called periodically to monitor work request progress, and perform
409  *  the prefetch of data for a subsequent work request
410  */
411 void gpuProgressFn() {
412   if (wrQueue == NULL) {
413     printf("Error: work request queue not initialized\n"); 
414     return; 
415   }
416   
417   if (isEmpty(wrQueue)) {
418     //    flushPinnedMemQueue();    
419     return;
420   } 
421
422   // while (!isEmpty(wrQueue)) {
423     int returnVal; 
424     workRequest *head = firstElement(wrQueue); 
425     workRequest *second = secondElement(wrQueue);
426     workRequest *third = thirdElement(wrQueue); 
427     if (head->state == QUEUED) {
428 #ifdef GPU_PROFILE
429       gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
430       gpuEvents[timeIndex].eventType = DATA_SETUP; 
431       gpuEvents[timeIndex].ID = head->id; 
432       dataSetupIndex = timeIndex; 
433 #ifdef GPU_TRACE
434       gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
435       gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
436 #endif
437       timeIndex++; 
438 #endif
439       allocateBuffers(head); 
440       setupData(head); 
441       head->state = TRANSFERRING_IN; 
442     }  
443     if (head->state == TRANSFERRING_IN) {
444       if ((returnVal = cudaStreamQuery(data_in_stream)) == cudaSuccess) {
445 #ifdef GPU_PROFILE
446         gpuEvents[dataSetupIndex].endTime = cutGetTimerValue(timerHandle);
447 #ifdef GPU_TRACE
448         gpuEvents[dataSetupIndex].cmiendTime = CmiWallTimer();
449         traceUserBracketEvent(gpuEvents[dataSetupIndex].stage, 
450                               gpuEvents[dataSetupIndex].cmistartTime, 
451                               gpuEvents[dataSetupIndex].cmiendTime); 
452 #endif 
453 #endif
454         if (second != NULL /*&& (second->state == QUEUED)*/) {
455           allocateBuffers(second); 
456         }
457 #ifdef GPU_PROFILE
458         gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
459         gpuEvents[timeIndex].eventType = KERNEL_EXECUTION; 
460         gpuEvents[timeIndex].ID = head->id; 
461         runningKernelIndex = timeIndex; 
462 #ifdef GPU_TRACE
463         gpuEvents[timeIndex].stage = GPU_KERNEL_EXEC; 
464         gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
465 #endif
466         timeIndex++; 
467 #endif
468         //flushPinnedMemQueue();
469         kernelSelect(head); 
470         head->state = EXECUTING; 
471         if (second != NULL) {
472 #ifdef GPU_PROFILE
473           gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
474           gpuEvents[timeIndex].eventType = DATA_SETUP; 
475           gpuEvents[timeIndex].ID = second->id; 
476           dataSetupIndex = timeIndex; 
477 #ifdef GPU_TRACE
478           gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
479           gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
480 #endif
481           timeIndex++; 
482 #endif
483           setupData(second); 
484           second->state = TRANSFERRING_IN;
485         }
486       }
487       /*
488 #ifdef GPU_DEBUG
489       printf("Querying memory stream returned: %d %.2f\n", returnVal, 
490              cutGetTimerValue(timerHandle));
491 #endif  
492       */
493     }
494     if (head->state == EXECUTING) {
495       if ((returnVal = cudaStreamQuery(kernel_stream)) == cudaSuccess) {
496 #ifdef GPU_PROFILE
497         gpuEvents[runningKernelIndex].endTime = cutGetTimerValue(timerHandle); 
498 #ifdef GPU_TRACE
499         gpuEvents[runningKernelIndex].cmiendTime = CmiWallTimer();
500         traceUserBracketEvent(gpuEvents[runningKernelIndex].stage, 
501                               gpuEvents[runningKernelIndex].cmistartTime, 
502                               gpuEvents[runningKernelIndex].cmiendTime); 
503 #endif
504 #endif
505         if (second != NULL && second->state == QUEUED) {
506 #ifdef GPU_PROFILE
507           gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
508           gpuEvents[timeIndex].eventType = DATA_SETUP; 
509           gpuEvents[timeIndex].ID = second->id; 
510           dataSetupIndex = timeIndex; 
511 #ifdef GPU_TRACE
512           gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
513           gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
514 #endif
515           timeIndex++; 
516 #endif
517           allocateBuffers(second); 
518           setupData(second); 
519           second->state = TRANSFERRING_IN;      
520         } 
521         if (second != NULL && second->state == TRANSFERRING_IN) {
522           if (cudaStreamQuery(data_in_stream) == cudaSuccess) {
523 #ifdef GPU_PROFILE
524             gpuEvents[dataSetupIndex].endTime = cutGetTimerValue(timerHandle); 
525 #ifdef GPU_TRACE
526             gpuEvents[dataSetupIndex].cmiendTime = CmiWallTimer();
527             traceUserBracketEvent(gpuEvents[dataSetupIndex].stage, 
528                                   gpuEvents[dataSetupIndex].cmistartTime, 
529                                   gpuEvents[dataSetupIndex].cmiendTime); 
530 #endif
531 #endif
532             if (third != NULL /*&& (third->state == QUEUED)*/) {
533               allocateBuffers(third); 
534             }
535 #ifdef GPU_PROFILE
536             gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
537             gpuEvents[timeIndex].eventType = KERNEL_EXECUTION; 
538             gpuEvents[timeIndex].ID = second->id; 
539             runningKernelIndex = timeIndex; 
540 #ifdef GPU_TRACE
541             gpuEvents[timeIndex].stage = GPU_KERNEL_EXEC; 
542             gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
543 #endif
544             timeIndex++; 
545 #endif
546             //      flushPinnedMemQueue();          
547             kernelSelect(second); 
548             second->state = EXECUTING; 
549             if (third != NULL) {
550 #ifdef GPU_PROFILE
551               gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
552               gpuEvents[timeIndex].eventType = DATA_SETUP; 
553               gpuEvents[timeIndex].ID = third->id; 
554               dataSetupIndex = timeIndex; 
555 #ifdef GPU_TRACE
556               gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
557               gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
558 #endif
559               timeIndex++; 
560 #endif
561               setupData(third); 
562               third->state = TRANSFERRING_IN;   
563             }
564           }
565         }
566 #ifdef GPU_PROFILE
567         gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
568         gpuEvents[timeIndex].eventType = DATA_CLEANUP; 
569         gpuEvents[timeIndex].ID = head->id; 
570         dataCleanupIndex = timeIndex;   
571 #ifdef GPU_TRACE
572         gpuEvents[timeIndex].stage = GPU_MEM_CLEANUP; 
573         gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
574 #endif
575         timeIndex++; 
576 #endif
577         copybackData(head);
578         head->state = TRANSFERRING_OUT;
579       }
580       /*
581 #ifdef GPU_DEBUG
582       printf("Querying kernel completion returned: %d %.2f\n", returnVal,
583              cutGetTimerValue(timerHandle));
584 #endif  
585       */
586     }
587
588     if (head->state == TRANSFERRING_OUT) {
589       if (cudaStreamQuery(data_out_stream) == cudaSuccess) {
590         freeMemory(head); 
591 #ifdef GPU_PROFILE
592         gpuEvents[dataCleanupIndex].endTime = cutGetTimerValue(timerHandle);
593 #ifdef GPU_TRACE
594         gpuEvents[dataCleanupIndex].cmiendTime = CmiWallTimer();
595         traceUserBracketEvent(gpuEvents[dataCleanupIndex].stage, 
596                               gpuEvents[dataCleanupIndex].cmistartTime, 
597                               gpuEvents[dataCleanupIndex].cmiendTime); 
598 #endif
599 #endif
600         dequeue(wrQueue);
601         CUDACallbackManager(head->callbackFn);
602       }
603     }
604     // }
605 }
606
607 /* exitHybridAPI
608  *  cleans up and deletes memory allocated for the queue and the CUDA streams
609  */
610 void exitHybridAPI() {
611   deleteWRqueue(wrQueue); 
612   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(kernel_stream)); 
613   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(data_in_stream)); 
614   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(data_out_stream)); 
615
616 #ifdef GPU_PROFILE
617   for (int i=0; i<timeIndex; i++) {
618     switch (gpuEvents[i].eventType) {
619     case DATA_SETUP:
620       printf("Kernel %d data setup", gpuEvents[i].ID); 
621       break;
622     case DATA_CLEANUP:
623       printf("Kernel %d data cleanup", gpuEvents[i].ID); 
624       break; 
625     case KERNEL_EXECUTION:
626       printf("Kernel %d execution", gpuEvents[i].ID); 
627       break;
628     default:
629       printf("Error, invalid timer identifier\n"); 
630     }
631     printf(" %.2f:%.2f\n", gpuEvents[i].startTime-gpuEvents[0].startTime, gpuEvents[i].endTime-gpuEvents[0].startTime); 
632   }
633
634   CUT_SAFE_CALL(cutStopTimer(timerHandle));
635   CUT_SAFE_CALL(cutDeleteTimer(timerHandle));  
636
637 #endif
638 }