 
 template<unsigned int TILE_W, unsigned int TILE_H>
 __global__ static
-void d_convolve3x3Shared(ImageView<float> src, ImageView<float> dst
-                   )
+void d_convolve3x3Shared(ImageView<float> src, ImageView<float> dst)
 {
     const unsigned int tx = threadIdx.x;
     const unsigned int ty = threadIdx.y;

        h_dest = dest;
        SAIGA_ASSERT(h_dest == h_ref);
    }

    {
        float time;

    const unsigned int x = x_tile + tx;
    const unsigned int y = y_tile + ty;
    dst(x,y) = sum;
}



template<unsigned int TILE_W, unsigned int TILE_H>
__global__ static
void d_convolve3x3Shared2(ImageView<float> src, ImageView<float> dst)
{
    const unsigned int tx = threadIdx.x;
    const unsigned int ty = threadIdx.y;
    const unsigned int t = ty * TILE_W + tx;

    const unsigned int x_tile = blockIdx.x * TILE_W;
    const unsigned int y_tile = blockIdx.y * TILE_H;
    const unsigned int blockStartX = x_tile - 1;
    const unsigned int blockStartY = y_tile - 1;


    const unsigned int TILE_SIZE =  TILE_H * TILE_W;
    const unsigned int TILE_SIZE_WITH_BORDER = (TILE_H+2) * (TILE_W+2);
    __shared__ float sbuffer[TILE_H + 2][TILE_W + 2];


    //copy main data


    __syncthreads();


    float sum = 0;
#if 1
    for(int dy = -1; dy <= 1; ++dy){
        for(int dx = -1; dx <= 1; ++dx){
            int x = tx + 1 + dx;
            int y = ty + 1 + dy;
            sum += sbuffer[y][x];
        }
    }
#endif

    const unsigned int x = x_tile + tx;
    const unsigned int y = y_tile + ty;
    dst(x,y) = sum;
}


template<unsigned int TILE_W, unsigned int TILE_H>
__global__ static
void d_copySharedSync(ImageView<float> src, ImageView<float> dst)
{
    const unsigned int tx = threadIdx.x;
    const unsigned int ty = threadIdx.y;
    const unsigned int t = ty * TILE_W + tx;
    const unsigned int x_tile = blockIdx.x * TILE_W;
    const unsigned int y_tile = blockIdx.y * TILE_H;
    const unsigned int x = x_tile + tx;
    const unsigned int y = y_tile + ty;

    __shared__ float sbuffer[TILE_H][TILE_W];

    sbuffer[ty][tx]  = src(x,y);
    __syncthreads();
    dst(x,y) = sbuffer[ty][tx];
}

void convolutionTest3x3(){
