/tomo/pyhst

#define subh_cache(subh) cache[subh]

    __local float *mcache = shared;

    __local float *cache = shared + BLOCK_SIZE * BLOCK_SIZE;                                            // cache[BLOCK_SIZE_Y][12 * BLOCK_SIZE_X + 1];

    const int sidx = (sbidx * 8 + stidx);

    const int idx = bidx + sidx;

    const float x = idx + off_x;

    const float y = idy + off_y;

    const int num_blocks = (num_proj / BLOCK_SIZE) +  ((num_proj % BLOCK_SIZE)?1:0);

    

    for (int proj_block = 0; proj_block < num_blocks; proj_block += BLOCK_SIZE) {

        const int proj_offset = proj_block * BLOCK_SIZE;

        float4 all = c_all[proj_offset + BLOCK_SIZE * tidy + tidx];

 

        float minh = floor(all.z + bx * all.x - by * all.y + all.w);

        mcache[BLOCK_SIZE * tidy + tidx] = 193.f * tidx - 4 * minh;

 

        barrier(CLK_LOCAL_MEM_FENCE);

 

        const int sub_blocks = min(BLOCK_SIZE, num_blocks - proj_block);

        

        for (int p = 0; p < sub_blocks; p++) {

            const int proje = proj_offset + p * BLOCK_SIZE;

            const int proj = proje + tidy;

            

            float minh = 193.f * tidy - mcache[p * BLOCK_SIZE + tidy];

            for (int i = 0; i < 12; i++) {

                int pos = i * BLOCK_SIZE + tidx;

                cache[tidy * (12 * BLOCK_SIZE + 1) + pos] = read_imagef(sinogram, volumeSampler, (float2)(0.25f * (minh + pos), proj + 0.5f)).x;

            }

 

            barrier(CLK_LOCAL_MEM_FENCE);

            

            for (int i = 0; i < BLOCK_SIZE; i++) {

//            float4 all = -4 * c_all[proje + i];

//            float subh = mad(193.f, i, mad(sidx, all.x, mad(sidy, all.y, all.w)));

              float4 all = c_all[proje + i]; all.y = -all.y;

//              float minh = 193 * i - 4 * floor(all.z + bx * all.x + by * all.y + all.w);

 

                float minh = mcache[p * BLOCK_SIZE + i];

 

                float subh = mad(4,  mad(x, all.x, mad(y, all.y, all.z)),  minh);

                res.x += subh_cache((int)subh);

                res.y += subh_cache((int)mad(4.f, all.y, subh));

                res.z += subh_cache((int)mad(8.f, all.y, subh));

                res.w += subh_cache((int)mad(12.f, all.y, subh));

            }

 

            barrier(CLK_LOCAL_MEM_FENCE);

        }

    const int sidx = (sbidx * 8 + 2 * stidx);

    const int idx = bidx + sidx;

    const float x = idx + off_x;

    const float y = idy + off_y;

        float minh = floor(all.z + bx * all.x - by * all.y + all.w);

            float4 all = c_all[proje + i]; all.y = - all.y;

//          float minh = all.z + bx * all.x - by * all.y + all.w;

            float minh = floor(all.z + bx * all.x + by * all.y + all.w);

            h.x = mad(49.f, i, mad(x, all.x, mad(y, all.y, all.z))) - minh;

            

//            h.x = mad(49.f, i, mad(sidx, all.x, mad(sidy, all.y, all.w)));

            h.y = h.x + all.x;

/*

*/

 

 

#undef subh_cache

#define subh_cache(subh) cache[(int)(subh)]

 

#define PPT 6

__kernel void hst_backproject_6ppt_oversample4(const int num_proj, 

    const int num_bins, 

    const float off_x, 

    const float off_y, 

    __read_only image2d_t sinogram,

    __global float *slice,

    __constant float4 *c_all, 

    __local float *cache

    ) {

    const int tidx = get_local_id(0);

    const int tidy = get_local_id(1);

 

    const int bidx = PPT * get_group_id(0) * get_local_size(0);

    const int bidy = PPT * (get_global_id(1) - get_local_id(1));

 

    const float bx = bidx + off_x;

    const float by = bidy + off_y;

 

/*

    const int sbidx = tidy % 4;

    const int sbidy = tidy / 4;

    

    const int stidx = tidx % 4;

    const int stidy = tidx / 4;

 

    const int sidx = (sbidx * 4 + stidx);

    const int sidy = (sbidy * 4 + stidy);

*/

 

    const int sidx = tidx;

    const int sidy = tidy;

 

 

    const float sx = sidx;

    const float sy = sidy;

 

    const int idx = bidx + sidx;

    const int idy = bidy + sidy;

 

    const float x = idx + off_x;

    const float y = idy + off_y;

 

    const int slice_width = PPT * get_global_size(0);

 

    float res[PPT][PPT] = {0};

    for (int proj = 0; proj < num_proj; proj++) {

        float4 all = c_all[proj]; all.y = -all.y;

        float minh = floor(all.z + bx * all.x + by * all.y + all.w);

 

#pragma unroll 5

        for (int i = 0; i < 5; i++) {

            int pos = i * get_local_size(0) * get_local_size(0) + tidy * get_local_size(0) + tidx;

            cache[pos] = read_imagef(sinogram, volumeSampler, (float2)(minh + 0.25f*pos, proj + .5f)).x;

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

 

 

        all *= 4;

        float h = mad(x, all.x, mad(y, all.y, all.z)) - 4 * minh;

 

#pragma unroll 6

        for (int i = 0; i < PPT; i++) {

#pragma unroll 6

            for (int j = 0; j < PPT; j++) {

                res[i][j] += subh_cache(h + (i * get_local_size(0)) * all.y + (j * get_local_size(0)) * all.x);

            }

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

    }

 

#pragma unroll 6

    for (int i = 0; i < PPT; i++) {

#pragma unroll 6

        for (int j = 0; j < PPT; j++) {

            slice[(idy + get_local_size(1) * i) * slice_width +  idx + get_local_size(0) * j] = res[i][j];

        }

    }

}

 

 

__kernel void hst_backproject_6ppt_tex(const int num_proj, 

    const int num_bins, 

    const float off_x, 

    const float off_y, 

    __read_only image2d_t sinogram,

    __global float *slice,

    __constant float4 *c_all/*, 

    __local float *shared*/

    ) {

 

    const int tidx = get_local_id(0);

    const int tidy = get_local_id(1);

 

    const int sidx = tidx;

    const int sidy = tidy;

 

    const int idx = PPT * get_group_id(0) * get_local_size(0) + sidx;

    const int idy = PPT * (get_global_id(1) - get_local_id(1)) + sidy;

 

    const float x = idx + off_x;

    const float y = idy + off_y;

 

    const int slice_width = PPT * get_global_size(0);

 

    float res[PPT][PPT] = {0};

 

    for (int proj = 0; proj < num_proj; proj++) {

        float4 all = c_all[proj];

        all.y = - all.y;

        float subh = mad(x, all.x, mad(y, all.y, all.z));

 

#pragma unroll 6

        for (int i = 0; i < PPT; i++) {

#pragma unroll 6

            for (int j = 0; j < PPT; j++) {

                res[i][j] += read_imagef(sinogram, volumeSampler, (float2)(subh + i * get_local_size(1) * all.y + j * get_local_size(0) * all.x, proj + .5f)).x;

            }

        }

    }

 

 

#pragma unroll 6

    for (int i = 0; i < PPT; i++) {

#pragma unroll 6

        for (int j = 0; j < PPT; j++) {

            slice[(idy + get_local_size(1) * i) * slice_width +  idx + get_local_size(0) * j] = res[i][j];

        }

    }

}

 

 

/*

__kernel void hst_backproject_6ppt(const int num_proj, 

    const int num_bins, 

    const float off_x, 

    const float off_y, 

    __read_only image2d_t sinogram,

    __global float *slice,

    __constant float4 *c_all, 

    __local float *cache) {

 

 

    const int mode = (get_group_id(1) * get_num_groups(0) + get_group_id(0)) % 5;

    if (mode < 2) {

        hst_backproject_test_oversample4_2(num_proj, num_bins, off_x, off_y, sinogram, slice, c_all, cache);

    } else {

        hst_backproject_test_oversample4_1(num_proj, num_bins, off_x, off_y, sinogram, slice, c_all, cache);

    }

}

*/

 

 

#undef PPT

#define PPT 8

/*

__kernel void hst_backproject_8ppt_vliw_oversample4(const int num_proj, 

    const int num_bins, 

    const float off_x, 

    const float off_y, 

    __read_only image2d_t sinogram,

    __global float *slice,

    __constant float4 *c_all, 

    __local float *cache

    ) {

    const int tidx = get_local_id(0);

    const int tidy = get_local_id(1);

 

    const int bidx = PPT * get_group_id(0) * get_local_size(0);

    const int bidy = PPT * (get_global_id(1) - get_local_id(1));

 

    const float bx = bidx + off_x;

    const float by = bidy + off_y;

 

    const int sidx = tidx;

    const int sidy = tidy;

 

    const int idx = bidx + sidx;

    const int idy = bidy + sidy;

 

    const float x = idx + off_x;

    const float y = idy + off_y;

 

    const int slice_width = PPT * get_global_size(0);

 

    float res[PPT][PPT] = {0};

    for (int proj = 0; proj < num_proj; proj++) {

        float4 all = c_all[proj]; all.y = -all.y;

        float minh = floor(all.z + bx * all.x + by * all.y + all.w);

 

#pragma unroll 6

        for (int i = 0; i < 6; i++) {

            int pos = i * get_local_size(0) * get_local_size(0) + tidy * get_local_size(0) + tidx;

            cache[pos] = read_imagef(sinogram, volumeSampler, (float2)(minh + 0.25f*pos, proj + .5f)).x;

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

 

 

        all *= 4;

        float h = mad(x, all.x, mad(y, all.y, all.z)) - 4 * minh;

 

#pragma unroll 8

        for (int i = 0; i < PPT; i++) {

#pragma unroll 8

            for (int j = 0; j < PPT; j++) {

                res[i][j] += subh_cache(h + (i * get_local_size(0)) * all.y + (j * get_local_size(0)) * all.x);

            }

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

    }

 

#pragma unroll 8

    for (int i = 0; i < PPT; i++) {

#pragma unroll 8

        for (int j = 0; j < PPT; j++) {

            slice[(idy + get_local_size(1) * i) * slice_width +  idx + get_local_size(0) * j] = res[i][j];

        }

    }

}

*/

 

__kernel void hst_backproject_8ppt_vliw_oversample4(const int num_proj, 

    const int num_bins, 

    const float off_x, 

    const float off_y, 

    __read_only image2d_t sinogram,

    __global float *slice,

    __constant float4 *c_all, 

    __local float *cache

    ) {

    const int tidx = get_local_id(0);

    const int tidy = get_local_id(1);

 

    const int bidx = PPT * get_group_id(0) * get_local_size(0);

    const int bidy = PPT * (get_global_id(1) - get_local_id(1));

 

    const float bx = bidx + off_x;

    const float by = bidy + off_y;

 

    const int sidx = tidx;

    const int sidy = tidy;

 

    const int idx = bidx + sidx;

    const int idy = bidy + sidy;

 

    const float x = idx + off_x;

    const float y = idy + off_y;

 

    const int slice_width = PPT * get_global_size(0);

 

float res_0_0 = 0;

float res_0_1 = 0;

float res_0_2 = 0;

float res_0_3 = 0;

float res_0_4 = 0;

float res_0_5 = 0;

float res_0_6 = 0;

float res_0_7 = 0;

float res_1_0 = 0;

float res_1_1 = 0;

float res_1_2 = 0;

float res_1_3 = 0;

float res_1_4 = 0;

float res_1_5 = 0;

float res_1_6 = 0;

float res_1_7 = 0;

float res_2_0 = 0;

float res_2_1 = 0;

float res_2_2 = 0;

float res_2_3 = 0;

float res_2_4 = 0;

float res_2_5 = 0;

float res_2_6 = 0;

float res_2_7 = 0;

float res_3_0 = 0;

float res_3_1 = 0;

float res_3_2 = 0;

float res_3_3 = 0;

float res_3_4 = 0;

float res_3_5 = 0;

float res_3_6 = 0;

float res_3_7 = 0;

float res_4_0 = 0;

float res_4_1 = 0;

float res_4_2 = 0;

float res_4_3 = 0;

float res_4_4 = 0;

float res_4_5 = 0;

float res_4_6 = 0;

float res_4_7 = 0;

float res_5_0 = 0;

float res_5_1 = 0;

float res_5_2 = 0;

float res_5_3 = 0;

float res_5_4 = 0;

float res_5_5 = 0;

float res_5_6 = 0;

float res_5_7 = 0;

float res_6_0 = 0;

float res_6_1 = 0;

float res_6_2 = 0;

float res_6_3 = 0;

float res_6_4 = 0;

float res_6_5 = 0;

float res_6_6 = 0;

float res_6_7 = 0;

float res_7_0 = 0;

float res_7_1 = 0;

float res_7_2 = 0;

float res_7_3 = 0;

float res_7_4 = 0;

float res_7_5 = 0;

float res_7_6 = 0;

float res_7_7 = 0;

 

    for (int proj = 0; proj < num_proj; proj++) {

        float4 all = c_all[proj]; all.y = -all.y;

        float minh = floor(all.z + bx * all.x + by * all.y + all.w);

 

#pragma unroll 6

        for (int i = 0; i < 6; i++) {

            int pos = i * get_local_size(0) * get_local_size(0) + tidy * get_local_size(0) + tidx;

            cache[pos] = read_imagef(sinogram, volumeSampler, (float2)(minh + 0.25f*pos, proj + .5f)).x;

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

 

 

        all *= 4;

        float h = mad(x, all.x, mad(y, all.y, all.z)) - 4 * minh;

 

res_0_0 += subh_cache(h + (0 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x);

res_0_1 += subh_cache(h + (0 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x);

res_0_2 += subh_cache(h + (0 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x);

res_0_3 += subh_cache(h + (0 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x);

res_0_4 += subh_cache(h + (0 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x);

res_0_5 += subh_cache(h + (0 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x);

res_0_6 += subh_cache(h + (0 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x);

res_0_7 += subh_cache(h + (0 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x);

res_1_0 += subh_cache(h + (1 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x);

res_1_1 += subh_cache(h + (1 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x);

res_1_2 += subh_cache(h + (1 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x);

res_1_3 += subh_cache(h + (1 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x);

res_1_4 += subh_cache(h + (1 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x);

res_1_5 += subh_cache(h + (1 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x);

res_1_6 += subh_cache(h + (1 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x);

res_1_7 += subh_cache(h + (1 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x);

res_2_0 += subh_cache(h + (2 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x);

res_2_1 += subh_cache(h + (2 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x);

res_2_2 += subh_cache(h + (2 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x);

res_2_3 += subh_cache(h + (2 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x);

res_2_4 += subh_cache(h + (2 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x);

res_2_5 += subh_cache(h + (2 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x);

res_2_6 += subh_cache(h + (2 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x);

res_2_7 += subh_cache(h + (2 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x);

res_3_0 += subh_cache(h + (3 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x);

res_3_1 += subh_cache(h + (3 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x);

res_3_2 += subh_cache(h + (3 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x);

res_3_3 += subh_cache(h + (3 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x);

res_3_4 += subh_cache(h + (3 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x);

res_3_5 += subh_cache(h + (3 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x);

res_3_6 += subh_cache(h + (3 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x);

res_3_7 += subh_cache(h + (3 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x);

res_4_0 += subh_cache(h + (4 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x);

res_4_1 += subh_cache(h + (4 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x);

res_4_2 += subh_cache(h + (4 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x);

res_4_3 += subh_cache(h + (4 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x);

res_4_4 += subh_cache(h + (4 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x);

res_4_5 += subh_cache(h + (4 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x);

res_4_6 += subh_cache(h + (4 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x);

res_4_7 += subh_cache(h + (4 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x);

res_5_0 += subh_cache(h + (5 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x);

res_5_1 += subh_cache(h + (5 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x);

res_5_2 += subh_cache(h + (5 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x);

res_5_3 += subh_cache(h + (5 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x);

res_5_4 += subh_cache(h + (5 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x);

res_5_5 += subh_cache(h + (5 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x);

res_5_6 += subh_cache(h + (5 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x);

res_5_7 += subh_cache(h + (5 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x);

res_6_0 += subh_cache(h + (6 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x);

res_6_1 += subh_cache(h + (6 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x);

res_6_2 += subh_cache(h + (6 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x);

res_6_3 += subh_cache(h + (6 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x);

res_6_4 += subh_cache(h + (6 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x);

res_6_5 += subh_cache(h + (6 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x);

res_6_6 += subh_cache(h + (6 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x);

res_6_7 += subh_cache(h + (6 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x);

res_7_0 += subh_cache(h + (7 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x);

res_7_1 += subh_cache(h + (7 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x);

res_7_2 += subh_cache(h + (7 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x);

res_7_3 += subh_cache(h + (7 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x);

res_7_4 += subh_cache(h + (7 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x);

res_7_5 += subh_cache(h + (7 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x);

res_7_6 += subh_cache(h + (7 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x);

res_7_7 += subh_cache(h + (7 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x);

 

        barrier(CLK_LOCAL_MEM_FENCE);

    }

 

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 0] = res_0_0;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 1] = res_0_1;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 2] = res_0_2;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 3] = res_0_3;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 4] = res_0_4;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 5] = res_0_5;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 6] = res_0_6;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 7] = res_0_7;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 0] = res_1_0;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 1] = res_1_1;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 2] = res_1_2;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 3] = res_1_3;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 4] = res_1_4;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 5] = res_1_5;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 6] = res_1_6;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 7] = res_1_7;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 0] = res_2_0;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 1] = res_2_1;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 2] = res_2_2;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 3] = res_2_3;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 4] = res_2_4;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 5] = res_2_5;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 6] = res_2_6;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 7] = res_2_7;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 0] = res_3_0;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 1] = res_3_1;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 2] = res_3_2;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 3] = res_3_3;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 4] = res_3_4;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 5] = res_3_5;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 6] = res_3_6;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 7] = res_3_7;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 0] = res_4_0;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 1] = res_4_1;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 2] = res_4_2;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 3] = res_4_3;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 4] = res_4_4;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 5] = res_4_5;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 6] = res_4_6;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 7] = res_4_7;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 0] = res_5_0;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 1] = res_5_1;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 2] = res_5_2;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 3] = res_5_3;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 4] = res_5_4;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 5] = res_5_5;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 6] = res_5_6;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 7] = res_5_7;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 0] = res_6_0;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 1] = res_6_1;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 2] = res_6_2;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 3] = res_6_3;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 4] = res_6_4;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 5] = res_6_5;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 6] = res_6_6;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 7] = res_6_7;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 0] = res_7_0;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 1] = res_7_1;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 2] = res_7_2;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 3] = res_7_3;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 4] = res_7_4;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 5] = res_7_5;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 6] = res_7_6;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 7] = res_7_7;

}

 

/*

__kernel void hst_backproject_8ppt_vliw_linear(const int num_proj, 

    const int num_bins, 

    const float off_x, 

    const float off_y, 

    __read_only image2d_t sinogram,

    __global float *slice,

    __constant float4 *c_all, 

    __local float *share

    ) {

    __local float2 *cache = (__local float2*)share;

    const int tidx = get_local_id(0);

    const int tidy = get_local_id(1);

 

    const int bidx = PPT * get_group_id(0) * get_local_size(0);

    const int bidy = PPT * (get_global_id(1) - get_local_id(1));

 

    const float bx = bidx + off_x;

    const float by = bidy + off_y;

 

    const int sidx = tidx;

    const int sidy = tidy;

 

    const int idx = bidx + sidx;

    const int idy = bidy + sidy;

 

    const float x = idx + off_x;

    const float y = idy + off_y;

 

    const int slice_width = PPT * get_global_size(0);

 

    float res[PPT][PPT] = {0};

    for (int proj = 0; proj < num_proj; proj++) {

        float4 all = c_all[proj]; all.y = -all.y;

        float minh = floor(all.z + bx * all.x + by * all.y + all.w);

 

#pragma unroll 2

        for (int i = 0; i < 2; i++) {

            int pos = i * get_local_size(0) * get_local_size(0) + tidy * get_local_size(0) + tidx;

            cache[pos].x = read_imagef(sinogram, volumeSampler, (float2)(minh + pos, proj + .5f)).x;

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

 

        for (int i = 0; i < 2; i++) {

            int pos = i * get_local_size(0) * get_local_size(0) + tidy * get_local_size(0) + tidx;

            cache[pos].y = cache[pos + 1].x - cache[pos].x;

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

 

        float h = mad(x, all.x, mad(y, all.y, all.z)) - minh;

        float2 v;

        float subh, isubh;

        

#pragma unroll 8

        for (int i = 0; i < PPT; i++) {

#pragma unroll 8

            for (int j = 0; j < PPT; j++) {

                subh = h + (i * get_local_size(0)) * all.y + (j * get_local_size(0)) * all.x;

                isubh = floor(subh);

                v = subh_cache(isubh);

                res[i][j] += v.x + (subh - isubh)*v.y;

            }

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

    }

 

#pragma unroll 8

    for (int i = 0; i < PPT; i++) {

#pragma unroll 8

        for (int j = 0; j < PPT; j++) {

            slice[(idy + get_local_size(1) * i) * slice_width +  idx + get_local_size(0) * j] = res[i][j];

        }

    }

}

*/

 

/*

__kernel void hst_backproject_8ppt_vliw_linear(const int num_proj, 

    const int num_bins, 

    const float off_x, 

    const float off_y, 

    __read_only image2d_t sinogram,

    __global float *slice,

    __constant float4 *c_all, 

    __local float *share

    ) {

    __local float2 *cache = (__local float2*)share;

    const int tidx = get_local_id(0);

    const int tidy = get_local_id(1);

 

    const int bidx = PPT * get_group_id(0) * get_local_size(0);

    const int bidy = PPT * (get_global_id(1) - get_local_id(1));

 

    const float bx = bidx + off_x;

    const float by = bidy + off_y;

 

    const int sidx = tidx;

    const int sidy = tidy;

 

    const int idx = bidx + sidx;

    const int idy = bidy + sidy;

 

    const float x = idx + off_x;

    const float y = idy + off_y;

 

    const int slice_width = PPT * get_global_size(0);

 

float res_0_0 = 0;

float res_0_1 = 0;

float res_0_2 = 0;

float res_0_3 = 0;

float res_0_4 = 0;

float res_0_5 = 0;

float res_0_6 = 0;

float res_0_7 = 0;

float res_1_0 = 0;

float res_1_1 = 0;

float res_1_2 = 0;

float res_1_3 = 0;

float res_1_4 = 0;

float res_1_5 = 0;

float res_1_6 = 0;

float res_1_7 = 0;

float res_2_0 = 0;

float res_2_1 = 0;

float res_2_2 = 0;

float res_2_3 = 0;

float res_2_4 = 0;

float res_2_5 = 0;

float res_2_6 = 0;

float res_2_7 = 0;

float res_3_0 = 0;

float res_3_1 = 0;

float res_3_2 = 0;

float res_3_3 = 0;

float res_3_4 = 0;

float res_3_5 = 0;

float res_3_6 = 0;

float res_3_7 = 0;

float res_4_0 = 0;

float res_4_1 = 0;

float res_4_2 = 0;

float res_4_3 = 0;

float res_4_4 = 0;

float res_4_5 = 0;

float res_4_6 = 0;

float res_4_7 = 0;

float res_5_0 = 0;

float res_5_1 = 0;

float res_5_2 = 0;

float res_5_3 = 0;

float res_5_4 = 0;

float res_5_5 = 0;

float res_5_6 = 0;

float res_5_7 = 0;

float res_6_0 = 0;

float res_6_1 = 0;

float res_6_2 = 0;

float res_6_3 = 0;

float res_6_4 = 0;

float res_6_5 = 0;

float res_6_6 = 0;

float res_6_7 = 0;

float res_7_0 = 0;

float res_7_1 = 0;

float res_7_2 = 0;

float res_7_3 = 0;

float res_7_4 = 0;

float res_7_5 = 0;

float res_7_6 = 0;

float res_7_7 = 0;

 

 

 

    for (int proj = 0; proj < num_proj; proj++) {

        float4 all = c_all[proj]; all.y = -all.y;

        float minh = floor(all.z + bx * all.x + by * all.y + all.w);

 

#pragma unroll 2

        for (int i = 0; i < 2; i++) {

            int pos = i * get_local_size(0) * get_local_size(0) + tidy * get_local_size(0) + tidx;

            cache[pos].x = read_imagef(sinogram, volumeSampler, (float2)(minh + pos, proj + .5f)).x;

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

 

        for (int i = 0; i < 2; i++) {

            int pos = i * get_local_size(0) * get_local_size(0) + tidy * get_local_size(0) + tidx;

            cache[pos].y = cache[pos + 1].x - cache[pos].x;

        }

 

        barrier(CLK_LOCAL_MEM_FENCE);

 

        float h = mad(x, all.x, mad(y, all.y, all.z)) - minh;

        float2 v;

        float subh, isubh;

 

subh = h + (0 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_0_0 += v.x + (subh - isubh) * v.y;

subh = h + (0 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_0_1 += v.x + (subh - isubh) * v.y;

subh = h + (0 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_0_2 += v.x + (subh - isubh) * v.y;

subh = h + (0 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_0_3 += v.x + (subh - isubh) * v.y;

subh = h + (0 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_0_4 += v.x + (subh - isubh) * v.y;

subh = h + (0 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_0_5 += v.x + (subh - isubh) * v.y;

subh = h + (0 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_0_6 += v.x + (subh - isubh) * v.y;

subh = h + (0 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_0_7 += v.x + (subh - isubh) * v.y;

subh = h + (1 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_1_0 += v.x + (subh - isubh) * v.y;

subh = h + (1 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_1_1 += v.x + (subh - isubh) * v.y;

subh = h + (1 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_1_2 += v.x + (subh - isubh) * v.y;

subh = h + (1 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_1_3 += v.x + (subh - isubh) * v.y;

subh = h + (1 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_1_4 += v.x + (subh - isubh) * v.y;

subh = h + (1 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_1_5 += v.x + (subh - isubh) * v.y;

subh = h + (1 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_1_6 += v.x + (subh - isubh) * v.y;

subh = h + (1 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_1_7 += v.x + (subh - isubh) * v.y;

subh = h + (2 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_2_0 += v.x + (subh - isubh) * v.y;

subh = h + (2 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_2_1 += v.x + (subh - isubh) * v.y;

subh = h + (2 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_2_2 += v.x + (subh - isubh) * v.y;

subh = h + (2 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_2_3 += v.x + (subh - isubh) * v.y;

subh = h + (2 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_2_4 += v.x + (subh - isubh) * v.y;

subh = h + (2 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_2_5 += v.x + (subh - isubh) * v.y;

subh = h + (2 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_2_6 += v.x + (subh - isubh) * v.y;

subh = h + (2 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_2_7 += v.x + (subh - isubh) * v.y;

subh = h + (3 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_3_0 += v.x + (subh - isubh) * v.y;

subh = h + (3 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_3_1 += v.x + (subh - isubh) * v.y;

subh = h + (3 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_3_2 += v.x + (subh - isubh) * v.y;

subh = h + (3 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_3_3 += v.x + (subh - isubh) * v.y;

subh = h + (3 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_3_4 += v.x + (subh - isubh) * v.y;

subh = h + (3 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_3_5 += v.x + (subh - isubh) * v.y;

subh = h + (3 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_3_6 += v.x + (subh - isubh) * v.y;

subh = h + (3 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_3_7 += v.x + (subh - isubh) * v.y;

subh = h + (4 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_4_0 += v.x + (subh - isubh) * v.y;

subh = h + (4 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_4_1 += v.x + (subh - isubh) * v.y;

subh = h + (4 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_4_2 += v.x + (subh - isubh) * v.y;

subh = h + (4 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_4_3 += v.x + (subh - isubh) * v.y;

subh = h + (4 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_4_4 += v.x + (subh - isubh) * v.y;

subh = h + (4 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_4_5 += v.x + (subh - isubh) * v.y;

subh = h + (4 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_4_6 += v.x + (subh - isubh) * v.y;

subh = h + (4 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_4_7 += v.x + (subh - isubh) * v.y;

subh = h + (5 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_5_0 += v.x + (subh - isubh) * v.y;

subh = h + (5 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_5_1 += v.x + (subh - isubh) * v.y;

subh = h + (5 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_5_2 += v.x + (subh - isubh) * v.y;

subh = h + (5 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_5_3 += v.x + (subh - isubh) * v.y;

subh = h + (5 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_5_4 += v.x + (subh - isubh) * v.y;

subh = h + (5 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_5_5 += v.x + (subh - isubh) * v.y;

subh = h + (5 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_5_6 += v.x + (subh - isubh) * v.y;

subh = h + (5 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_5_7 += v.x + (subh - isubh) * v.y;

subh = h + (6 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_6_0 += v.x + (subh - isubh) * v.y;

subh = h + (6 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_6_1 += v.x + (subh - isubh) * v.y;

subh = h + (6 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_6_2 += v.x + (subh - isubh) * v.y;

subh = h + (6 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_6_3 += v.x + (subh - isubh) * v.y;

subh = h + (6 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_6_4 += v.x + (subh - isubh) * v.y;

subh = h + (6 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_6_5 += v.x + (subh - isubh) * v.y;

subh = h + (6 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_6_6 += v.x + (subh - isubh) * v.y;

subh = h + (6 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_6_7 += v.x + (subh - isubh) * v.y;

subh = h + (7 * get_local_size(0)) * all.y + (0 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_7_0 += v.x + (subh - isubh) * v.y;

subh = h + (7 * get_local_size(0)) * all.y + (1 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_7_1 += v.x + (subh - isubh) * v.y;

subh = h + (7 * get_local_size(0)) * all.y + (2 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_7_2 += v.x + (subh - isubh) * v.y;

subh = h + (7 * get_local_size(0)) * all.y + (3 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_7_3 += v.x + (subh - isubh) * v.y;

subh = h + (7 * get_local_size(0)) * all.y + (4 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_7_4 += v.x + (subh - isubh) * v.y;

subh = h + (7 * get_local_size(0)) * all.y + (5 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_7_5 += v.x + (subh - isubh) * v.y;

subh = h + (7 * get_local_size(0)) * all.y + (6 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_7_6 += v.x + (subh - isubh) * v.y;

subh = h + (7 * get_local_size(0)) * all.y + (7 * get_local_size(0)) * all.x; isubh = floor(subh); v = subh_cache(isubh); res_7_7 += v.x + (subh - isubh) * v.y;

 

        barrier(CLK_LOCAL_MEM_FENCE);

    }

 

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 0] = res_0_0;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 1] = res_0_1;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 2] = res_0_2;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 3] = res_0_3;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 4] = res_0_4;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 5] = res_0_5;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 6] = res_0_6;

slice[(idy + get_local_size(1) * 0) * slice_width +  idx + get_local_size(0) * 7] = res_0_7;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 0] = res_1_0;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 1] = res_1_1;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 2] = res_1_2;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 3] = res_1_3;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 4] = res_1_4;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 5] = res_1_5;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 6] = res_1_6;

slice[(idy + get_local_size(1) * 1) * slice_width +  idx + get_local_size(0) * 7] = res_1_7;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 0] = res_2_0;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 1] = res_2_1;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 2] = res_2_2;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 3] = res_2_3;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 4] = res_2_4;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 5] = res_2_5;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 6] = res_2_6;

slice[(idy + get_local_size(1) * 2) * slice_width +  idx + get_local_size(0) * 7] = res_2_7;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 0] = res_3_0;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 1] = res_3_1;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 2] = res_3_2;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 3] = res_3_3;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 4] = res_3_4;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 5] = res_3_5;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 6] = res_3_6;

slice[(idy + get_local_size(1) * 3) * slice_width +  idx + get_local_size(0) * 7] = res_3_7;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 0] = res_4_0;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 1] = res_4_1;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 2] = res_4_2;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 3] = res_4_3;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 4] = res_4_4;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 5] = res_4_5;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 6] = res_4_6;

slice[(idy + get_local_size(1) * 4) * slice_width +  idx + get_local_size(0) * 7] = res_4_7;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 0] = res_5_0;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 1] = res_5_1;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 2] = res_5_2;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 3] = res_5_3;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 4] = res_5_4;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 5] = res_5_5;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 6] = res_5_6;

slice[(idy + get_local_size(1) * 5) * slice_width +  idx + get_local_size(0) * 7] = res_5_7;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 0] = res_6_0;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 1] = res_6_1;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 2] = res_6_2;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 3] = res_6_3;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 4] = res_6_4;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 5] = res_6_5;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 6] = res_6_6;

slice[(idy + get_local_size(1) * 6) * slice_width +  idx + get_local_size(0) * 7] = res_6_7;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 0] = res_7_0;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 1] = res_7_1;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 2] = res_7_2;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 3] = res_7_3;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 4] = res_7_4;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 5] = res_7_5;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 6] = res_7_6;

slice[(idy + get_local_size(1) * 7) * slice_width +  idx + get_local_size(0) * 7] = res_7_7;

}

*/

 

__kernel void hst_backproject_6ppt_vliw_tex(const int num_proj, 

    const int num_bins, 

    const float off_x, 

    const float off_y, 

    __read_only image2d_t sinogram,

    __global float *slice,

    __constant float4 *c_all/*, 

    __local float *shared*/

    ) {

 

    const int tidx = get_local_id(0);

    const int tidy = get_local_id(1);

 

    const int sidx = tidx;

    const int sidy = tidy;

 

    const int idx = PPT * get_group_id(0) * get_local_size(0) + sidx;

    const int idy = PPT * (get_global_id(1) - get_local_id(1)) + sidy;

 

    const float x = idx + off_x;

    const float y = idy + off_y;

 

    const int slice_width = PPT * get_global_size(0);

 

float res_0_0 = 0;

float res_0_1 = 0;

float res_0_2 = 0;

float res_0_3 = 0;

float res_0_4 = 0;

float res_0_5 = 0;

float res_1_0 = 0;

float res_1_1 = 0;

float res_1_2 = 0;

float res_1_3 = 0;

float res_1_4 = 0;

float res_1_5 = 0;

float res_2_0 = 0;

float res_2_1 = 0;

float res_2_2 = 0;

float res_2_3 = 0;

float res_2_4 = 0;

float res_2_5 = 0;

float res_3_0 = 0;

float res_3_1 = 0;

float res_3_2 = 0;

float res_3_3 = 0;

float res_3_4 = 0;

float res_3_5 = 0;

float res_4_0 = 0;

float res_4_1 = 0;

float res_4_2 = 0;

float res_4_3 = 0;

float res_4_4 = 0;

float res_4_5 = 0;

float res_5_0 = 0;

float res_5_1 = 0;

float res_5_2 = 0;

float res_5_3 = 0;

float res_5_4 = 0;

float res_5_5 = 0;

 

    for (int proj = 0; proj < num_proj; proj++) {

        float4 all = c_all[proj];

        all.y = - all.y;

        float subh = mad(x, all.x, mad(y, all.y, all.z));

 

res_0_0 += read_imagef(sinogram, volumeSampler, (float2)(subh + 0 * BLOCK_SIZE * all.y + 0 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_0_1 += read_imagef(sinogram, volumeSampler, (float2)(subh + 0 * BLOCK_SIZE * all.y + 1 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_0_2 += read_imagef(sinogram, volumeSampler, (float2)(subh + 0 * BLOCK_SIZE * all.y + 2 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_0_3 += read_imagef(sinogram, volumeSampler, (float2)(subh + 0 * BLOCK_SIZE * all.y + 3 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_0_4 += read_imagef(sinogram, volumeSampler, (float2)(subh + 0 * BLOCK_SIZE * all.y + 4 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_0_5 += read_imagef(sinogram, volumeSampler, (float2)(subh + 0 * BLOCK_SIZE * all.y + 5 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_1_0 += read_imagef(sinogram, volumeSampler, (float2)(subh + 1 * BLOCK_SIZE * all.y + 0 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_1_1 += read_imagef(sinogram, volumeSampler, (float2)(subh + 1 * BLOCK_SIZE * all.y + 1 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_1_2 += read_imagef(sinogram, volumeSampler, (float2)(subh + 1 * BLOCK_SIZE * all.y + 2 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_1_3 += read_imagef(sinogram, volumeSampler, (float2)(subh + 1 * BLOCK_SIZE * all.y + 3 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_1_4 += read_imagef(sinogram, volumeSampler, (float2)(subh + 1 * BLOCK_SIZE * all.y + 4 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_1_5 += read_imagef(sinogram, volumeSampler, (float2)(subh + 1 * BLOCK_SIZE * all.y + 5 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_2_0 += read_imagef(sinogram, volumeSampler, (float2)(subh + 2 * BLOCK_SIZE * all.y + 0 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_2_1 += read_imagef(sinogram, volumeSampler, (float2)(subh + 2 * BLOCK_SIZE * all.y + 1 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_2_2 += read_imagef(sinogram, volumeSampler, (float2)(subh + 2 * BLOCK_SIZE * all.y + 2 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_2_3 += read_imagef(sinogram, volumeSampler, (float2)(subh + 2 * BLOCK_SIZE * all.y + 3 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_2_4 += read_imagef(sinogram, volumeSampler, (float2)(subh + 2 * BLOCK_SIZE * all.y + 4 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_2_5 += read_imagef(sinogram, volumeSampler, (float2)(subh + 2 * BLOCK_SIZE * all.y + 5 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_3_0 += read_imagef(sinogram, volumeSampler, (float2)(subh + 3 * BLOCK_SIZE * all.y + 0 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_3_1 += read_imagef(sinogram, volumeSampler, (float2)(subh + 3 * BLOCK_SIZE * all.y + 1 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_3_2 += read_imagef(sinogram, volumeSampler, (float2)(subh + 3 * BLOCK_SIZE * all.y + 2 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_3_3 += read_imagef(sinogram, volumeSampler, (float2)(subh + 3 * BLOCK_SIZE * all.y + 3 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_3_4 += read_imagef(sinogram, volumeSampler, (float2)(subh + 3 * BLOCK_SIZE * all.y + 4 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_3_5 += read_imagef(sinogram, volumeSampler, (float2)(subh + 3 * BLOCK_SIZE * all.y + 5 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_4_0 += read_imagef(sinogram, volumeSampler, (float2)(subh + 4 * BLOCK_SIZE * all.y + 0 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_4_1 += read_imagef(sinogram, volumeSampler, (float2)(subh + 4 * BLOCK_SIZE * all.y + 1 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_4_2 += read_imagef(sinogram, volumeSampler, (float2)(subh + 4 * BLOCK_SIZE * all.y + 2 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_4_3 += read_imagef(sinogram, volumeSampler, (float2)(subh + 4 * BLOCK_SIZE * all.y + 3 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_4_4 += read_imagef(sinogram, volumeSampler, (float2)(subh + 4 * BLOCK_SIZE * all.y + 4 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_4_5 += read_imagef(sinogram, volumeSampler, (float2)(subh + 4 * BLOCK_SIZE * all.y + 5 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_5_0 += read_imagef(sinogram, volumeSampler, (float2)(subh + 5 * BLOCK_SIZE * all.y + 0 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_5_1 += read_imagef(sinogram, volumeSampler, (float2)(subh + 5 * BLOCK_SIZE * all.y + 1 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_5_2 += read_imagef(sinogram, volumeSampler, (float2)(subh + 5 * BLOCK_SIZE * all.y + 2 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_5_3 += read_imagef(sinogram, volumeSampler, (float2)(subh + 5 * BLOCK_SIZE * all.y + 3 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_5_4 += read_imagef(sinogram, volumeSampler, (float2)(subh + 5 * BLOCK_SIZE * all.y + 4 * BLOCK_SIZE * all.x, proj + .5f)).x;

res_5_5 += read_imagef(sinogram, volumeSampler, (float2)(subh + 5 * BLOCK_SIZE * all.y + 5 * BLOCK_SIZE * all.x, proj + .5f)).x;

 

    }

 

 

slice[(idy + BLOCK_SIZE * 0) * slice_width + idx + BLOCK_SIZE * 0] = res_0_0;

slice[(idy + BLOCK_SIZE * 0) * slice_width + idx + BLOCK_SIZE * 1] = res_0_1;

slice[(idy + BLOCK_SIZE * 0) * slice_width + idx + BLOCK_SIZE * 2] = res_0_2;

slice[(idy + BLOCK_SIZE * 0) * slice_width + idx + BLOCK_SIZE * 3] = res_0_3;

slice[(idy + BLOCK_SIZE * 0) * slice_width + idx + BLOCK_SIZE * 4] = res_0_4;

slice[(idy + BLOCK_SIZE * 0) * slice_width + idx + BLOCK_SIZE * 5] = res_0_5;

slice[(idy + BLOCK_SIZE * 1) * slice_width + idx + BLOCK_SIZE * 0] = res_1_0;

slice[(idy + BLOCK_SIZE * 1) * slice_width + idx + BLOCK_SIZE * 1] = res_1_1;

slice[(idy + BLOCK_SIZE * 1) * slice_width + idx + BLOCK_SIZE * 2] = res_1_2;

slice[(idy + BLOCK_SIZE * 1) * slice_width + idx + BLOCK_SIZE * 3] = res_1_3;

slice[(idy + BLOCK_SIZE * 1) * slice_width + idx + BLOCK_SIZE * 4] = res_1_4;

slice[(idy + BLOCK_SIZE * 1) * slice_width + idx + BLOCK_SIZE * 5] = res_1_5;

slice[(idy + BLOCK_SIZE * 2) * slice_width + idx + BLOCK_SIZE * 0] = res_2_0;

slice[(idy + BLOCK_SIZE * 2) * slice_width + idx + BLOCK_SIZE * 1] = res_2_1;

slice[(idy + BLOCK_SIZE * 2) * slice_width + idx + BLOCK_SIZE * 2] = res_2_2;

slice[(idy + BLOCK_SIZE * 2) * slice_width + idx + BLOCK_SIZE * 3] = res_2_3;

slice[(idy + BLOCK_SIZE * 2) * slice_width + idx + BLOCK_SIZE * 4] = res_2_4;

slice[(idy + BLOCK_SIZE * 2) * slice_width + idx + BLOCK_SIZE * 5] = res_2_5;

slice[(idy + BLOCK_SIZE * 3) * slice_width + idx + BLOCK_SIZE * 0] = res_3_0;

slice[(idy + BLOCK_SIZE * 3) * slice_width + idx + BLOCK_SIZE * 1] = res_3_1;

slice[(idy + BLOCK_SIZE * 3) * slice_width + idx + BLOCK_SIZE * 2] = res_3_2;

slice[(idy + BLOCK_SIZE * 3) * slice_width + idx + BLOCK_SIZE * 3] = res_3_3;

slice[(idy + BLOCK_SIZE * 3) * slice_width + idx + BLOCK_SIZE * 4] = res_3_4;

slice[(idy + BLOCK_SIZE * 3) * slice_width + idx + BLOCK_SIZE * 5] = res_3_5;

slice[(idy + BLOCK_SIZE * 4) * slice_width + idx + BLOCK_SIZE * 0] = res_4_0;

slice[(idy + BLOCK_SIZE * 4) * slice_width + idx + BLOCK_SIZE * 1] = res_4_1;

slice[(idy + BLOCK_SIZE * 4) * slice_width + idx + BLOCK_SIZE * 2] = res_4_2;

slice[(idy + BLOCK_SIZE * 4) * slice_width + idx + BLOCK_SIZE * 3] = res_4_3;

slice[(idy + BLOCK_SIZE * 4) * slice_width + idx + BLOCK_SIZE * 4] = res_4_4;

slice[(idy + BLOCK_SIZE * 4) * slice_width + idx + BLOCK_SIZE * 5] = res_4_5;

slice[(idy + BLOCK_SIZE * 5) * slice_width + idx + BLOCK_SIZE * 0] = res_5_0;

slice[(idy + BLOCK_SIZE * 5) * slice_width + idx + BLOCK_SIZE * 1] = res_5_1;

slice[(idy + BLOCK_SIZE * 5) * slice_width + idx + BLOCK_SIZE * 2] = res_5_2;

slice[(idy + BLOCK_SIZE * 5) * slice_width + idx + BLOCK_SIZE * 3] = res_5_3;

slice[(idy + BLOCK_SIZE * 5) * slice_width + idx + BLOCK_SIZE * 4] = res_5_4;

slice[(idy + BLOCK_SIZE * 5) * slice_width + idx + BLOCK_SIZE * 5] = res_5_5;

}