To get this branch, use:
bzr branch
http://suren.me/webbzr/normxcorr/trunk
« back to all changes in this revision
Viewing changes to dict_hw/src/dict_hw.cu
- browse files at revision 29
- compare with another revision
- download diff
- download tarball
- view history from revision 29
- Committer: Suren A. Chilingaryan
- Date: 2010-04-22 13:42:41 UTC
- Revision ID: csa@dside.dyndns.org-20100422134241-fv5m2ufk8n2tc9h5
Implementation of image and fragment modes, support for non-cacheable grids
added
removed
dict_hw/src/dict_hw.cu
20
20
21
21
static int device_number = 0;
22
22
static int devices[MAX_DEVICES];
23
static size_t device_memory = 2147483648;
23
24
24
25
#ifdef DICT_HW_MEASURE_TIMINGS
25
26
static int hardware_detection_time = 0;
26
27
#endif /* DICT_HW_MEASURE_TIMINGS */
27
28
28
29
int dictDetectHardware() {
30
int memory;
29
31
int deviceCount;
30
32
cudaDeviceProp deviceProp;
31
33
41
43
for (int i = 0; i < deviceCount; i++) {
42
44
cudaGetDeviceProperties(&deviceProp, i);
43
45
if ((deviceProp.major > 1)||((deviceProp.major == 1)&&(deviceProp.minor > 2))) {
44
devices[device_number++] = i;
46
memory = deviceProp.totalGlobalMem;
47
if (memory > 268435455) { // 256 MB
48
if (memory < device_memory) device_memory = memory;
49
devices[device_number++] = i;
50
}
45
51
}
46
52
}
47
53
84
90
int time;
85
91
struct timeval tv1, tv2;
86
92
87
int init_time = ctx->time[0]+ctx->time[1]+ctx->time[2]+ctx->time[3]+ctx->time[4]+ctx->time[12]+ctx->time[13];
93
int init_time = ctx->time[0]+ctx->time[1]+ctx->time[2]+ctx->time[3]+ctx->time[12]+ctx->time[13];
88
94
int load_time = ctx->time[14] + ctx->time[15];
89
95
int process_time = ctx->time[5];
96
97
if (ctx->use_cache) init_time += ctx->time[6];
98
90
99
reportMessage("");
91
100
reportMessage("Library timings");
92
101
reportMessage("---------------");
98
107
print_timing(" Setting Current Points", ctx->time[3]);
99
108
print_timing(" Loading Template Image", ctx->time[12]);
100
109
print_timing(" Reducting Template Image", ctx->time[13]);
101
print_timing(" Processing Template Image", ctx->time[4]);
110
if (ctx->use_cache) {
111
print_timing(" Processing Template Image", ctx->time[6]);
112
}
102
113
print_timing("Loading Images", load_time);
103
114
print_timing(" Load", ctx->time[14]);
104
115
print_timing(" Reduction", ctx->time[15]);
105
116
print_timing("Processing Images", process_time);
106
print_timing(" Copy Fragment", ctx->time[6]);
117
if (!ctx->use_cache) {
118
print_timing(" Processing Template Image", ctx->time[6]);
119
}
120
print_timing(" Copy Image ", ctx->time[4]);
107
121
print_timing(" Load Fragment", ctx->time[7]);
108
122
print_timing(" Preprocessing", ctx->time[8]);
109
123
print_timing(" FFT", ctx->time[9]);
173
187
base_blocks = calc_blocks(4 * ps->corr_size + 1, BLOCK_SIZE_1D);
174
188
ps->base_blocks_power = get_power(base_blocks);
175
189
176
err = fftInit(ps);
190
ps->use_cache = 1;
191
err = fftInit(ps, device_memory);
192
if ((err == DICT_ERROR_CUDA_MALLOC)&&(ps->use_cache)) {
193
ps->use_cache = 0;
194
err = fftInit(ps, device_memory);
195
}
177
196
178
197
#ifdef DICT_HW_MEASURE_TIMINGS
179
198
gettimeofday(&tv2, NULL);
180
199
ps->time[1] += (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
181
200
#endif /* DICT_HW_MEASURE_TIMINGS */
182
201
202
if (ps->use_cache) {
203
reportMessage("Caching is enabled");
204
} else {
205
reportMessage("Caching is disabled");
206
}
207
183
208
return err;
184
209
}
185
210
201
226
}
202
227
203
228
int dictSetDimensions(DICTContext ps, int width, int height) {
204
ps->width = width;
205
ps->height = height;
229
if ((width != ps->width)||(height != ps->height)) {
230
ps->width = width;
231
ps->height = height;
232
233
return fftSetupDimensions(ps);
234
}
206
235
207
236
return 0;
208
237
}
287
316
return 0;
288
317
}
289
318
290
291
319
int dictLoadTemplateFragment(DICTContext ps, int icp, int ncp, const unsigned char *img) {
292
320
int err;
293
321
294
#ifdef DICT_HW_MEASURE_TIMINGS
295
struct timeval tv1, tv2;
296
gettimeofday(&tv1, NULL);
297
#endif /* DICT_HW_MEASURE_TIMINGS */
298
299
err = fftLoadBaseFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), img);
300
301
#ifdef DICT_HW_MEASURE_TIMINGS
302
gettimeofday(&tv2, NULL);
303
ps->time[4] += (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
304
#endif /* DICT_HW_MEASURE_TIMINGS */
322
err = fftLoadBaseImage(ps, img);
323
if (!err) err = fftLoadBaseFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), img);
305
324
306
325
return err;
307
326
}
309
328
int dictLoadTemplateImage(DICTContext ps, const unsigned char *img, int width, int height) {
310
329
int err;
311
330
331
int icp;
332
312
333
#ifdef DICT_HW_MEASURE_TIMINGS
313
334
struct timeval tv1, tv2;
314
335
gettimeofday(&tv1, NULL);
315
336
#endif /* DICT_HW_MEASURE_TIMINGS */
316
337
338
dictSetDimensions(ps, width, height);
339
340
/*
317
341
ps->width = width;
318
342
ps->height = height;
319
343
323
347
int base_size = 4 * ps->corr_size + 1;
324
348
int base_size2 = base_size * base_size;
325
349
326
if (width * height > ps->ncp * size2) {
350
if ((ps->use_cache)&&(width * height > ps->ncp * size2)) {
327
351
ps->mode = 0;
328
352
} else {
329
353
ps->mode = 1;
330
354
}
331
355
332
// if not enoguh space for caching enable anyway ?
333
if (width * height > ps->ncp * base_size2) {
356
if ((ps->use_cache)&&(width * height > ps->ncp * base_size2)) {
334
357
ps->base_mode = 0;
335
358
} else {
336
359
ps->base_mode = 1;
337
if (!ps->mode) {
338
ps->minx = 0;
339
ps->maxx = width - 1;
340
ps->miny = 0;
341
ps->maxy = height - 1;
342
}
343
}
344
345
for (int icp = 0; icp < ps->ncp; icp+=CP_BLOCK) {
346
err = fftLoadBaseFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), img);
347
if (err) break;
348
}
349
350
if ((ps->base_mode)&&(!ps->mode)) {
351
// Correcting difference of area size between base and data images
352
ps->minx += ps->corr_size;
353
ps->miny += ps->corr_size;
354
ps->maxx -= ps->corr_size;
355
ps->maxy -= ps->corr_size;
356
357
width = ceil(ps->maxx) - floor(ps->minx);
358
height = ceil(ps->maxy) - floor(ps->miny);
359
360
// printf("%ux%u=%u %u\n", width, height, width*height, ps->ncp * size2);
361
if (width * height < ps->ncp * size2) {
362
ps->mode = 1;
363
}
360
}
361
*/
362
363
err = fftLoadBaseImage(ps, img);
364
365
if (ps->use_cache) {
366
for (icp = 0; (icp < ps->ncp)&&(!err); icp+=CP_BLOCK) {
367
err = fftLoadBaseFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), img);
368
}
364
369
}
365
370
366
371
#ifdef DICT_HW_MEASURE_TIMINGS
367
372
gettimeofday(&tv2, NULL);
368
ps->time[4] += (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
373
ps->time[6] += (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
369
374
#endif /* DICT_HW_MEASURE_TIMINGS */
370
375
376
/*
371
377
if (ps->mode) {
372
reportMessage("Running in the image mode");
378
if (ps->use_cache) {
379
reportMessage("Running in the image mode, caching enabled");
380
} else {
381
reportMessage("Running in the image mode, caching disabled");
382
}
373
383
} else {
374
reportMessage("Running in the fragment mode");
384
if (ps->base_mode) {
385
if (ps->use_cache) {
386
reportMessage("Running in the image/fragment mode, caching enabled");
387
} else {
388
reportMessage("Running in the image/fragment mode, caching disabled");
389
}
390
} else {
391
if (ps->use_cache) {
392
reportMessage("Running in the fragment mode, caching enabled");
393
} else {
394
reportMessage("Running in the fragment mode, caching disabled");
395
}
396
}
375
397
}
398
*/
376
399
377
return 0;
400
return err;
378
401
}
379
402
380
403
381
404
int dictLoadFragment(DICTContext ps, int icp, int ncp, const unsigned char *input) {
382
405
int err;
406
int load_base = !ps->use_cache;
383
407
cudaStream_t stream = NULL;
384
408
385
409
#ifdef DICT_HW_MEASURE_TIMINGS
388
412
gettimeofday(&tv1, NULL);
389
413
#endif /* DICT_HW_MEASURE_TIMINGS */
390
414
391
err = fftCopyFragment(ps, icp, ncp, input);
392
if (err) return err;
415
if (load_base) {
416
// We are enforcing image mode if the cache is disabled
417
err = fftLoadBaseFragment(ps, icp, ncp, NULL);
418
if (err) return err;
419
}
393
420
394
421
#ifdef DICT_HW_MEASURE_TIMINGS
395
422
gettimeofday(&tv2, NULL);
396
423
time[1] = (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
397
424
#endif /* DICT_HW_MEASURE_TIMINGS */
398
425
399
err = fftLoadFragment(ps, icp, ncp, input, stream);
400
if (err) return err;
426
err = fftLoadFragment(ps, icp, ncp, input);
427
if (err) return err;
401
428
402
429
#ifdef DICT_HW_MEASURE_TIMINGS
403
430
gettimeofday(&tv2, NULL);
404
431
time[2] = (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
405
432
#endif /* DICT_HW_MEASURE_TIMINGS */
406
433
407
err = fftPreprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
408
if (err) return err;
434
err = fftPreprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
435
if (err) return err;
409
436
410
437
#ifdef DICT_HW_MEASURE_TIMINGS
411
438
gettimeofday(&tv2, NULL);
412
439
time[3] = (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
413
440
#endif /* DICT_HW_MEASURE_TIMINGS */
414
441
415
err = fftProcessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
416
if (err) return err;
442
err = fftProcessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
443
if (err) return err;
417
444
418
445
#ifdef DICT_HW_MEASURE_TIMINGS
419
446
gettimeofday(&tv2, NULL);
420
447
time[4] = (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
421
448
#endif /* DICT_HW_MEASURE_TIMINGS */
422
449
423
err = fftPostprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
424
if (err) return err;
450
err = fftPostprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
451
if (err) return err;
425
452
426
453
#ifdef DICT_HW_MEASURE_TIMINGS
427
454
gettimeofday(&tv2, NULL);
432
459
time[i] -= time[i-1];
433
460
}
434
461
435
time[0] = (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
436
437
for (int i = 0; i < 6; i++) {
462
for (int i = 1; i < 6; i++) {
438
463
ps->time[i+5] += time[i];
439
464
}
440
465
#endif /* DICT_HW_MEASURE_TIMINGS */
447
472
int err;
448
473
int ncp = ps->ncp;
449
474
450
/*
451
475
#ifdef DICT_HW_MEASURE_TIMINGS
452
int time[16];
453
476
struct timeval tv1, tv2;
454
477
gettimeofday(&tv1, NULL);
478
#endif /* DICT_HW_MEASURE_TIMINGS */
479
480
err = fftLoadImage(ps, img);
481
if (err) return err;
482
483
#ifdef DICT_HW_MEASURE_TIMINGS
455
484
gettimeofday(&tv2, NULL);
456
time[0] = (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
457
printf("Pre: %li, Nope: %li, Comp: %li, Post: %li, Copy/Load: %li, -: %li\n", time[0], time[1], time[2], time[3], time[4], time[5]);
458
459
cudaStream_t stream[2];
460
for (int i = 0; i < 2; ++i) {
461
cudaStreamCreate(&stream[i]);
462
}
463
for (int i = 0; i < 2; ++i) {
464
cudaStreamDestroy(stream[i]);
465
}
466
#endif
467
*/
485
ps->time[4] += (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
486
#endif /* DICT_HW_MEASURE_TIMINGS */
468
487
469
488
for (int icp = 0; icp < ncp; icp+=CP_BLOCK) {
470
489
err = dictLoadFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), img);
471
490
if (err) return err;
472
491
}
473
492
474
/*
475
err = fftCopyFragment(ps, icp, ncp, input);
476
if (err) return err;
477
478
err = fftLoadFragment(ps, icp, ncp, input, stream);
479
if (err) return err;
480
481
err = fftPreprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
482
if (err) return err;
483
484
err = fftProcessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
485
if (err) return err;
486
487
err = fftPostprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream);
488
if (err) return err;
489
}
490
*/
491
492
493
/*
494
err = fftCopyFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), input, stream[0], NULL);
495
err = fftLoadFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), input, stream[0], NULL);
496
for (int i = 0; icp < ps->ncp; icp+=CP_BLOCK,i++) {
497
err = fftPreprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream[i%2], NULL);
498
err = fftComputeFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream[i%2], NULL);
499
err = fftPostprocessFragment(ps, icp, min2(CP_BLOCK, ps->ncp - icp), stream[i%2], NULL);
500
if (next_icp < ps->ncp) {
501
err = fftCopyFragment(ps, next_icp, min2(CP_BLOCK, ps->ncp - next_icp), input, stream[(i+1)%2], NULL);
502
err = fftLoadFragment(ps, next_icp, min2(CP_BLOCK, ps->ncp - next_icp), input, stream[(i+1)%2], NULL);
503
}
504
}
505
506
*/
493
#ifdef DICT_HW_MEASURE_TIMINGS
494
gettimeofday(&tv2, NULL);
495
ps->time[5] += (tv2.tv_sec - tv1.tv_sec) * 1000000 + (tv2.tv_usec - tv1.tv_usec);
496
#endif /* DICT_HW_MEASURE_TIMINGS */
497
507
498
return 0;
508
499
}
509
500
Loggerhead is a web-based interface for Breezy
Version: 1.20.0