Fixed a bug in allocation of GPU buffers. Those buffers which were not
[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 index value is invalid, use an available ID  
198       if (index < 0 || index >= NUM_BUFFERS) {
199         int found = 0; 
200         for (int j=nextBuffer; j<NUM_BUFFERS*2; j++) {
201           if (devBuffers[j] == NULL) {
202             index = j;
203             found = 1; 
204             break;
205           }
206         }
207
208         /* if no index was found, try to search for a value at the
209          * beginning of the system addressed space 
210          */
211         
212         if (!found) {
213           for (int j=NUM_BUFFERS; j<nextBuffer; j++) {
214             if (devBuffers[j] == NULL) {        
215               index = j;
216               found = 1; 
217               break;
218             }
219           }
220         }
221
222         /* if no index was found, print an error */
223         if (!found) {
224           printf("Error: devBuffers is full \n");
225         }
226
227         nextBuffer = index+1; 
228         if (nextBuffer == NUM_BUFFERS * 2) {
229           nextBuffer = NUM_BUFFERS; 
230         }
231         
232         bufferInfo[i].bufferID = index; 
233
234       }      
235       
236       // allocate if the buffer for the corresponding index is NULL 
237       if (devBuffers[index] == NULL) {
238 #ifdef GPU_PRINT_BUFFER_ALLOCATE
239         double mil = 1e6;
240         printf("*** ALLOCATE buffer 0x%x size %f mb\n", devBuffers[index], 1.0*size/mil);
241 #endif
242
243         CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **) &devBuffers[index], size));
244 #ifdef GPU_DEBUG
245         printf("buffer %d allocated at time %.2f size: %d error string: %s\n", 
246                index, cutGetTimerValue(timerHandle), size, 
247                cudaGetErrorString( cudaGetLastError() ) );
248 #endif
249       }
250     }
251   }
252 }
253
254
255 /* setupData
256  *  copy data to the GPU before kernel execution 
257  */
258 void setupData(workRequest *wr) {
259   dataInfo *bufferInfo = wr->bufferInfo; 
260
261   if (bufferInfo != NULL) {
262     for (int i=0; i<wr->nBuffers; i++) {
263       int index = bufferInfo[i].bufferID; 
264       int size = bufferInfo[i].size; 
265       hostBuffers[index] = bufferInfo[i].hostBuffer; 
266       
267       /* allocate if the buffer for the corresponding index is NULL */
268       /*
269       if (devBuffers[index] == NULL) {
270         CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **) &devBuffers[index], size));
271 #ifdef GPU_DEBUG
272         printf("buffer %d allocated %.2f\n", index,
273                cutGetTimerValue(timerHandle)); 
274 #endif
275       }
276       */
277       
278       if (bufferInfo[i].transferToDevice) {
279         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpyAsync(devBuffers[index], 
280           hostBuffers[index], size, cudaMemcpyHostToDevice, data_in_stream));
281 #ifdef GPU_DEBUG
282         printf("transferToDevice bufId: %d at time %.2f size: %d " 
283                "error string: %s\n", index, cutGetTimerValue(timerHandle), 
284                size, cudaGetErrorString( cudaGetLastError() )); 
285 #endif  
286         /*
287         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(devBuffers[index], 
288           hostBuffers[index], size, cudaMemcpyHostToDevice));
289         */
290
291       }
292     }
293   }
294
295
296 /* copybackData
297  *  transfer data from the GPU to the CPU after a work request is done 
298  */ 
299 void copybackData(workRequest *wr) {
300   dataInfo *bufferInfo = wr->bufferInfo; 
301
302   if (bufferInfo != NULL) {
303     int nBuffers = wr->nBuffers; 
304     
305     for (int i=0; i<nBuffers; i++) {
306       int index = bufferInfo[i].bufferID; 
307       int size = bufferInfo[i].size; 
308       
309       if (bufferInfo[i].transferFromDevice) {
310 #ifdef GPU_DEBUG
311         printf("transferFromDevice: %d at time %.2f size: %d "
312                "error string: %s\n", index, cutGetTimerValue(timerHandle), 
313                size, cudaGetErrorString( cudaGetLastError() )); 
314 #endif
315         
316         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpyAsync(hostBuffers[index], 
317           devBuffers[index], size, cudaMemcpyDeviceToHost,
318           data_out_stream));
319         
320         /*
321         CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(hostBuffers[index], 
322           devBuffers[index], size, cudaMemcpyDeviceToHost));
323         */
324       }
325     }     
326   }
327 }
328
329 /* frees GPU memory for buffers specified by the user; also frees the
330  *  work request's bufferInfo array
331  */
332 void freeMemory(workRequest *wr) {
333   dataInfo *bufferInfo = wr->bufferInfo;   
334   int nBuffers = wr->nBuffers; 
335   if (bufferInfo != NULL) {
336     for (int i=0; i<nBuffers; i++) {    
337       int index = bufferInfo[i].bufferID; 
338       if (bufferInfo[i].freeBuffer) {
339 #ifdef GPU_PRINT_BUFFER_ALLOCATE
340         printf("*** FREE buffer 0x%x\n", devBuffers[index]);
341 #endif
342
343 #ifdef GPU_DEBUG
344         printf("buffer %d freed at time %.2f error string: %s\n", 
345                index, cutGetTimerValue(timerHandle),  
346                cudaGetErrorString( cudaGetLastError() ));
347 #endif 
348         CUDA_SAFE_CALL_NO_SYNC(cudaFree(devBuffers[index])); 
349         devBuffers[index] = NULL; 
350       }
351     }
352     free(bufferInfo); 
353   }
354 }
355
356 /* kernelSelect
357  * a switch statement defined by the user to allow the library to execute
358  * the correct kernel 
359  */ 
360 void kernelSelect(workRequest *wr);
361
362 /* initHybridAPI
363  *   initializes the work request queue, host/device buffer pointer
364  *   arrays, and CUDA streams
365  */
366 void initHybridAPI(int myPe) {
367
368   int deviceCount;
369   cudaGetDeviceCount(&deviceCount);
370
371   cudaSetDevice(myPe % deviceCount); 
372
373   initWRqueue(&wrQueue);
374
375   /* allocate host/device buffers array (both user and
376      system-addressed) */
377   hostBuffers = (void **) malloc(NUM_BUFFERS * 2 * sizeof(void *)); 
378   devBuffers = (void **) malloc(NUM_BUFFERS * 2 * sizeof(void *)); 
379
380   /* initialize device array to NULL */ 
381   for (int i=0; i<NUM_BUFFERS*2; i++) {
382     devBuffers[i] = NULL; 
383   }
384   
385   CUDA_SAFE_CALL_NO_SYNC(cudaStreamCreate(&kernel_stream)); 
386   CUDA_SAFE_CALL_NO_SYNC(cudaStreamCreate(&data_in_stream)); 
387   CUDA_SAFE_CALL_NO_SYNC(cudaStreamCreate(&data_out_stream)); 
388
389 #ifdef GPU_PROFILE
390   CUT_SAFE_CALL(cutCreateTimer(&timerHandle));
391   CUT_SAFE_CALL(cutStartTimer(timerHandle));
392 #endif
393
394   nextBuffer = NUM_BUFFERS;  
395
396 #ifdef GPU_TRACE
397   traceRegisterUserEvent("GPU Memory Setup", GPU_MEM_SETUP);
398   traceRegisterUserEvent("GPU Kernel Execution", GPU_KERNEL_EXEC);
399   traceRegisterUserEvent("GPU Memory Cleanup", GPU_MEM_CLEANUP);
400 #endif
401 }
402
403 /* gpuProgressFn
404  *  called periodically to monitor work request progress, and perform
405  *  the prefetch of data for a subsequent work request
406  */
407 void gpuProgressFn() {
408   if (wrQueue == NULL) {
409     printf("Error: work request queue not initialized\n"); 
410     return; 
411   }
412   if (isEmpty(wrQueue)) {
413     //    flushPinnedMemQueue();    
414     return;
415   } 
416   int returnVal; 
417   workRequest *head = firstElement(wrQueue); 
418   workRequest *second = secondElement(wrQueue);
419   workRequest *third = thirdElement(wrQueue); 
420   if (head->state == QUEUED) {
421 #ifdef GPU_PROFILE
422     gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
423     gpuEvents[timeIndex].eventType = DATA_SETUP; 
424     gpuEvents[timeIndex].ID = head->id; 
425     dataSetupIndex = timeIndex; 
426 #ifdef GPU_TRACE
427     gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
428     gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
429 #endif
430     timeIndex++; 
431 #endif
432     allocateBuffers(head); 
433     setupData(head); 
434     head->state = TRANSFERRING_IN; 
435   }  
436   if (head->state == TRANSFERRING_IN) {
437     if ((returnVal = cudaStreamQuery(data_in_stream)) == cudaSuccess) {
438 #ifdef GPU_PROFILE
439       gpuEvents[dataSetupIndex].endTime = cutGetTimerValue(timerHandle);
440 #ifdef GPU_TRACE
441       gpuEvents[dataSetupIndex].cmiendTime = CmiWallTimer();
442       traceUserBracketEvent(gpuEvents[dataSetupIndex].stage, 
443                             gpuEvents[dataSetupIndex].cmistartTime, 
444                             gpuEvents[dataSetupIndex].cmiendTime); 
445 #endif 
446 #endif
447       if (second != NULL /*&& (second->state == QUEUED)*/) {
448         allocateBuffers(second); 
449       }
450 #ifdef GPU_PROFILE
451       gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
452       gpuEvents[timeIndex].eventType = KERNEL_EXECUTION; 
453       gpuEvents[timeIndex].ID = head->id; 
454       runningKernelIndex = timeIndex; 
455 #ifdef GPU_TRACE
456       gpuEvents[timeIndex].stage = GPU_KERNEL_EXEC; 
457       gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
458 #endif
459       timeIndex++; 
460 #endif
461       //flushPinnedMemQueue();
462       kernelSelect(head); 
463       head->state = EXECUTING; 
464       if (second != NULL) {
465 #ifdef GPU_PROFILE
466         gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
467         gpuEvents[timeIndex].eventType = DATA_SETUP; 
468         gpuEvents[timeIndex].ID = second->id; 
469         dataSetupIndex = timeIndex; 
470 #ifdef GPU_TRACE
471         gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
472         gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
473 #endif
474         timeIndex++; 
475 #endif
476         setupData(second); 
477         second->state = TRANSFERRING_IN;
478       }
479     }
480       /*
481 #ifdef GPU_DEBUG
482       printf("Querying memory stream returned: %d %.2f\n", returnVal, 
483              cutGetTimerValue(timerHandle));
484 #endif  
485       */
486   }
487   if (head->state == EXECUTING) {
488     if ((returnVal = cudaStreamQuery(kernel_stream)) == cudaSuccess) {
489 #ifdef GPU_PROFILE
490       gpuEvents[runningKernelIndex].endTime = cutGetTimerValue(timerHandle); 
491 #ifdef GPU_TRACE
492       gpuEvents[runningKernelIndex].cmiendTime = CmiWallTimer();
493       traceUserBracketEvent(gpuEvents[runningKernelIndex].stage, 
494                             gpuEvents[runningKernelIndex].cmistartTime, 
495                             gpuEvents[runningKernelIndex].cmiendTime); 
496 #endif
497 #endif
498       if (second != NULL && second->state == QUEUED) {
499 #ifdef GPU_PROFILE
500         gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
501         gpuEvents[timeIndex].eventType = DATA_SETUP; 
502         gpuEvents[timeIndex].ID = second->id; 
503         dataSetupIndex = timeIndex; 
504 #ifdef GPU_TRACE
505         gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
506         gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
507 #endif
508         timeIndex++; 
509 #endif
510         allocateBuffers(second); 
511         setupData(second); 
512         second->state = TRANSFERRING_IN;        
513       } 
514       if (second != NULL && second->state == TRANSFERRING_IN) {
515         if (cudaStreamQuery(data_in_stream) == cudaSuccess) {
516 #ifdef GPU_PROFILE
517           gpuEvents[dataSetupIndex].endTime = cutGetTimerValue(timerHandle); 
518 #ifdef GPU_TRACE
519           gpuEvents[dataSetupIndex].cmiendTime = CmiWallTimer();
520           traceUserBracketEvent(gpuEvents[dataSetupIndex].stage, 
521                                 gpuEvents[dataSetupIndex].cmistartTime, 
522                                 gpuEvents[dataSetupIndex].cmiendTime); 
523 #endif
524 #endif
525           if (third != NULL /*&& (third->state == QUEUED)*/) {
526             allocateBuffers(third); 
527           }
528 #ifdef GPU_PROFILE
529           gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
530           gpuEvents[timeIndex].eventType = KERNEL_EXECUTION; 
531           gpuEvents[timeIndex].ID = second->id; 
532           runningKernelIndex = timeIndex; 
533 #ifdef GPU_TRACE
534           gpuEvents[timeIndex].stage = GPU_KERNEL_EXEC; 
535           gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
536 #endif
537           timeIndex++; 
538 #endif
539           //        flushPinnedMemQueue();          
540           kernelSelect(second); 
541           second->state = EXECUTING; 
542           if (third != NULL) {
543 #ifdef GPU_PROFILE
544             gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
545             gpuEvents[timeIndex].eventType = DATA_SETUP; 
546             gpuEvents[timeIndex].ID = third->id; 
547             dataSetupIndex = timeIndex; 
548 #ifdef GPU_TRACE
549             gpuEvents[timeIndex].stage = GPU_MEM_SETUP; 
550             gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
551 #endif
552             timeIndex++; 
553 #endif
554             setupData(third); 
555             third->state = TRANSFERRING_IN;     
556           }
557         }
558       }
559 #ifdef GPU_PROFILE
560       gpuEvents[timeIndex].startTime = cutGetTimerValue(timerHandle); 
561       gpuEvents[timeIndex].eventType = DATA_CLEANUP; 
562       gpuEvents[timeIndex].ID = head->id; 
563       dataCleanupIndex = timeIndex;     
564 #ifdef GPU_TRACE
565       gpuEvents[timeIndex].stage = GPU_MEM_CLEANUP; 
566       gpuEvents[timeIndex].cmistartTime = CmiWallTimer();
567 #endif
568       timeIndex++; 
569 #endif
570       copybackData(head);
571       head->state = TRANSFERRING_OUT;
572     }
573       /*
574 #ifdef GPU_DEBUG
575       printf("Querying kernel completion returned: %d %.2f\n", returnVal,
576              cutGetTimerValue(timerHandle));
577 #endif  
578       */
579   }
580   if (head->state == TRANSFERRING_OUT) {
581     if (cudaStreamQuery(data_out_stream) == cudaSuccess) {
582       freeMemory(head); 
583 #ifdef GPU_PROFILE
584       gpuEvents[dataCleanupIndex].endTime = cutGetTimerValue(timerHandle);
585 #ifdef GPU_TRACE
586       gpuEvents[dataCleanupIndex].cmiendTime = CmiWallTimer();
587       traceUserBracketEvent(gpuEvents[dataCleanupIndex].stage, 
588                             gpuEvents[dataCleanupIndex].cmistartTime, 
589                             gpuEvents[dataCleanupIndex].cmiendTime); 
590 #endif
591 #endif
592       dequeue(wrQueue);
593       CUDACallbackManager(head->callbackFn);
594     }
595   }
596   
597 }
598
599 /* exitHybridAPI
600  *  cleans up and deletes memory allocated for the queue and the CUDA streams
601  */
602 void exitHybridAPI() {
603   deleteWRqueue(wrQueue); 
604   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(kernel_stream)); 
605   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(data_in_stream)); 
606   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(data_out_stream)); 
607
608 #ifdef GPU_PROFILE
609   for (int i=0; i<timeIndex; i++) {
610     switch (gpuEvents[i].eventType) {
611     case DATA_SETUP:
612       printf("Kernel %d data setup", gpuEvents[i].ID); 
613       break;
614     case DATA_CLEANUP:
615       printf("Kernel %d data cleanup", gpuEvents[i].ID); 
616       break; 
617     case KERNEL_EXECUTION:
618       printf("Kernel %d execution", gpuEvents[i].ID); 
619       break;
620     default:
621       printf("Error, invalid timer identifier\n"); 
622     }
623     printf(" %.2f:%.2f\n", gpuEvents[i].startTime-gpuEvents[0].startTime, gpuEvents[i].endTime-gpuEvents[0].startTime); 
624   }
625
626   CUT_SAFE_CALL(cutStopTimer(timerHandle));
627   CUT_SAFE_CALL(cutDeleteTimer(timerHandle));  
628
629 #endif
630 }