 }
 
 surface<void, cudaSurfaceType2D> image;
-//texture<char, 4, cudaReadModeNormalizedFloat> image;
 
 __global__ void toGrayKernel(int image_size)
 {

  if (x >= image_size || y >= image_size)
    return;

  //int index = y * image_size + x;

  //char4 color = out[index];
  char4 color;
  surf2Dread(&color, image, x * 4, y);


  color.y = gray;
  color.z = gray;
  surf2Dwrite(color, image, x * 4, y);
}

void toGray(cudaGraphicsResource_t *inputImage, int image_size)
{
  HANDLE_ERROR(cudaGraphicsMapResources(1, inputImage));
  cudaArray_t input_array;
  HANDLE_ERROR(
      cudaGraphicsSubResourceGetMappedArray(&input_array, *inputImage, 0, 0));
  cudaChannelFormatDesc channeldesc;
  HANDLE_ERROR(cudaGetChannelDesc(&channeldesc, input_array));

  HANDLE_ERROR(cudaBindSurfaceToArray(image, input_array, channeldesc));

  dim3 dimBlock(32, 32, 1);
  dim3 dimGrid(divUp(image_size, dimBlock.x), divUp(image_size, dimBlock.y), 1);

  toGrayKernel << <dimGrid, dimBlock>>> (image_size);
  HANDLE_ERROR(cudaThreadSynchronize());
  HANDLE_ERROR(cudaGraphicsUnmapResources(1, inputImage));
}

