Minor changes in code structure.
[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   if (isEmpty(wrQueue)) {
417     //    flushPinnedMemQueue();    
418     return;
419   } 
420   int returnVal; 
421   workRequest *head = firstElement(wrQueue); 
422   workRequest *second = secondElement(wrQueue);
423   workRequest *third = thirdElement(wrQueue); 
424   if (head->state == QUEUED) {
425 #ifdef GPU_PROFILE
426     gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
427     gpuEvents[timeIndex].eventType = DATA_SETUP; 
428     gpuEvents[timeIndex].ID = head->id; 
429     dataSetupIndex = timeIndex; 
430 #ifdef GPU_TRACE
431     gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
432     gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
433 #endif
434     timeIndex++; 
435 #endif
436     allocateBuffers(head); 
437     setupData(head); 
438     head->state = TRANSFERRING_IN; 
439   }  
440   if (head->state == TRANSFERRING_IN) {
441     if ((returnVal = cudaStreamQuery(data_in_stream)) == cudaSuccess) {
442 #ifdef GPU_PROFILE
443       gpuEvents[dataSetupIndex].endTime = cutGetTimerValue(timerHandle);
444 #ifdef GPU_TRACE
445       gpuEvents[dataSetupIndex].cmiendTime = CmiWallTimer();
446       traceUserBracketEvent(gpuEvents[dataSetupIndex].stage, 
447                             gpuEvents[dataSetupIndex].cmistartTime, 
448                             gpuEvents[dataSetupIndex].cmiendTime); 
449 #endif 
450 #endif
451       if (second != NULL /*&& (second->state == QUEUED)*/) {
452         allocateBuffers(second); 
453       }
454 #ifdef GPU_PROFILE
455       gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
456       gpuEvents[timeIndex].eventType = KERNEL_EXECUTION; 
457       gpuEvents[timeIndex].ID = head->id; 
458       runningKernelIndex = timeIndex; 
459 #ifdef GPU_TRACE
460       gpuEvents[timeIndex].stage = GPU_KERNEL_EXEC; 
461       gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
462 #endif
463       timeIndex++; 
464 #endif
465       //flushPinnedMemQueue();
466       kernelSelect(head); 
467       head->state = EXECUTING; 
468       if (second != NULL) {
469 #ifdef GPU_PROFILE
470         gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
471         gpuEvents[timeIndex].eventType = DATA_SETUP; 
472         gpuEvents[timeIndex].ID = second->id; 
473         dataSetupIndex = timeIndex; 
474 #ifdef GPU_TRACE
475         gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
476         gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
477 #endif
478         timeIndex++; 
479 #endif
480         setupData(second); 
481         second->state = TRANSFERRING_IN;
482       }
483     }
484       /*
485 #ifdef GPU_DEBUG
486       printf("Querying memory stream returned: %d %.2f\n", returnVal, 
487              cutGetTimerValue(timerHandle));
488 #endif  
489       */
490   }
491   if (head->state == EXECUTING) {
492     if ((returnVal = cudaStreamQuery(kernel_stream)) == cudaSuccess) {
493 #ifdef GPU_PROFILE
494       gpuEvents[runningKernelIndex].endTime = cutGetTimerValue(timerHandle); 
495 #ifdef GPU_TRACE
496       gpuEvents[runningKernelIndex].cmiendTime = CmiWallTimer();
497       traceUserBracketEvent(gpuEvents[runningKernelIndex].stage, 
498                             gpuEvents[runningKernelIndex].cmistartTime, 
499                             gpuEvents[runningKernelIndex].cmiendTime); 
500 #endif
501 #endif
502       if (second != NULL && second->state == QUEUED) {
503 #ifdef GPU_PROFILE
504         gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
505         gpuEvents[timeIndex].eventType = DATA_SETUP; 
506         gpuEvents[timeIndex].ID = second->id; 
507         dataSetupIndex = timeIndex; 
508 #ifdef GPU_TRACE
509         gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
510         gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
511 #endif
512         timeIndex++; 
513 #endif
514         allocateBuffers(second); 
515         setupData(second); 
516         second->state = TRANSFERRING_IN;        
517       } 
518       if (second != NULL && second->state == TRANSFERRING_IN) {
519         if (cudaStreamQuery(data_in_stream) == cudaSuccess) {
520 #ifdef GPU_PROFILE
521           gpuEvents[dataSetupIndex].endTime = cutGetTimerValue(timerHandle); 
522 #ifdef GPU_TRACE
523           gpuEvents[dataSetupIndex].cmiendTime = CmiWallTimer();
524           traceUserBracketEvent(gpuEvents[dataSetupIndex].stage, 
525                                 gpuEvents[dataSetupIndex].cmistartTime, 
526                                 gpuEvents[dataSetupIndex].cmiendTime); 
527 #endif
528 #endif
529           if (third != NULL /*&& (third->state == QUEUED)*/) {
530             allocateBuffers(third); 
531           }
532 #ifdef GPU_PROFILE
533           gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
534           gpuEvents[timeIndex].eventType = KERNEL_EXECUTION; 
535           gpuEvents[timeIndex].ID = second->id; 
536           runningKernelIndex = timeIndex; 
537 #ifdef GPU_TRACE
538           gpuEvents[timeIndex].stage = GPU_KERNEL_EXEC; 
539           gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
540 #endif
541           timeIndex++; 
542 #endif
543           //        flushPinnedMemQueue();          
544           kernelSelect(second); 
545           second->state = EXECUTING; 
546           if (third != NULL) {
547 #ifdef GPU_PROFILE
548             gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
549             gpuEvents[timeIndex].eventType = DATA_SETUP; 
550             gpuEvents[timeIndex].ID = third->id; 
551             dataSetupIndex = timeIndex; 
552 #ifdef GPU_TRACE
553             gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
554             gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
555 #endif
556             timeIndex++; 
557 #endif
558             setupData(third); 
559             third->state = TRANSFERRING_IN;     
560           }
561         }
562       }
563 #ifdef GPU_PROFILE
564       gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
565       gpuEvents[timeIndex].eventType = DATA_CLEANUP; 
566       gpuEvents[timeIndex].ID = head->id; 
567       dataCleanupIndex = timeIndex;     
568 #ifdef GPU_TRACE
569       gpuEvents[timeIndex].stage = GPU_MEM_CLEANUP; 
570       gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
571 #endif
572       timeIndex++; 
573 #endif
574       copybackData(head);
575       head->state = TRANSFERRING_OUT;
576     }
577       /*
578 #ifdef GPU_DEBUG
579       printf("Querying kernel completion returned: %d %.2f\n", returnVal,
580              cutGetTimerValue(timerHandle));
581 #endif  
582       */
583   }
584   if (head->state == TRANSFERRING_OUT) {
585     if (cudaStreamQuery(data_out_stream) == cudaSuccess) {
586       freeMemory(head); 
587 #ifdef GPU_PROFILE
588       gpuEvents[dataCleanupIndex].endTime = cutGetTimerValue(timerHandle);
589 #ifdef GPU_TRACE
590       gpuEvents[dataCleanupIndex].cmiendTime = CmiWallTimer();
591       traceUserBracketEvent(gpuEvents[dataCleanupIndex].stage, 
592                             gpuEvents[dataCleanupIndex].cmistartTime, 
593                             gpuEvents[dataCleanupIndex].cmiendTime); 
594 #endif
595 #endif
596       dequeue(wrQueue);
597       CUDACallbackManager(head->callbackFn);
598     }
599   }
600   
601 }
602
603 /* exitHybridAPI
604  *  cleans up and deletes memory allocated for the queue and the CUDA streams
605  */
606 void exitHybridAPI() {
607   deleteWRqueue(wrQueue); 
608   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(kernel_stream)); 
609   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(data_in_stream)); 
610   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(data_out_stream)); 
611
612 #ifdef GPU_PROFILE
613   for (int i=0; i<timeIndex; i++) {
614     switch (gpuEvents[i].eventType) {
615     case DATA_SETUP:
616       printf("Kernel %d data setup", gpuEvents[i].ID); 
617       break;
618     case DATA_CLEANUP:
619       printf("Kernel %d data cleanup", gpuEvents[i].ID); 
620       break; 
621     case KERNEL_EXECUTION:
622       printf("Kernel %d execution", gpuEvents[i].ID); 
623       break;
624     default:
625       printf("Error, invalid timer identifier\n"); 
626     }
627     printf(" %.2f:%.2f\n", gpuEvents[i].startTime-gpuEvents[0].startTime, gpuEvents[i].endTime-gpuEvents[0].startTime); 
628   }
629
630   CUT_SAFE_CALL(cutStopTimer(timerHandle));
631   CUT_SAFE_CALL(cutDeleteTimer(timerHandle));  
632
633 #endif
634 }