/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_cuda/hst_cuda_bp_kepler.h

  • Committer: Suren A. Chilingaryan
  • Date: 2017-09-28 10:34:47 UTC
  • Revision ID: csa@suren.me-20170928103447-dggjgnuxmymgew2a
Quality fixes (tex-based)

Show diffs side-by-side

added added

removed removed

Lines of Context:
hst_cuda/hst_cuda_bp_kepler.h
1
 
#if SLICE_BLOCK == 4
2
 
# define shfl_sum(val, k) ({ \
3
 
    val.x += __shfl_xor(val.x, k, 16); \
4
 
    val.y += __shfl_xor(val.y, k, 16); \
5
 
    val.z += __shfl_xor(val.z, k, 16); \
6
 
    val.w += __shfl_xor(val.w, k, 16); \
7
 
})
8
 
#elif SLICE_BLOCK == 2
9
 
# define shfl_sum(val, k) ({ \
10
 
    val.x += __shfl_xor(val.x, k, 16); \
11
 
    val.y += __shfl_xor(val.y, k, 16); \
12
 
})
13
 
#else
14
 
# define shfl_sum(val, k) ({ \
15
 
    val += __shfl_xor(val, k, 16); \
16
 
})
17
 
#endif
18
 
 
19
 
 
20
1
#if defined(HST_NEWTEX4)
21
2
# define BIN_DIM 64
22
3
# define BIN_STEPS 2
39
20
# endif
40
21
#endif
41
22
 
 
23
#if SLICE_BLOCK == 4
 
24
# define shfl_sum(val, k) ({ \
 
25
    val.x += __shfl_xor(val.x, k, PROJ_STEP); \
 
26
    val.y += __shfl_xor(val.y, k, PROJ_STEP); \
 
27
    val.z += __shfl_xor(val.z, k, PROJ_STEP); \
 
28
    val.w += __shfl_xor(val.w, k, PROJ_STEP); \
 
29
})
 
30
#elif SLICE_BLOCK == 2
 
31
# define shfl_sum(val, k) ({ \
 
32
    val.x += __shfl_xor(val.x, k, PROJ_STEP); \
 
33
    val.y += __shfl_xor(val.y, k, PROJ_STEP); \
 
34
})
 
35
#else
 
36
# define shfl_sum(val, k) ({ \
 
37
    val += __shfl_xor(val, k, PROJ_STEP); \
 
38
})
 
39
#endif
 
40
 
 
41
 
 
42
 
42
43
#ifdef HYBRID_KEPLER
43
44
__device__
44
45
#else
55
56
hst_kepler_kernel
56
57
#endif
57
58
#if defined(HST_FLOAT_LOOPS)
58
 
(cudaTextureObject_t texptr, float num_proj, int num_bins, vfloat *d_SLICE, float apos_off_x, float apos_off_y,
 
59
(cudaTextureObject_t texptr, const float * __restrict__ g_all, float num_proj, int num_bins, vfloat *d_SLICE, float apos_off_x, float apos_off_y,
59
60
#else
60
 
(cudaTextureObject_t texptr, int num_proj, int num_bins, vfloat *d_SLICE, float apos_off_x, float apos_off_y,
 
61
(cudaTextureObject_t texptr, const float * __restrict__ g_all, int num_proj, int num_bins, vfloat *d_SLICE, float apos_off_x, float apos_off_y,
61
62
#endif
62
63
#ifdef HYBRID_KEPLER
63
64
const int bidx, const int bidy) {
118
119
    const float x = bidx + sidx + apos_off_x;
119
120
    const float y = bidy + sidy + apos_off_y;
120
121
 
 
122
 
121
123
#if defined(HST_FLOAT_LOOPS)
122
124
    for (float projf = proj + 0.5f; projf < num_proj; projf += PROJ_STEP) {
123
125
//        float fidx = projf + exp2(23.f); const int idx = (*(int*)(&fidx)) - 0x4B000000;
150
152
#ifdef HST_SHFL_SUM
151
153
            vfloat val = buf[ridy][ridx];
152
154
# pragma unroll
153
 
            for (int k=PROJ_DIM; k>=1; k/=2) 
 
155
            for (int k=(PROJ_DIM/2); k>=1; k/=2) 
154
156
                shfl_sum(val, k);
155
157
#else // HST_SHFL_SUM
156
158
# pragma unroll