/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_kernels.cl

  • Committer: Suren A. Chilingaryan
  • Date: 2017-12-20 15:28:19 UTC
  • Revision ID: csa@suren.me-20171220152819-byiryz4mmovcg8u2
New remmaping, small adjustment of configuration

Show diffs side-by-side

added added

removed removed

Lines of Context:
hst_opencl/hst_opencl_kernels.cl
33
33
# define SET_SLICE(slice, idx, idy, sum) ({ slice[(idy)*(slice_width) + (idx)] = (sum); })
34
34
#endif
35
35
 
 
36
#define REMAP64                 // No performance impact is measured
 
37
 
36
38
#define HST_CACHE_CONST
37
39
#define HST_GMEM_CONST
38
40
#define CACHE_BLOCK 256
628
630
    slice[slice_width * (idy + 1) + idx + 1] = res.w;
629
631
}
630
632
 
631
 
 
632
633
__kernel void hst_backproject_gpu(const int num_proj,
633
634
    const int num_bins,
634
635
    const float off_x,
643
644
    const int q2x = q2 % 2, q2y = q2 / 2;
644
645
    const int q4 = get_local_id(0) / 4;
645
646
    const int q4x = q4 % 2, q4y = q4 / 2;
 
647
 
 
648
# ifdef REMAP64
 
649
    const int q8 = get_local_id(1) % 4;
 
650
    const int q8x = q8 % 2, q8y = q8 / 2;
 
651
 
 
652
    const int q16 = get_local_id(1) / 4;
 
653
    const int q16x = q16 % 2, q16y = q16 / 2;
 
654
 
 
655
    const int tidx = 8 * q16x + 4 * q8x + 2 * q4x + q2x;
 
656
    const int tidy = 8 * q16y + 4 * q8y + 2 * q4y + q2y;
 
657
# else
646
658
    const int q16 = get_local_id(1);
647
659
    const int q16x = q16 % 4, q16y = q16 / 4;
648
660
 
649
661
    const int tidx = 4 * q16x + 2 * q4x + q2x;
650
662
    const int tidy = 4 * q16y + 2 * q4y + q2y;
 
663
# endif
651
664
 
652
665
    const int idx = get_group_id(0) * get_local_size(0) + tidx;
653
666
    const int idy = get_group_id(1) * get_local_size(1) + tidy;