                         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 );
}

#ifdef USE_DOG_ARRAY
__global__
void filter_gauss_vert_v11_dog( cudaTextureObject_t src_data,
                                Plane2D_float       dst_data,

                      _octaves[octave]._data_tex[level-1],
                      _octaves[octave].getDogSurface( ),
                      level-1 );
            }
        }
    }
