text: // fareed_barrier.cu.cc
  #ifdef GOOGLE_CUDA
  #define EIGEN_USE_GPU
  #include "fareed_barrier.h"
  #include "tensorflow/core/util/gpu_kernel_helper.h"
  
  using namespace tensorflow;
  
  using GPUDevice = Eigen::GpuDevice;
  
  // Define the CUDA kernel.
  template <typename T>
  __global__ void ExampleCudaKernel(const int size, const T* in, T* out) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < size;
         i += blockDim.x * gridDim.x) {
      out[i] = 2 * __ldg(in + i);
    }
  }
  
  // Define the GPU implementation that launches the CUDA kernel.
  template <typename T>
  void ExampleFunctor<GPUDevice, T>::operator()(
      const GPUDevice& d, int size, const T* in, T* out) {
    // Launch the cuda kernel.
    //
    // See core/util/gpu_kernel_helper.h for example of computing
    // block count and thread_per_block count.
    int block_count = 1024;
    int thread_per_block = 20;
    ExampleCudaKernel<T>
        <<<block_count, thread_per_block, 0, d.stream()>>>(size, in, out);
  }
  
  // Explicitly instantiate functors for the types of OpKernels registered.
  template struct ExampleFunctor<GPUDevice, float>;
  template struct ExampleFunctor<GPUDevice, int32>;
  
  #endif  // GOOGLE_CUDA