#include <algorithm>
#include <vector>

#include "caffe/layer.hpp"
#include "caffe/vision_layers.hpp"

namespace caffe {

template <typename Dtype>
__global__ void PairLabelForward(const int n, 
    const Dtype* in1, const Dtype* in2, Dtype* out) {
  CUDA_KERNEL_LOOP(index, n) {
    out[index] = (in1[index] == in2[index]) ? 1 : 0;
  }
}

template <typename Dtype>
void PairLabelLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
    vector<Blob<Dtype>*>* top) {
  const Dtype* bottom_data1 = bottom[0]->gpu_data();
  const Dtype* bottom_data2 = bottom[1]->gpu_data();
  Dtype* top_data = (*top)[0]->mutable_gpu_data();
  const int count = bottom[0]->count();
  // NOLINT_NEXT_LINE(whitespace/operators)
  PairLabelForward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
      count, bottom_data1, bottom_data2, top_data);
  CUDA_POST_KERNEL_CHECK;
}


INSTANTIATE_CLASS(PairLabelLayer);


}  // namespace caffe
