                     ( _octaves[octave]._interm_data_tex,
                       _octaves[octave].getData( level ) );
             } else {
-#ifdef USE_DOG_ARRAY
                 filter_gauss_vert_v11_dog
                     <<<d_grid,d_block>>>
                     ( _octaves[octave]._interm_data_tex,

                        idx*4, idy, level,
                        cudaBoundaryModeZero );
}
#else // not USE_DOG_ARRAY
__global__
void filter_gauss_vert_v11_dog( cudaTextureObject_t src_data,
                                Plane2D_float       dst_data,
                                cudaTextureObject_t top_data,
                                Plane2D_float       dog_data )
{
    float b = filter_gauss_vert_v11_sub( src_data, dst_data );

    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const int idy = blockIdx.y * blockDim.y + threadIdx.y;

    float a;
    a = tex2D<float>( top_data, idx, idy );
    a = fabs( a - b );

    const int width  = dog_data.getWidth();
    const int height = dog_data.getHeight();
    if( idx >= width ) return;
    if( idy >= height ) return;

    dog_data.ptr(idy)[idx] = a;
}
#endif // not USE_DOG_ARRAY

/*************************************************************
 * V11: host side

    filter_gauss_vert_v11_sub( src_data, dst_data );
}

__global__
void filter_gauss_vert_v11_dog( cudaTextureObject_t src_data,
                                Plane2D_float       dst_data,
