/tomo/pyhst

To get this branch, use:
bzr branch http://darksoft.org/webbzr/tomo/pyhst

« back to all changes in this revision

Viewing changes to hst_opencl/hst_opencl.c

  • Committer: Suren A. Chilingaryan
  • Date: 2012-05-10 15:06:33 UTC
  • Revision ID: csa@dside.dyndns.org-20120510150633-56gdy6t3tflz2gab
OpenCL clean-up

Show diffs side-by-side

added added

removed removed

Lines of Context:
hst_opencl/hst_opencl.c
39
39
#include "hst_opencl_extensions.h"
40
40
 
41
41
#define PARALLEL_PER_GPU 1
42
 
#define SHARED_CONTEXT
 
42
 
 
43
#ifdef HW_IGNORE_OLD_HARDWARE
 
44
# define SHARED_CONTEXT
 
45
#endif /* HW_IGNORE_OLD_HARDWARE */
 
46
 
43
47
//#define BLOCKED_MEMCPY
44
48
 
45
49
 
49
53
 
50
54
#define HST_CHECK_MEMORY 512    // minimal size in megabytes
51
55
#define HST_OPENCL_CONST_HACK
52
 
#define HST_OPENCL_BUILD_FLAGS NULL //"-cl-denorms-are-zero -cl-mad-enable -cl-no-signed-zeros -cl-unsafe-math-optimizations -cl-finite-math-only"
 
56
#define HST_OPENCL_BUILD_FLAGS NULL 
 
57
    //"-cl-denorms-are-zero -cl-mad-enable -cl-no-signed-zeros -cl-unsafe-math-optimizations -cl-finite-math-only" 
 
58
    //"-cl-nv-maxrregcount=48"
53
59
//#define HST_OPENCL_LOAD_PTX "/tmp/hst_opencl.ptx"
54
60
//#define HST_OPENCL_SAVE_PTX "/tmp/hst_opencl.ptx"
55
61
#define MAX_PROJECTIONS 4000
158
164
    float off_y;
159
165
    
160
166
    int ppt;            // point per thread
 
167
    int bp_block;
 
168
    int extra_wraps;    // multiplier
161
169
 
162
170
    int blocked_mode;
163
171
    int blocked_fft;
347
355
        ctx->arch = HST_DEVICE_INTEL_CPU;
348
356
    } else ctx->arch = HST_DEVICE_UNKNOWN;
349
357
    
350
 
 
 
358
    ctx->bp_block = BLOCK_SIZE;
 
359
    
351
360
#ifdef HST_MULTI_POINT_RECONSTRUCTION
352
361
    switch (ctx->arch) {
353
 
        case HST_DEVICE_FERMI:
354
362
        case HST_DEVICE_KEPLER:
 
363
            ctx->ppt = 6;
 
364
            ctx->bp_block = 8;
 
365
            break;
355
366
        case HST_DEVICE_CYPRESS:
 
367
            if (setup->oversampling > 1) {
 
368
                ctx->ppt = 8;
 
369
                ctx->bp_block = 8;
 
370
            } else if (setup->oversampling == 0) {
 
371
                ctx->ppt = 6;
 
372
            }
 
373
            break;
356
374
        case HST_DEVICE_GCN:
 
375
        case HST_DEVICE_FERMI:
357
376
            ctx->ppt = 2;
358
377
            break;
359
378
        default:
363
382
    ctx->ppt = 1;
364
383
#endif /* HST_MULTI_POINT_RECONSTRUCTION */
365
384
 
 
385
    ctx->extra_wraps = 1;
 
386
 
366
387
 
367
388
    if (num_proj%2) num_proj++;
368
389
 
378
399
    dim_proj = 2 * BLOCK_SIZE * hw_calc_blocks(num_proj, 2 * BLOCK_SIZE, NULL);
379
400
 
380
401
#ifdef SHARED_CONTEXT
381
 
    ctx->clctx = clctx;
382
 
#else /* SHARED_CONTEXT */
383
 
    ctx->clctx = clCreateContext(0, 1, &ctx->device_id, NULL, NULL, &err);
384
 
    CL_CHECK_ERROR(err);
 
402
    if (clctx) ctx->clctx = clctx;
 
403
    else {
 
404
#endif /* SHARED_CONTEXT */
 
405
        ctx->clctx = clCreateContext(0, 1, &ctx->device_id, NULL, NULL, &err);
 
406
        CL_CHECK_ERROR(err);
 
407
#ifdef SHARED_CONTEXT
 
408
    }
385
409
#endif /* SHARED_CONTEXT */
386
410
#ifdef PYHST_MEASURE_TIMINGS
387
411
    flags = CL_QUEUE_PROFILING_ENABLE;
409
433
        all[i].x = setup->cos_s[i];
410
434
        all[i].y = setup->sin_s[i];
411
435
        all[i].z = setup->axis_position_corr_s[i];
412
 
        all[i].w = floor(ctx->ppt * MIN(MIN(BLOCK_SIZE * all[i].x, - BLOCK_SIZE * all[i].y), MIN(0., BLOCK_SIZE * all[i].x - BLOCK_SIZE * all[i].y)));
 
436
        all[i].w = floor(ctx->ppt * ctx->bp_block * MIN(MIN(all[i].x, - all[i].y), MIN(0., all[i].x - all[i].y)));
413
437
    }
414
438
    if (setup->num_projections < dim_proj) memset(all + setup->num_projections, 0, (dim_proj - setup->num_projections) * sizeof(cl_float4));
