template <int ksize, typename T, typename D, typename B>
    __global__ void linearRowFilter(const DevMem2D_<T> src, PtrStep_<D> dst, int anchor, const B b)
    {
        typedef typename SmemType<T>::smem_t smem_t;

        __shared__ smem_t smem[BLOCK_DIM_Y * BLOCK_DIM_X * 3];

        const int x = BLOCK_DIM_X * blockIdx.x + threadIdx.x;
        const int y = BLOCK_DIM_Y * blockIdx.y + threadIdx.y;

        smem_t* sDataRow = smem + threadIdx.y * BLOCK_DIM_X * 3;

        if (y < src.rows)
        {
            const T* rowSrc = src.ptr(y);

            sDataRow[threadIdx.x                  ] = b.at_low(x - BLOCK_DIM_X, rowSrc);
            sDataRow[threadIdx.x + BLOCK_DIM_X    ] = b.at_high(x, rowSrc);
            sDataRow[threadIdx.x + BLOCK_DIM_X * 2] = b.at_high(x + BLOCK_DIM_X, rowSrc);

            __syncthreads();

            if (x < src.cols)
            {
                typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type sum_t;
                sum_t sum = VecTraits<sum_t>::all(0);

                sDataRow += threadIdx.x + BLOCK_DIM_X - anchor;

                #pragma unroll
                for(int i = 0; i < ksize; ++i)
                    sum = sum + sDataRow[i] * cLinearKernel[i];

                dst.ptr(y)[x] = saturate_cast<D>(sum);
            }
        }
    }

#pragma unroll
for(int i = 0; i < ksize; ++i)
     sum = sum + sDataRow[i] * cLinearKernel[i];

./gauss image.jpg 500 out2.bmp fast
Renderer: Gallium 0.4 on NV98 (nouveau)
Driver: 2.1 Mesa 7.12-devel (git-d4fcf67)
nv50_screen_get_shader_param:197 - unknown PIPE_SHADER_CAP 19
Error while compiling shader blurx.frag: 0:15(6): error: type mismatch
0:18(25): error: Could not implicitly convert operands to relational operator
0:18(25): error: loop condition must be scalar boolean
0:20(13): error: Could not implicitly convert operands to arithmetic operator
0:21(15): error: Could not implicitly convert operands to arithmetic operator

neko@gdetotut ~/Загрузки/GPGPU gauss $ ./gauss image.jpg 500 out1.bmp
Renderer: ATI Mobility Radeon HD 4500 Series  (ATI Technologies Inc.)
Driver: 3.3.11251 Compatibility Profile Context
Error while compiling shader blurx.frag: Fragment shader failed to compile with the following errors:
WARNING: 0:14: warning(#254) Signature matching could not take type conversions into account in implicit GLSL version number 110 Gauss
ERROR: 0:15: error(#160) Cannot convert from 'uniform float' to 'mediump int'
ERROR: error(#273) 1 compilation errors.  No code generated

$ nice --20 ./gauss image.jpg 500 out1.bmp
Renderer: GeForce GTS 450/PCI/SSE2 (NVIDIA Corporation)
Driver: 4.2.0 NVIDIA 290.06
Time: 0.150000 sec

$ nice --20 ./gauss image.jpg 500 out2.bmp fast
Renderer: GeForce GTS 450/PCI/SSE2 (NVIDIA Corporation)
Driver: 4.2.0 NVIDIA 290.06
Time: 0.020000 sec

$ ./gauss image.jpg 500 out1.bmp
Renderer: Mesa DRI Intel(R) Sandybridge Mobile  (Tungsten Graphics, Inc)
Driver: 2.1 Mesa 7.11.1
Time: 0.130000 sec