CUDA: Small fix in the delayed free.
[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 extern int CmiMyPe();
27
28 /* initial size of the user-addressed portion of host/device buffer
29  * arrays; the system-addressed portion of host/device buffer arrays
30  * (used when there is no need to share buffers between work requests)
31  * will be equivalant in size.  
32  */ 
33 #define NUM_BUFFERS 256
34 #define MAX_PINNED_REQ 64  
35 #define MAX_DELAYED_FREE_REQS 64  
36
37 /* a flag which tells the system to record the time for invocation and
38  *  completion of GPU events: memory allocation, transfer and
39  *  kernel execution
40  */  
41 //#define GPU_PROFILE
42 //#define GPU_DEBUG
43 //#define GPU_TRACE
44 //#define _DEBUG
45
46 /* work request queue */
47 workRequestQueue *wrQueue = NULL; 
48
49 /* pending page-locked memory allocation requests */
50 unsigned int pinnedMemQueueIndex = 0; 
51 pinnedMemReq pinnedMemQueue[MAX_PINNED_REQ];
52
53 unsigned int currentDfr = 0;
54 void *delayedFreeReqs[MAX_DELAYED_FREE_REQS];
55
56 #ifdef GPU_MEMPOOL
57 #define GPU_MEMPOOL_NUM_SLOTS 15
58
59 CkVec<BufferPool> memPoolFreeBufs;
60 CkVec<int> memPoolBoundaries;
61 //int memPoolBoundaries[GPU_MEMPOOL_NUM_SLOTS];
62 #endif
63
64 /* The runtime system keeps track of all allocated buffers on the GPU.
65  * The following arrays contain pointers to host (CPU) data and the
66  * corresponding data on the device (GPU). 
67  */ 
68
69 /* host buffers  */ 
70 void **hostBuffers = NULL; 
71
72 /* device buffers */
73 void **devBuffers = NULL; 
74
75 /* used to assign bufferIDs automatically by the system if the user 
76    specifies an invalid bufferID */
77 unsigned int nextBuffer; 
78
79 unsigned int timerHandle; 
80
81 #ifdef GPU_PROFILE
82
83 /* event types */
84 #define DATA_SETUP          1            
85 #define KERNEL_EXECUTION    2
86 #define DATA_CLEANUP        3
87
88 typedef struct gpuEventTimer {
89   float startTime; 
90   float endTime; 
91   int eventType;
92   int ID; 
93 #ifdef GPU_TRACE
94   int stage; 
95   double cmistartTime; 
96   double cmiendTime; 
97 #endif
98 } gpuEventTimer; 
99
100 gpuEventTimer gpuEvents[QUEUE_SIZE_INIT * 3]; 
101 unsigned int timeIndex = 0; 
102 unsigned int runningKernelIndex = 0; 
103 unsigned int dataSetupIndex = 0; 
104 unsigned int dataCleanupIndex = 0; 
105
106 #if defined GPU_TRACE || defined GPU_INSTRUMENT_WRS
107 extern "C" double CmiWallTimer(); 
108 #endif
109
110 #ifdef GPU_TRACE
111 extern "C" int traceRegisterUserEvent(const char*x, int e);
112 extern "C" void traceUserBracketEvent(int e, double beginT, double endT);
113
114 #define GPU_MEM_SETUP 8800
115 #define GPU_KERNEL_EXEC 8801
116 #define GPU_MEM_CLEANUP 8802
117
118 #endif
119
120 #endif
121
122 #ifdef GPU_INSTRUMENT_WRS
123 CkVec<CkVec<CkVec<RequestTimeInfo> > > avgTimes;
124 bool initialized_instrument;
125 bool initializedInstrument();
126 #endif
127
128 /* There are separate CUDA streams for kernel execution, data transfer
129  *  into the device, and data transfer out. This allows prefetching of
130  *  data for a subsequent kernel while the previous kernel is
131  *  executing and transferring data out of the device. 
132  */
133 cudaStream_t kernel_stream; 
134 cudaStream_t data_in_stream;
135 cudaStream_t data_out_stream; 
136
137 /* pinnedMallocHost
138  *
139  * schedules a pinned memory allocation so that it does not impede
140  * concurrent asynchronous execution 
141  *
142  */
143 void pinnedMallocHost(pinnedMemReq *reqs) {
144
145   if ( (cudaStreamQuery(kernel_stream) == cudaSuccess) &&
146        (cudaStreamQuery(data_in_stream) == cudaSuccess) &&
147        (cudaStreamQuery(data_out_stream) == cudaSuccess) ) {    
148
149
150
151     for (int i=0; i<reqs->nBuffers; i++) {
152       CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **) reqs->hostPtrs[i], 
153                                             reqs->sizes[i])); 
154     }
155
156     free(reqs->hostPtrs);
157     free(reqs->sizes);
158
159     CUDACallbackManager(reqs->callbackFn);
160
161   }
162   else {
163     pinnedMemQueue[pinnedMemQueueIndex].hostPtrs = reqs->hostPtrs;
164     pinnedMemQueue[pinnedMemQueueIndex].sizes = reqs->sizes; 
165     pinnedMemQueue[pinnedMemQueueIndex].callbackFn = reqs->callbackFn;     
166     pinnedMemQueueIndex++;
167     if (pinnedMemQueueIndex == MAX_PINNED_REQ) {
168       printf("Error: pinned memory request buffer is overflowing\n"); 
169     }
170   }
171 }
172
173 void delayedFree(void *ptr){
174   if(currentDfr == MAX_DELAYED_FREE_REQS){
175     printf("Ran out of DFR queue space. Increase MAX_DELAYED_FREE_REQS\n");
176     exit(-1);
177   }
178   else{
179     delayedFreeReqs[currentDfr] = ptr;
180   }
181   currentDfr++;
182 }
183
184 void flushDelayedFrees(){
185   for(int i = 0; i < currentDfr; i++){
186     if(delayedFreeReqs[i] == NULL){
187       printf("recorded NULL ptr in delayedFree()");
188       exit(-1);
189     }
190     cudaFreeHost(delayedFreeReqs[i]);
191   }
192   currentDfr = 0; 
193 }
194
195 /* flushPinnedMemQueue
196  *
197  * executes pending pinned memory allocation requests
198  *
199  */
200 void flushPinnedMemQueue() {
201
202   for (int i=0; i<pinnedMemQueueIndex; i++) {
203     pinnedMemReq *req = &pinnedMemQueue[i]; 
204     for (int j=0; j<req->nBuffers; j++) {
205       CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **) req->hostPtrs[j], 
206                                             req->sizes[j])); 
207     }
208     free(req->hostPtrs);
209     free(req->sizes);
210     CUDACallbackManager(pinnedMemQueue[i].callbackFn);    
211   }
212   pinnedMemQueueIndex = 0; 
213
214 }
215
216 /* allocateBuffers
217  *
218  * allocates a work request's data on the GPU
219  *
220  * used to allocate memory for work request data in advance in order
221  * to allow overlapping the work request's data transfer to the GPU
222  * with the execution of the previous kernel; the allocation needs to
223  * take place before the kernel starts executing in order to allow overlap
224  *
225  */
226
227 void allocateBuffers(workRequest *wr) {
228   dataInfo *bufferInfo = wr->bufferInfo; 
229
230   if (bufferInfo != NULL) {
231
232     for (int i=0; i<wr->nBuffers; i++) {
233       int index = bufferInfo[i].bufferID; 
234       int size = bufferInfo[i].size; 
235
236       // if index value is invalid, use an available ID  
237       if (index < 0 || index >= NUM_BUFFERS) {
238         int found = 0; 
239         for (int j=nextBuffer; j<NUM_BUFFERS*2; j++) {
240           if (devBuffers[j] == NULL) {
241             index = j;
242             found = 1; 
243             break;
244           }
245         }
246
247         /* if no index was found, try to search for a value at the
248          * beginning of the system addressed space 
249          */
250         
251         if (!found) {
252           for (int j=NUM_BUFFERS; j<nextBuffer; j++) {
253             if (devBuffers[j] == NULL) {        
254               index = j;
255               found = 1; 
256               break;
257             }
258           }
259         }
260
261         /* if no index was found, print an error */
262         if (!found) {
263           printf("Error: devBuffers is full \n");
264         }
265
266         nextBuffer = index+1; 
267         if (nextBuffer == NUM_BUFFERS * 2) {
268           nextBuffer = NUM_BUFFERS; 
269         }
270         
271         bufferInfo[i].bufferID = index; 
272
273       }      
274       
275       // allocate if the buffer for the corresponding index is NULL 
276       if (devBuffers[index] == NULL && size > 0) {
277 #ifdef GPU_PRINT_BUFFER_ALLOCATE
278         double mil = 1e3;
279         printf("*** ALLOCATE buffer 0x%x (%d) size %f kb\n", devBuffers[index], index, 1.0*size/mil);
280
281 #endif
282
283         CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **) &devBuffers[index], size));
284 #ifdef GPU_DEBUG
285         printf("buffer %d allocated at time %.2f size: %d error string: %s\n", 
286                index, cutGetTimerValue(timerHandle), size, 
287                cudaGetErrorString( cudaGetLastError() ) );
288 #endif
289       }
290     }
291   }
292 }
293
294
295 /* setupData
296  *  copy data to the GPU before kernel execution 
297  */
298 void setupData(workRequest *wr) {
299   dataInfo *bufferInfo = wr->bufferInfo; 
300
301   if (bufferInfo != NULL) {
302     for (int i=0; i<wr->nBuffers; i++) {
303       int index = bufferInfo[i].bufferID; 
304       int size = bufferInfo[i].size; 
305       hostBuffers[index] = bufferInfo[i].hostBuffer; 
306       
307       /* allocate if the buffer for the corresponding index is NULL */
308       /*
309       if (devBuffers[index] == NULL) {
310         CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **) &devBuffers[index], size));
311 #ifdef GPU_DEBUG
312         printf("buffer %d allocated %.2f\n", index,
313                cutGetTimerValue(timerHandle)); 
314 #endif
315       }
316       */
317       
318       if (bufferInfo[i].transferToDevice && size > 0) {
319         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpyAsync(devBuffers[index], 
320           hostBuffers[index], size, cudaMemcpyHostToDevice, data_in_stream));
321 #ifdef GPU_DEBUG
322         printf("transferToDevice bufId: %d at time %.2f size: %d " 
323                "error string: %s\n", index, cutGetTimerValue(timerHandle), 
324                size, cudaGetErrorString( cudaGetLastError() )); 
325 #endif  
326         /*
327         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(devBuffers[index], 
328           hostBuffers[index], size, cudaMemcpyHostToDevice));
329         */
330
331       }
332     }
333   }
334
335
336 /* copybackData
337  *  transfer data from the GPU to the CPU after a work request is done 
338  */ 
339 void copybackData(workRequest *wr) {
340   dataInfo *bufferInfo = wr->bufferInfo; 
341
342   if (bufferInfo != NULL) {
343     int nBuffers = wr->nBuffers; 
344     
345     for (int i=0; i<nBuffers; i++) {
346       int index = bufferInfo[i].bufferID; 
347       int size = bufferInfo[i].size; 
348       
349       if (bufferInfo[i].transferFromDevice && size > 0) {
350 #ifdef GPU_DEBUG
351         printf("transferFromDevice: %d at time %.2f size: %d "
352                "error string: %s\n", index, cutGetTimerValue(timerHandle), 
353                size, cudaGetErrorString( cudaGetLastError() )); 
354 #endif
355         
356         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpyAsync(hostBuffers[index], 
357           devBuffers[index], size, cudaMemcpyDeviceToHost,
358           data_out_stream));
359         
360         /*
361         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(hostBuffers[index], 
362           devBuffers[index], size, cudaMemcpyDeviceToHost));
363         */
364       }
365     }     
366   }
367 }
368
369 /* frees GPU memory for buffers specified by the user; also frees the
370  *  work request's bufferInfo array
371  */
372 void freeMemory(workRequest *wr) {
373   dataInfo *bufferInfo = wr->bufferInfo;   
374   int nBuffers = wr->nBuffers; 
375   if (bufferInfo != NULL) {
376     for (int i=0; i<nBuffers; i++) {    
377       int index = bufferInfo[i].bufferID; 
378       if (bufferInfo[i].freeBuffer) {
379 #ifdef GPU_PRINT_BUFFER_ALLOCATE
380         printf("*** FREE buffer 0x%x (%d)\n", devBuffers[index], index);
381 #endif
382
383 #ifdef GPU_DEBUG
384         printf("buffer %d freed at time %.2f error string: %s\n", 
385                index, cutGetTimerValue(timerHandle),  
386                cudaGetErrorString( cudaGetLastError() ));
387 #endif 
388         CUDA_SAFE_CALL_NO_SYNC(cudaFree(devBuffers[index])); 
389         devBuffers[index] = NULL; 
390       }
391     }
392     free(bufferInfo); 
393   }
394 }
395
396 /* 
397  * a switch statement defined by the user to allow the library to execute
398  * the correct kernel 
399  */ 
400 void kernelSelect(workRequest *wr);
401
402 /* initHybridAPI
403  *   initializes the work request queue, host/device buffer pointer
404  *   arrays, and CUDA streams
405  */
406 void initHybridAPI(int myPe) {
407
408   int deviceCount;
409   cudaGetDeviceCount(&deviceCount);
410
411   cudaSetDevice(myPe % deviceCount); 
412
413   initWRqueue(&wrQueue);
414
415   /* allocate host/device buffers array (both user and
416      system-addressed) */
417   hostBuffers = (void **) malloc(NUM_BUFFERS * 2 * sizeof(void *)); 
418   devBuffers = (void **) malloc(NUM_BUFFERS * 2 * sizeof(void *)); 
419
420   /* initialize device array to NULL */ 
421   for (int i=0; i<NUM_BUFFERS*2; i++) {
422     devBuffers[i] = NULL; 
423   }
424   
425   CUDA_SAFE_CALL_NO_SYNC(cudaStreamCreate(&kernel_stream)); 
426   CUDA_SAFE_CALL_NO_SYNC(cudaStreamCreate(&data_in_stream)); 
427   CUDA_SAFE_CALL_NO_SYNC(cudaStreamCreate(&data_out_stream)); 
428
429 #ifdef GPU_PROFILE
430   CUT_SAFE_CALL(cutCreateTimer(&timerHandle));
431   CUT_SAFE_CALL(cutStartTimer(timerHandle));
432 #endif
433
434   nextBuffer = NUM_BUFFERS;  
435
436 #ifdef GPU_TRACE
437   traceRegisterUserEvent("GPU Memory Setup", GPU_MEM_SETUP);
438   traceRegisterUserEvent("GPU Kernel Execution", GPU_KERNEL_EXEC);
439   traceRegisterUserEvent("GPU Memory Cleanup", GPU_MEM_CLEANUP);
440 #endif
441
442 #ifdef GPU_MEMPOOL
443   int nslots = GPU_MEMPOOL_NUM_SLOTS;
444   int *sizes;
445   sizes = (int *)malloc(sizeof(int)*nslots); 
446
447   memPoolBoundaries.reserve(GPU_MEMPOOL_NUM_SLOTS);
448   memPoolBoundaries.length() = GPU_MEMPOOL_NUM_SLOTS;
449
450   int bufSize = GPU_MEMPOOL_MIN_BUFFER_SIZE;
451   for(int i = 0; i < GPU_MEMPOOL_NUM_SLOTS; i++){
452     memPoolBoundaries[i] = bufSize;
453     bufSize = bufSize << 1;
454   }
455
456   //1K
457   sizes[0] = 512; 
458   //2K
459   sizes[1] = 512;
460   //4K
461   sizes[2] = 64;
462   //8K
463   sizes[3] = 64;
464   //16K
465   sizes[4] = 32;
466   //32K
467   sizes[5] = 32;
468   //64K
469   sizes[6] = 32;
470   //128K
471   sizes[7] = 32;
472   //256K
473   sizes[8] = 32;
474   //512K
475   sizes[9] = 32;
476   //1M
477   sizes[10] = 170;
478   //2M
479   sizes[11] = 16;
480   //4M
481   sizes[12] = 4;
482   //8M
483   sizes[13] = 2;
484   //16M
485   sizes[14] = 2; 
486
487   createPool(sizes, nslots, memPoolFreeBufs);
488   printf("[%d] done creating buffer pool\n", CmiMyPe());
489
490 #endif
491
492 #ifdef GPU_INSTRUMENT_WRS
493   initialized_instrument = false;
494 #endif
495 }
496
497 /* gpuProgressFn
498  *  called periodically to monitor work request progress, and perform
499  *  the prefetch of data for a subsequent work request
500  */
501 void gpuProgressFn() {
502   if (wrQueue == NULL) {
503     printf("Error: work request queue not initialized\n"); 
504     return; 
505   }
506   if (isEmpty(wrQueue)) {
507     flushPinnedMemQueue();    
508     flushDelayedFrees();
509     return;
510   } 
511   int returnVal; 
512   workRequest *head = firstElement(wrQueue); 
513   workRequest *second = secondElement(wrQueue);
514   workRequest *third = thirdElement(wrQueue); 
515
516   if (head->state == QUEUED) {
517 #ifdef GPU_PROFILE
518     gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
519     gpuEvents[timeIndex].eventType = DATA_SETUP; 
520     gpuEvents[timeIndex].ID = head->id; 
521     dataSetupIndex = timeIndex; 
522 #ifdef GPU_TRACE
523     gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
524     gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
525 #endif
526     timeIndex++; 
527 #endif
528
529 #ifdef GPU_INSTRUMENT_WRS
530     head->startTime = CmiWallTimer(); 
531 #endif
532
533     allocateBuffers(head); 
534     setupData(head); 
535     head->state = TRANSFERRING_IN; 
536   }  
537   if (head->state == TRANSFERRING_IN) {
538     if ((returnVal = cudaStreamQuery(data_in_stream)) == cudaSuccess) {
539 #ifdef GPU_PROFILE
540       gpuEvents[dataSetupIndex].endTime = cutGetTimerValue(timerHandle);
541 #ifdef GPU_TRACE
542       gpuEvents[dataSetupIndex].cmiendTime = CmiWallTimer();
543       traceUserBracketEvent(gpuEvents[dataSetupIndex].stage, 
544                             gpuEvents[dataSetupIndex].cmistartTime, 
545                             gpuEvents[dataSetupIndex].cmiendTime); 
546 #endif 
547 #endif
548
549 #ifdef GPU_INSTRUMENT_WRS
550       {
551         if(initializedInstrument()){
552           double tt = CmiWallTimer()-(head->startTime);
553           int index = head->chareIndex;
554           char type = head->compType;
555           char phase = head->compPhase;
556
557           CkVec<RequestTimeInfo> &vec = avgTimes[index][type];
558           if(vec.length() <= phase){
559             vec.growAtLeast(phase);
560           }
561           vec[phase].transferTime += tt;
562         }
563       }
564 #endif
565
566       if (second != NULL /*&& (second->state == QUEUED)*/) {
567         allocateBuffers(second); 
568       }
569 #ifdef GPU_PROFILE
570       gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
571       gpuEvents[timeIndex].eventType = KERNEL_EXECUTION; 
572       gpuEvents[timeIndex].ID = head->id; 
573       runningKernelIndex = timeIndex; 
574 #ifdef GPU_TRACE
575       gpuEvents[timeIndex].stage = GPU_KERNEL_EXEC; 
576       gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
577 #endif
578       timeIndex++; 
579 #endif
580 #ifdef GPU_INSTRUMENT_WRS
581       head->startTime = CmiWallTimer(); 
582 #endif
583
584       //flushPinnedMemQueue();
585       flushDelayedFrees();
586       kernelSelect(head); 
587
588       head->state = EXECUTING; 
589       if (second != NULL) {
590 #ifdef GPU_PROFILE
591         gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
592         gpuEvents[timeIndex].eventType = DATA_SETUP; 
593         gpuEvents[timeIndex].ID = second->id; 
594         dataSetupIndex = timeIndex; 
595 #ifdef GPU_TRACE
596         gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
597         gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
598 #endif
599         timeIndex++; 
600 #endif
601
602 #ifdef GPU_INSTRUMENT_WRS
603         second->startTime = CmiWallTimer();
604 #endif
605         setupData(second); 
606         second->state = TRANSFERRING_IN;
607       }
608     }
609       /*
610 #ifdef GPU_DEBUG
611       printf("Querying memory stream returned: %d %.2f\n", returnVal, 
612              cutGetTimerValue(timerHandle));
613 #endif  
614       */
615   }
616   if (head->state == EXECUTING) {
617     if ((returnVal = cudaStreamQuery(kernel_stream)) == cudaSuccess) {
618 #ifdef GPU_PROFILE
619       gpuEvents[runningKernelIndex].endTime = cutGetTimerValue(timerHandle); 
620 #ifdef GPU_TRACE
621       gpuEvents[runningKernelIndex].cmiendTime = CmiWallTimer();
622       traceUserBracketEvent(gpuEvents[runningKernelIndex].stage, 
623                             gpuEvents[runningKernelIndex].cmistartTime, 
624                             gpuEvents[runningKernelIndex].cmiendTime); 
625 #endif
626 #endif
627 #ifdef GPU_INSTRUMENT_WRS
628       {
629         if(initializedInstrument()){
630           double tt = CmiWallTimer()-(head->startTime);
631           int index = head->chareIndex;
632           char type = head->compType;
633           char phase = head->compPhase;
634
635           CkVec<RequestTimeInfo> &vec = avgTimes[index][type];
636           if(vec.length() <= phase){
637             vec.growAtLeast(phase);
638           }
639           vec[phase].kernelTime += tt;
640         }
641       }
642 #endif
643
644       if (second != NULL && second->state == QUEUED) {
645 #ifdef GPU_PROFILE
646         gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
647         gpuEvents[timeIndex].eventType = DATA_SETUP; 
648         gpuEvents[timeIndex].ID = second->id; 
649         dataSetupIndex = timeIndex; 
650 #ifdef GPU_TRACE
651         gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
652         gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
653 #endif
654         timeIndex++; 
655 #endif
656
657 #ifdef GPU_INSTRUMENT_MS
658         second->startTime = CmiWallTimer();
659 #endif
660         
661         allocateBuffers(second); 
662         setupData(second); 
663         second->state = TRANSFERRING_IN;        
664       } 
665       if (second != NULL && second->state == TRANSFERRING_IN) {
666         if (cudaStreamQuery(data_in_stream) == cudaSuccess) {
667 #ifdef GPU_PROFILE
668           gpuEvents[dataSetupIndex].endTime = cutGetTimerValue(timerHandle); 
669 #ifdef GPU_TRACE
670           gpuEvents[dataSetupIndex].cmiendTime = CmiWallTimer();
671           traceUserBracketEvent(gpuEvents[dataSetupIndex].stage, 
672                                 gpuEvents[dataSetupIndex].cmistartTime, 
673                                 gpuEvents[dataSetupIndex].cmiendTime); 
674 #endif
675 #endif
676 #ifdef GPU_INSTRUMENT_WRS
677           {
678             if(initializedInstrument()){
679               double tt = CmiWallTimer()-(head->startTime);
680               int index = second->chareIndex;
681               char type = second->compType;
682               char phase = second->compPhase;
683
684               CkVec<RequestTimeInfo> &vec = avgTimes[index][type];
685               if(vec.length() <= phase){
686                 vec.growAtLeast(phase);
687               }
688               vec[phase].transferTime += tt;
689             }
690           }
691 #endif
692
693           if (third != NULL /*&& (third->state == QUEUED)*/) {
694             allocateBuffers(third); 
695           }
696 #ifdef GPU_PROFILE
697           gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
698           gpuEvents[timeIndex].eventType = KERNEL_EXECUTION; 
699           gpuEvents[timeIndex].ID = second->id; 
700           runningKernelIndex = timeIndex; 
701 #ifdef GPU_TRACE
702           gpuEvents[timeIndex].stage = GPU_KERNEL_EXEC; 
703           gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
704 #endif
705           timeIndex++; 
706 #endif
707 #ifdef GPU_INSTRUMENT_WRS
708           second->startTime = CmiWallTimer();
709 #endif
710           //        flushPinnedMemQueue();          
711           flushDelayedFrees();
712           kernelSelect(second); 
713           second->state = EXECUTING; 
714           if (third != NULL) {
715 #ifdef GPU_PROFILE
716             gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
717             gpuEvents[timeIndex].eventType = DATA_SETUP; 
718             gpuEvents[timeIndex].ID = third->id; 
719             dataSetupIndex = timeIndex; 
720 #ifdef GPU_TRACE
721             gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
722             gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
723 #endif
724             timeIndex++; 
725 #endif
726
727 #ifdef GPU_INSTRUMENT_WRS
728             third->startTime = CmiWallTimer();
729 #endif
730             setupData(third); 
731             third->state = TRANSFERRING_IN;     
732           }
733         }
734       }
735 #ifdef GPU_PROFILE
736       gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
737       gpuEvents[timeIndex].eventType = DATA_CLEANUP; 
738       gpuEvents[timeIndex].ID = head->id; 
739       dataCleanupIndex = timeIndex;     
740 #ifdef GPU_TRACE
741       gpuEvents[timeIndex].stage = GPU_MEM_CLEANUP; 
742       gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
743 #endif
744       timeIndex++; 
745 #endif
746 #ifdef GPU_INSTRUMENT_WRS
747       head->startTime = CmiWallTimer(); 
748 #endif
749       copybackData(head);
750       head->state = TRANSFERRING_OUT;
751     }
752       /*
753 #ifdef GPU_DEBUG
754       printf("Querying kernel completion returned: %d %.2f\n", returnVal,
755              cutGetTimerValue(timerHandle));
756 #endif  
757       */
758   }
759   if (head->state == TRANSFERRING_OUT) {
760     if (cudaStreamQuery(data_out_stream) == cudaSuccess && cudaStreamQuery(kernel_stream) == cudaSuccess){
761       freeMemory(head); 
762 #ifdef GPU_PROFILE
763       gpuEvents[dataCleanupIndex].endTime = cutGetTimerValue(timerHandle);
764 #ifdef GPU_TRACE
765       gpuEvents[dataCleanupIndex].cmiendTime = CmiWallTimer();
766       traceUserBracketEvent(gpuEvents[dataCleanupIndex].stage, 
767                             gpuEvents[dataCleanupIndex].cmistartTime, 
768                             gpuEvents[dataCleanupIndex].cmiendTime); 
769 #endif
770 #endif
771 #ifdef GPU_INSTRUMENT_WRS
772       {
773         if(initializedInstrument()){
774           double tt = CmiWallTimer()-(head->startTime);
775           int index = head->chareIndex;
776           char type = head->compType;
777           char phase = head->compPhase;
778
779           CkVec<RequestTimeInfo> &vec = avgTimes[index][type];
780           if(vec.length() <= phase){
781             vec.growAtLeast(phase);
782           }
783           vec[phase].cleanupTime += tt;
784           vec[phase].n++;
785         }
786       }
787 #endif
788
789       dequeue(wrQueue);
790       CUDACallbackManager(head->callbackFn);
791     }
792   }
793 }
794
795 #ifdef GPU_MEMPOOL
796 void releasePool(CkVec<BufferPool> &pools);
797 #endif
798 /* exitHybridAPI
799  *  cleans up and deletes memory allocated for the queue and the CUDA streams
800  */
801 void exitHybridAPI() {
802   deleteWRqueue(wrQueue); 
803   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(kernel_stream)); 
804   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(data_in_stream)); 
805   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(data_out_stream)); 
806
807 #ifdef GPU_PROFILE
808   for (int i=0; i<timeIndex; i++) {
809     switch (gpuEvents[i].eventType) {
810     case DATA_SETUP:
811       printf("Kernel %d data setup", gpuEvents[i].ID); 
812       break;
813     case DATA_CLEANUP:
814       printf("Kernel %d data cleanup", gpuEvents[i].ID); 
815       break; 
816     case KERNEL_EXECUTION:
817       printf("Kernel %d execution", gpuEvents[i].ID); 
818       break;
819     default:
820       printf("Error, invalid timer identifier\n"); 
821     }
822     printf(" %.2f:%.2f\n", gpuEvents[i].startTime-gpuEvents[0].startTime, gpuEvents[i].endTime-gpuEvents[0].startTime); 
823   }
824
825   CUT_SAFE_CALL(cutStopTimer(timerHandle));
826   CUT_SAFE_CALL(cutDeleteTimer(timerHandle));  
827
828 #endif
829
830 #ifdef GPU_MEMPOOL
831   releasePool(memPoolFreeBufs);
832 #endif
833
834 }
835
836 #ifdef GPU_MEMPOOL
837 void releasePool(CkVec<BufferPool> &pools){
838   for(int i = 0; i < pools.length(); i++){
839     CUDA_SAFE_CALL_NO_SYNC(cudaFreeHost((void *)pools[i].head));
840   }
841   pools.free();
842 }
843
844 // Create a pool with nslots slots.
845 // There are nbuffers[i] buffers for each buffer size corresponding to entry i
846 // FIXME - list the alignment/fragmentation issues with either of two allocation schemes:
847 // if a single, large buffer is allocated for each subpool
848 // if multiple smaller buffers are allocated for each subpool
849 void createPool(int *nbuffers, int nslots, CkVec<BufferPool> &pools){
850   //pools  = (BufferPool *)malloc(nslots*sizeof(BufferPool));
851   pools.reserve(nslots);
852   pools.length() = nslots;
853
854   for(int i = 0; i < nslots; i++){
855     int bufSize = memPoolBoundaries[i];
856     int numBuffers = nbuffers[i];
857     pools[i].size = bufSize;
858     
859     CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **)(&pools[i].head), 
860                                           (sizeof(Header)+bufSize)*numBuffers));
861     if(pools[i].head == NULL){
862       abort();
863     }
864
865     Header *hd = pools[i].head;
866     Header *previous = NULL;
867     char *memory;
868
869     for(int j = 0; j < numBuffers; j++){
870       hd->slot = i;
871       hd->next = previous;
872       previous = hd;
873       hd++; // move ptr past header
874       memory = (char *)hd;
875       memory += bufSize;
876       hd = (Header *)memory;
877     }
878
879     pools[i].head = previous;
880 #ifdef GPU_MEMPOOL_DEBUG
881     pools[i].num = numBuffers;
882 #endif
883   }
884 }
885
886 int findPool(int size){
887   int boundaryArrayLen = memPoolBoundaries.length();
888   if(size <= memPoolBoundaries[0]){
889     return (0);
890   }
891   else if(size > memPoolBoundaries[boundaryArrayLen-1]){
892     // create new slot
893     memPoolBoundaries.push_back(size);
894
895     BufferPool newpool;
896     CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **)&newpool.head, size+sizeof(Header)));
897     newpool.size = size;
898 #ifdef GPU_MEMPOOL_DEBUG
899     newpool.num = 1;
900 #endif
901     memPoolFreeBufs.push_back(newpool);
902
903     Header *hd = newpool.head;
904     hd->next = NULL;
905     hd->slot = boundaryArrayLen;
906
907     return boundaryArrayLen;
908   }
909   for(int i = 0; i < GPU_MEMPOOL_NUM_SLOTS-1; i++){
910     if(memPoolBoundaries[i] < size && size <= memPoolBoundaries[i+1]){
911       return (i+1);
912     }
913   }
914   return -1;
915 }
916
917 void *getBufferFromPool(int pool, int size){
918   Header *ret;
919   if(pool < 0 || pool >= memPoolFreeBufs.length() || memPoolFreeBufs[pool].head == NULL){
920 #ifdef GPU_MEMPOOL_DEBUG
921     printf("(%d) pool %d size: %d, num: %d\n", CmiMyPe(), pool, size, memPoolFreeBufs[pool].num);
922 #endif
923     abort();
924   }
925   else{
926     ret = memPoolFreeBufs[pool].head;
927     memPoolFreeBufs[pool].head = ret->next;
928 #ifdef GPU_MEMPOOL_DEBUG
929     ret->size = size;
930     memPoolFreeBufs[pool].num--;
931 #endif
932     return (void *)(ret+1);
933   }
934   return NULL;
935 }
936
937 void returnBufferToPool(int pool, Header *hd){
938   hd->next = memPoolFreeBufs[pool].head;
939   memPoolFreeBufs[pool].head = hd;
940 #ifdef GPU_MEMPOOL_DEBUG
941   memPoolFreeBufs[pool].num++;
942 #endif
943 }
944
945 void *hapi_poolMalloc(int size){
946   int pool = findPool(size);
947   void *buf = getBufferFromPool(pool, size);
948 #ifdef GPU_MEMPOOL_DEBUG
949   printf("(%d) hapi_malloc size %d pool %d left %d\n", CmiMyPe(), size, pool, memPoolFreeBufs[pool].num);
950 #endif
951   return buf;
952 }
953
954 void hapi_poolFree(void *ptr){
955   Header *hd = ((Header *)ptr)-1;
956   int pool = hd->slot;
957   returnBufferToPool(pool, hd);
958 #ifdef GPU_MEMPOOL_DEBUG
959   int size = hd->size;
960   printf("(%d) hapi_free size %d pool %d left %d\n", CmiMyPe(), size, pool, memPoolFreeBufs[pool].num);
961 #endif
962 }
963
964
965 #endif
966
967 #ifdef GPU_INSTRUMENT_WRS
968 void hapi_initInstrument(int numChares, char types){
969   avgTimes.reserve(numChares);
970   avgTimes.length() = numChares;
971   for(int i = 0; i < numChares; i++){
972     avgTimes[i].reserve(types);
973     avgTimes[i].length() = types;
974   }
975   initialized_instrument = true;
976 }
977
978 bool initializedInstrument(){
979   return initialized_instrument;
980 }
981
982 RequestTimeInfo &hapi_queryInstrument(int chare, char type, char phase){
983   return avgTimes[chare][type][phase];
984 }
985
986 void hapi_clearInstrument(){
987   for(int chare = 0; chare < avgTimes.length(); chare++){
988     for(int type = 0; type < avgTimes[chare].length(); type++){
989       for(int phase = 0; phase < avgTimes[chare][type].length(); phase++){
990         avgTimes[chare][type][phase].transferTime = 0.0;
991         avgTimes[chare][type][phase].kernelTime = 0.0;
992         avgTimes[chare][type][phase].cleanupTime = 0.0;
993         avgTimes[chare][type][phase].n = 0;
994       }
995       avgTimes[chare][type].length() = 0;
996     }
997     avgTimes[chare].length() = 0;
998   }
999   avgTimes.length() = 0;
1000   initialized_instrument = false;
1001 }
1002 #endif