415
439
    ctx->unified_table = clCreateBuffer(ctx->clctx, CL_MEM_READ_ONLY, dim_proj * sizeof(cl_float4), NULL, &err);
510
534
    }
511
535
#endif /* HST_OPENCL_CONST_HACK */
512
536
 
 
537
 
513
538
#ifdef PYHST_MEASURE_TIMINGS
514
539
    g_timer_stop(ctx->init_timer);
515
540
#endif /* PYHST_MEASURE_TIMINGS */
811
836
        slice = ctx->hostres_ptr;
812
837
    }
813
838
 
 
839
    local_work_size[0] = ctx->bp_block;
 
840
    local_work_size[1] = ctx->bp_block * ctx->extra_wraps;
814
841
 
815
842
    if ((ctx->blocked_mode)&&(have_opencl11)&&(!blocking)) {
816
843
        size_t last_batch, batch;
817
 
        size_t last_batch_size, batch_size = ctx->ppt * ctx->blocked_mode * BLOCK_SIZE;
 
844
        size_t last_batch_size, batch_size = ctx->ppt * ctx->blocked_mode * ctx->bp_block;
818
845
        size_t bp_dim[3] = {0};
819
846
        size_t bp_offset[3] = {0};
820
847
        cl_event bp_events[1 + num_y / batch_size];
822
849
        cl_event pad_events[1 + num_y / batch_size];
823
850
 
824
851
        bp_dim[0] = ctx->bp_dim[0];
825
 
        bp_dim[1] = ctx->blocked_mode * BLOCK_SIZE;
 
852
        bp_dim[1] = ctx->blocked_mode * ctx->bp_block * ctx->extra_wraps;
826
853
        slice_size[1] = batch_size;
827
854
 
828
855
        for (batch = 0, i = 0; batch < num_y; batch += batch_size, i++) {
831
858
 
832
859
            if ((batch + batch_size) > num_y) {
833
860
                batch_size = num_y - batch;
834
 
                bp_dim[1] = BLOCK_SIZE * hw_calc_blocks(batch_size, ctx->ppt * BLOCK_SIZE, NULL);
 
861
                bp_dim[1] = ctx->bp_block * hw_calc_blocks(batch_size, ctx->ppt * ctx->bp_block, NULL) * ctx->extra_wraps;
835
862
                slice_size[1] = batch_size;
836
863
            }
837
864
 
879
906
        }
880
907
#endif /* PYHST_MEASURE_TIMINGS */
881
908
    } else {
882
 
        CL_CHECK_ERROR(clEnqueueNDRangeKernel(ctx->queue, ctx->kernel[HST_KERNEL_BACKPROJECT], 2, NULL, ctx->bp_dim, local_work_size, 0, NULL, &event));
 
909
        size_t bp_dim[3] = {0};
 
910
        bp_dim[0] = ctx->bp_dim[0];
 
911
        bp_dim[1] = ctx->bp_dim[1] * ctx->extra_wraps;
 
912
 
 
913
        CL_CHECK_ERROR(clEnqueueNDRangeKernel(ctx->queue, ctx->kernel[HST_KERNEL_BACKPROJECT], 2, NULL, bp_dim, local_work_size, 0, NULL, &event));
883
914
        hst_measure_time(&ctx->timer_backprojection, ctx->queue, event);
884
915
 
885
916
#ifdef PYHST_MEASURE_TIMINGS
1168
1199
 
1169
1200
#ifdef SHARED_CONTEXT
1170
1201
    if (g_hst_platform.num_devices > 0) {
1171
 
        int max_pos = 0;
 
1202
        int max_pos = -1;
1172
1203
        cl_ulong max_mem = 0;
 
1204
        char device_name[HST_MAX_NAME_LENGTH + 1];
1173
1205
        
1174
1206
        for (i = 0; i < g_hst_platform.num_devices; i++) {
1175
1207
            cl_ulong device_mem;
1178
1210
            if (err != CL_SUCCESS) continue;
1179
1211
            
1180
1212
            if (device_mem > max_mem) {
 
1213
                if (amd_platform) {
 
1214
                    if (clGetDeviceInfo(g_hst_platform.device_ids[i], CL_DEVICE_NAME, HST_MAX_NAME_LENGTH, device_name, NULL)) continue;
 
1215
                    if (!strcasecmp(device_name, "Cypress")) continue;
 
1216
                }
1181
1217
                max_pos = i;
1182
1218
                max_mem = device_mem;
1183
1219
            }
1184
1220
        }
1185
 
 
1186
 
        clctx = clCreateContext(NULL, g_hst_platform.num_devices, g_hst_platform.device_ids, NULL, NULL, &err);
1187
 
            // find device with maximum memory
1188
 
        if (err == CL_SUCCESS) clqueue = clCreateCommandQueue(clctx, g_hst_platform.device_ids[max_pos], 0, &err);
 
1221
        
 
1222
        if (max_pos >= 0) {
 
1223
            clctx = clCreateContext(NULL, g_hst_platform.num_devices, g_hst_platform.device_ids, NULL, NULL, &err);
 
1224
                // find device with maximum memory
 
1225
            if (err == CL_SUCCESS) clqueue = clCreateCommandQueue(clctx, g_hst_platform.device_ids[max_pos], 0, &err);
 
1226
        }
1189
1227
    }
1190
1228
#endif /* SHARED_CONTEXT */
1191
1229