/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 docs/optimizations/gcn/hst_opencl_dma_8x8_6ppt.cl

  • 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:
docs/optimizations/gcn/hst_opencl_dma_8x8_6ppt.cl
 
1
#undef subh_cache
 
2
#define subh_cache(subh) cache[(int)(subh)]
 
3
 
 
4
#define PPT 6
 
5
 
 
6
__kernel void hst_backproject_test_oversample4_1(const int num_proj, 
 
7
    const int num_bins, 
 
8
    const float off_x, 
 
9
    const float off_y, 
 
10
    __read_only image2d_t sinogram,
 
11
    __global float *slice,
 
12
    __constant float4 *c_all, 
 
13
    __local float *cache
 
14
    ) {
 
15
 
 
16
    const int tidx = get_local_id(0);
 
17
    const int tidy = get_local_id(1);
 
18
 
 
19
    const int bidx = PPT * get_group_id(0) * BLOCK_SIZE;
 
20
    const int bidy = PPT * (get_global_id(1) - get_local_id(1) - get_group_id(1) * BLOCK_SIZE);
 
21
 
 
22
    const float bx = bidx + off_x;
 
23
    const float by = bidy + off_y;
 
24
 
 
25
/*
 
26
    const int sbidx = tidy % 4;
 
27
    const int sbidy = tidy / 4;
 
28
    
 
29
    const int stidx = tidx % 4;
 
30
    const int stidy = tidx / 4;
 
31
 
 
32
    const int sidx = (sbidx * 4 + stidx);
 
33
    const int sidy = (sbidy * 4 + stidy);
 
34
*/
 
35
 
 
36
    const int sidx = tidx;
 
37
    const int sidy = tidy - BLOCK_SIZE;
 
38
 
 
39
 
 
40
    const float sx = sidx;
 
41
    const float sy = sidy;
 
42
 
 
43
    const int idx = bidx + sidx;
 
44
    const int idy = bidy + sidy;
 
45
 
 
46
    const int slice_width = PPT * get_global_size(0);
 
47
 
 
48
    int buf = 0;
 
49
    int offset = 5 * BLOCK_SIZE * BLOCK_SIZE;
 
50
    float4 all = c_all[0];
 
51
    
 
52
    if (tidy < BLOCK_SIZE) {
 
53
        float minh = floor(all.z + bx * all.x - by * all.y + all.w);
 
54
 
 
55
#pragma unroll 5
 
56
        for (int i = 0; i < 5; i++) {
 
57
            int pos = i * BLOCK_SIZE * BLOCK_SIZE + tidy * BLOCK_SIZE + tidx;
 
58
            cache[pos] = read_imagef(sinogram, volumeSampler, (float2)(minh + 0.25f*pos, .5f)).x;
 
59
        }
 
60
    }
 
61
    
 
62
    float res[PPT][PPT] = {0};
 
63
    for (int proj = 0; proj < num_proj; proj++) {
 
64
        if (get_local_id(1) >= BLOCK_SIZE) {
 
65
            all = -4 * all; 
 
66
            all.x = -all.x;
 
67
 
 
68
            float subh = mad(sx, all.x, mad(sy, all.y, all.w));
 
69
 
 
70
 
 
71
#pragma unroll 6
 
72
            for (int i = 0; i < PPT; i++) {
 
73
#pragma unroll 6
 
74
                for (int j = 0; j < PPT; j++) {
 
75
                    res[i][j] += subh_cache(buf * offset + subh + (i * BLOCK_SIZE) * all.y + (j * BLOCK_SIZE) * all.x);
 
76
                }
 
77
            }
 
78
        }
 
79
        
 
80
        buf = !buf;
 
81
        if (proj == num_proj) break;
 
82
 
 
83
        all = c_all[proj + 1];
 
84
        if (tidy < BLOCK_SIZE) {
 
85
            float minh = floor(all.z + bx * all.x - by * all.y + all.w);
 
86
 
 
87
#pragma unroll 5
 
88
            for (int i = 0; i < 5; i++) {
 
89
                int pos = i * BLOCK_SIZE * BLOCK_SIZE + tidy * BLOCK_SIZE + tidx;
 
90
                cache[buf * offset + pos] = read_imagef(sinogram, volumeSampler, (float2)(minh + 0.25f*pos, proj + 1.5f)).x;
 
91
            }
 
92
        }
 
93
 
 
94
        barrier(CLK_LOCAL_MEM_FENCE);
 
95
    }
 
96
 
 
97
    if (get_local_id(1) >= BLOCK_SIZE) {
 
98
#pragma unroll 6
 
99
    for (int i = 0; i < PPT; i++) {
 
100
#pragma unroll 6
 
101
        for (int j = 0; j < PPT; j++) {
 
102
            slice[(idy + BLOCK_SIZE * i) * slice_width +  idx + BLOCK_SIZE * j] = res[i][j];
 
103
        }
 
104
    }
 
105
    }
 
106
}
 
