  */
 __device__ float4 rgb(int maxValue, int value)
 {
-    float ratio = 2.0f * value / maxValue;
-    float b = max(0.0f, 1 - ratio);  // NOLINT
-    float r = max(0.0f, ratio - 1);  // NOLINT
-    float g = max(0.0f, 1.0f - b - r);  // NOLINT
-
-    return make_float4(r, g, b, 1.0);
+  float ratio = 2.0f * value / maxValue;
+  float b = max(0.0f, 1 - ratio);     // NOLINT
+  float r = max(0.0f, ratio - 1);     // NOLINT
+  float g = max(0.0f, 1.0f - b - r);  // NOLINT
+
+  return make_float4(r, g, b, 1.0);
 }
 
 __global__ void gather(cudaSurfaceObject_t output, int imageSize,

namespace Placement
{

  Apollonius::Apollonius(std::shared_ptr<CudaArrayProvider> distancesImage,
    std::shared_ptr<CudaArrayProvider> outputImage,
    std::vector<Eigen::Vector4f> labelPositions,
    int labelCount)
    : distancesImage(distancesImage), outputImage(outputImage),
    labelCount(labelCount), seedBuffer(labelCount)
  {
    imageSize = outputImage->getWidth();

    for (size_t i = 0; i < seedBuffer.size(); ++i)
    {
      auto labelPosition = labelPositions[i];
      seedBuffer[i] = make_float4(labelPosition.x(), labelPosition.y(),
        labelPosition.z(), labelPosition.w());
      int index = static_cast<int>(labelPosition.y()) +
        static_cast<int>(labelPosition.z()) * imageSize;
      pixelIndexToLabelId[index] = labelPosition.x();
    }

    pixelCount = imageSize * imageSize;

    distancesImage->map();
    auto resDesc = distancesImage->getResourceDesc();

    struct cudaTextureDesc texDesc;
    memset(&texDesc, 0, sizeof(texDesc));
    texDesc.addressMode[0] = cudaAddressModeWrap;
    texDesc.addressMode[1] = cudaAddressModeWrap;
    texDesc.filterMode = cudaFilterModeLinear;
    texDesc.readMode = cudaReadModeElementType;
    texDesc.normalizedCoords = 0;

    cudaCreateTextureObject(&distancesTexture, &resDesc, &texDesc, NULL);

    distancesImage->unmap();
  }

  void Apollonius::run()
  {
    resize();
    outputImage->map();
    auto resDesc = outputImage->getResourceDesc();
    cudaCreateSurfaceObject(&outputSurface, &resDesc);

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

    runSeedKernel();
    runStepsKernels();
    runGatherKernel();

    outputImage->unmap();
  }

  thrust::device_vector<int> &Apollonius::getIds()
  {
    return seedIds;
  }

  std::vector<int> Apollonius::calculateOrdering()
  {
    extractUniqueBoundaryIndices();
    updateLabelSeeds();

    size_t iterationCount = 0;
    size_t labelCount = pixelIndexToLabelId.size();
    while (extractedIndices.size() < labelCount && iterationCount < labelCount)
    {
      run();
      extractUniqueBoundaryIndices();
      updateLabelSeeds();

      ++iterationCount;
    }

    return std::vector<int>(insertionOrder.begin(), insertionOrder.end());
  }

  void Apollonius::resize()
  {
    if (computeVector.size() != static_cast<unsigned long>(pixelCount))
    {
      computeVector.resize(pixelCount, pixelCount);
    }

    if (seedIds.size() != labelCount || seedIndices.size() != labelCount)
    {
      seedIds.resize(labelCount, -1);
      seedIndices.resize(labelCount, -1);
    }
  }


  void Apollonius::extractUniqueBoundaryIndices()
  {
    const unsigned int borderSize = 4 * imageSize;

    if (orderedIndices.size() < borderSize)
    {
      orderedIndices.resize(borderSize);
    }

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

    int *computePtr = thrust::raw_pointer_cast(computeVector.data());
    int *orderedIndicesPtr = thrust::raw_pointer_cast(orderedIndices.data());

    copyBorderIndex << <dimGrid, dimBlock >> >(imageSize, computePtr,
      orderedIndicesPtr);

    HANDLE_ERROR(cudaThreadSynchronize());

    thrust::device_vector<int>::iterator it_found =
      thrust::unique(orderedIndices.begin(), orderedIndices.end());
    thrust::host_vector<int> uniqueIndices(orderedIndices.begin(), it_found);

    for (unsigned int i = 0; i < uniqueIndices.size(); i++)
    {
      const int uniqueIndex = uniqueIndices[i];
      if (uniqueIndex > 0)
      {
        if (extractedIndices.insert(uniqueIndex).second)
        {
          insertionOrder.push_front(pixelIndexToLabelId[uniqueIndex]);
        }
      }
    }
  }

  void Apollonius::updateLabelSeeds()
  {
    thrust::host_vector<float4> labelsSeed = seedBuffer;
    for (size_t index = 0; index < labelsSeed.size(); ++index)
    {
      float4 seed = labelsSeed[index];
      const uint cindex =
        static_cast<int>(seed.y) + static_cast<int>(seed.z) * imageSize;

      if (extractedIndices.find(cindex) != extractedIndices.end())
      {
        if (seed.x > 0)
        {
          seed.x = -seed.x;
          labelsSeed[index] = seed;
        }
      }
    }

    seedBuffer = labelsSeed;
  }

}  // namespace Placement


    }
  }

  float4 color =
      labelId == -1 ? make_float4(1, 1, 1, 1) : rgb(labelCount, labelId);

  surf2Dwrite(color, output, x * sizeof(float4), y);
}
