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