 
     surf2DLayeredwrite( c, dog_data, idx*4, idy, level, cudaBoundaryModeZero );
 }
+#endif
 
 } // namespace gauss
 } // namespace v11

    cudaEvent_t  ev     = oct_obj.getEventGaussDone( level-1 );
    cudaStreamWaitEvent( stream, ev, 0 );

    gauss::v11::make_dog
        <<<grid,block,0,stream>>>
        ( oct_obj._data_tex[level],
          oct_obj._data_tex[level-1],
          oct_obj.getDogSurface( ),
          level-1 );
}

/*************************************************************

}


#if 1
__global__
void make_dog( Plane2D_float       this_data,
               Plane2D_float       top_data,
               cudaSurfaceObject_t dog_data,
               int                 level )
{
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const int idy = blockIdx.y * blockDim.y + threadIdx.y;

    const int cols = this_data.getWidth();
    const int rows = this_data.getHeight();
    
    const int r_x = clamp( idx, cols );
    const int r_y = clamp( idy, rows );

    const float b = this_data.ptr(r_y)[r_x];
    const float a = top_data .ptr(r_y)[r_x];
    const float c = b - a; // c = fabs( a - b );

    surf2DLayeredwrite( c, dog_data, idx*4, idy, level, cudaBoundaryModeZero );
}
#else
__global__
void make_dog( cudaTextureObject_t this_data,
               cudaTextureObject_t top_data,