107
 
 
108
 
 
109
__kernel void hst_backproject_test_oversample4_2(const int num_proj, 
 
110
    const int num_bins, 
 
111
    const float off_x, 
 
112
    const float off_y, 
 
113
    __read_only image2d_t sinogram,
 
114
    __global float *slice,
 
115
    __constant float4 *c_all, 
 
116
    __local float *shared
 
117
    ) {
 
118
 
 
119
    const int tidx = get_local_id(0);
 
120
    const int tidy = get_local_id(1);
 
121
 
 
122
    const int sidx = tidx;
 
123
    const int sidy = tidy;
 
124
 
 
125
    const int idx = PPT * get_group_id(0) * BLOCK_SIZE + sidx;
 
126
    const int idy = PPT * (get_global_id(1) - get_local_id(1) - get_group_id(1) * BLOCK_SIZE) + sidy;
 
127
 
 
128
    const float x = idx + off_x;
 
129
    const float y = idy + off_y;
 
130
 
 
131
    const int slice_width = PPT * get_global_size(0);
 
132
 
 
133
    float res[PPT][PPT] = {0};
 
134
    for (int proj = 0; proj < num_proj; proj++) {
 
135
        float4 all = c_all[proj];
 
136
        all.y = - all.y;
 
137
        float subh = mad(x, all.x, mad(y, all.y, all.z));
 
138
 
 
139
#pragma unroll 6
 
140
        for (int i = 0; i < PPT; i++) {
 
141
#pragma unroll 6
 
142
            for (int j = 0; j < PPT; j++) {
 
143
                res[i][j] += read_imagef(sinogram, volumeSampler, (float2)(subh + i * BLOCK_SIZE * all.y + j * BLOCK_SIZE * all.x, proj + .5f)).x;
 
144
            }
 
145
        }
 
146
    }
 
147
 
 
148
#pragma unroll 6
 
149
    for (int i = 0; i < PPT; i++) {
 
150
#pragma unroll 6
 
151
        for (int j = 0; j < PPT; j++) {
 
152
            slice[(idy + BLOCK_SIZE * i) * slice_width +  idx + BLOCK_SIZE * j] = res[i][j];
 
153
        }
 
154
    }
 
155
}
 
156
 
 
157
 
 
158
 
 
159
__kernel void hst_backproject_test_oversample4(const int num_proj, 
 
160
    const int num_bins, 
 
161
    const float off_x, 
 
162
    const float off_y, 
 
163
    __read_only image2d_t sinogram,
 
164
    __global float *slice,
 
165
    __constant float4 *c_all, 
 
166
    __local float *cache) {
 
167
 
 
168
    const int mode = 4;//(get_group_id(1) * get_num_groups(0) + get_group_id(0)) % 8;
 
169
    if (mode < 3) {
 
170
    if (get_local_id(1) < BLOCK_SIZE) {
 
171
        hst_backproject_test_oversample4_2(num_proj, num_bins, off_x, off_y, sinogram, slice, c_all, cache);
 
172
    }
 
173
    } else {
 
174
        hst_backproject_test_oversample4_1(num_proj, num_bins, off_x, off_y, sinogram, slice, c_all, cache);
 
175
    }
 
176
}
 
177