 //    break;
 //  }
 
-//  switch(interpolation)
-//  {
-//  default:
-//  case IU_INTERPOLATE_NEAREST:
-//  case IU_INTERPOLATE_LINEAR: // fallthrough intended
-    cuRemapKernel_8u_C1 <<< dimGridOut, dimBlock >>> (
-        dst->data(), dst->stride(), dst->width(), dst->height());
-//    break;
-//  }
+  switch(interpolation)
+  {
+  case IU_INTERPOLATE_LINEAR: // fallthrough intended
+    cuRemapLinearInterpKernel_8u_C1 <<< dimGridOut, dimBlock >>> (
+        dst->data(), dst->stride(), dst->width(), dst->height());
+    break;
+  default:
+  case IU_INTERPOLATE_NEAREST:
+    cuRemapPointInterpKernel_8u_C1 <<< dimGridOut, dimBlock >>> (
+        dst->data(), dst->stride(), dst->width(), dst->height());
+    break;
+  }
 
   cudaUnbindTexture(&tex1_8u_C1__);
   cudaUnbindTexture(&tex_remap_dx_32f_C1__);


//-----------------------------------------------------------------------------
// 8u_C1
__global__ void cuRemapKernel_8u_C1(unsigned char*dst, size_t stride, int width, int height)
{
  const int x = blockIdx.x*blockDim.x + threadIdx.x;
  const int y = blockIdx.y*blockDim.y + threadIdx.y;

  tex1_8u_C1__.addressMode[0] = cudaAddressModeClamp;
  tex1_8u_C1__.addressMode[1] = cudaAddressModeClamp;
  tex1_8u_C1__.normalized = false;
  tex1_8u_C1__.filterMode = cudaFilterModePoint;

  tex_remap_dx_32f_C1__.addressMode[0] = cudaAddressModeClamp;
  tex_remap_dx_32f_C1__.addressMode[1] = cudaAddressModeClamp;
