#include <iostream>
#include <typeinfo>

#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/RuntimeDeviceInformation.h>

#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
#include <vtkm/cont/serial/DeviceAdapterSerial.h>
#include <vtkm/cont/TryExecute.h>

#include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>

// Step 1: For each virtual-like method create a function (in a detail
// namespace) that takes a void pointer to a concrete class and then calls the
// method on the class after performing a reinterpret_cast. The function will
// certainly be templated, but should not be overloaded. You need to be able
// to fully specify the function without arguments.
template<typename PortalType>
__device__ __host__
vtkm::Id ArrayPortalVirtualGetNumberOfValues(void *portalPointer)
{
  PortalType* portalRef = reinterpret_cast<PortalType*>(portalPointer);
  return portalRef->GetNumberOfValues();
}


template<typename PortalType, typename ValueType>
__device__ __host__
ValueType ArrayPortalVirtualGet(void *portalPointer, vtkm::Id index)
{
  PortalType* portalRef = reinterpret_cast<PortalType*>(portalPointer);
  return static_cast<ValueType>(portalRef->Get(index));
}

template<typename PortalType, typename ValueType>
__device__ __host__
void ArrayPortalVirtualSet(void *portalPointer,
                           vtkm::Id index,
                           const ValueType &value)
{
  PortalType* portalRef = reinterpret_cast<PortalType*>(portalPointer);
  portalRef->Set(index, value);
}


template<typename PortalType, typename ValueType>
__global__ void construct_functions(void** numvalue, void** get, void** set)
{
  typedef vtkm::Id(*GetNumberOfValuesSignature)(void*);
  typedef ValueType(*GetSignature)(void *, vtkm::Id);
  typedef void(*SetSignature)(void *, vtkm::Id, const ValueType &);

  {
  GetNumberOfValuesSignature* n = reinterpret_cast<GetNumberOfValuesSignature*>(numvalue);
  *n = &ArrayPortalVirtualGetNumberOfValues<PortalType>;
  printf("From Device: n1 = %p \n", *n);
  }

  {
  GetSignature* n = reinterpret_cast<GetSignature*>(get);
  *n = &ArrayPortalVirtualGet<PortalType>;
  printf("From Device: n2 = %p \n", *n);
  }

  {
  SetSignature* n = reinterpret_cast<SetSignature*>(set);
  *n = &ArrayPortalVirtualSet<PortalType>;
  printf("From Device: n3 = %p \n", *n);
  }
}


// Step 2: Create copy functions for each portal type

template<typename PortalType, typename ValueType, typename NumValueSig, typename GetSig, typename SetSig>
inline
void CopyFunctionPointers(NumValueSig& destNumValuePointer,
                          GetSig& destGetPointer,
                          SetSig& destSetPointer,
                          vtkm::cont::DeviceAdapterTagCuda)
{
  NumValueSig device_NumValuePointer;
  cudaMalloc(&device_NumValuePointer, sizeof(NumValueSig));

  GetSig device_GetPointer;
  cudaMalloc(&device_GetPointer, sizeof(GetSig));

  SetSig device_SetPointer;
  cudaMalloc(&device_SetPointer, sizeof(SetSig));

  construct_functions<PortalType, ValueType><<<1,1>>>((void**)device_NumValuePointer,
                                                      (void**)device_GetPointer,
                                                      (void**)device_SetPointer);

  cudaDeviceSynchronize();
  cudaError_t error = cudaGetLastError();
  printf("CUDA state after construct launch: %s\n", cudaGetErrorString(error));

  cudaMemcpy((void**)&destNumValuePointer, (void*)device_NumValuePointer,
    sizeof(NumValueSig), cudaMemcpyDeviceToHost);

  cudaMemcpy((void**)&destGetPointer, (void*)device_GetPointer, sizeof(GetSig),
    cudaMemcpyDeviceToHost);

  cudaMemcpy((void**)&destSetPointer, (void*)device_SetPointer, sizeof(SetSig),
    cudaMemcpyDeviceToHost);

  cudaDeviceSynchronize();
  error = cudaGetLastError();
  printf("CUDA state after mempcy launch: %s\n", cudaGetErrorString(error));

  printf("From Host: n1 = %p\n", destNumValuePointer);
  printf("From Host: n2 = %p\n", destGetPointer);
  printf("From Host: n3 = %p\n", destSetPointer);

}

template<typename PortalType, typename ValueType, typename NumValueSig, typename GetSig, typename SetSig>
inline
void CopyFunctionPointers(NumValueSig& destNumValuePointer,
                          GetSig& destGetPointer,
                          SetSig& destSetPointer,
                          vtkm::cont::DeviceAdapterTagSerial)
{
  destNumValuePointer = &ArrayPortalVirtualGetNumberOfValues<PortalType>;
  destGetPointer = &ArrayPortalVirtualGet<PortalType, ValueType>;
  destSetPointer = &ArrayPortalVirtualSet<PortalType, ValueType>;
}

// Step 3: Create the struct or class that has virtual-like methods as normal.
// There should be no actual virtual methods.

template<typename ValueType>
class ArrayPortalVirtual
{

  // Step 4: For each virtual-like method, create a typedef for the signature
  // of the function accessor you created earlier in the class.
  typedef vtkm::Id(*GetNumberOfValuesSignature)(void*);
  typedef ValueType(*GetSignature)(void *, vtkm::Id);
  typedef void(*SetSignature)(void *, vtkm::Id, const ValueType &);

  // Step 5: For each virtual-like method, create a member variable in the
  // class with the type of the associated signature.

  GetNumberOfValuesSignature GetNumberOfValuesFunction;
  GetSignature GetFunction;
  SetSignature SetFunction;

  // Step 6: The class should also have a void* member variable for a
  // reference to the concrete class.

  void *SrcPortal;

public:

  // Step 7: Create a constructor for the class. The constructor should be
  // templated on the concrete class you are calling methods on as well as a
  // device adapter tag. One of the arguments should be a pointer to an
  // instance of the concrete class. The constructor should initialize its
  // function pointer members using CopyFunctionPointer.

  template<typename PortalType, typename Device>
   __host__
  ArrayPortalVirtual(PortalType *portal, Device)
    : SrcPortal(portal)
  {
    CopyFunctionPointers<PortalType,ValueType>( this->GetNumberOfValuesFunction,
                                                this->GetFunction,
                                                this->SetFunction,
                                                Device());
  }

  // Step 8: For each virtual-like method, create an implementation of that
  // method that calls the associated function pointer.
  __device__ __host__
  vtkm::Id GetNumberOfValues() const
  {
    return this->GetNumberOfValuesFunction(this->SrcPortal);
  }
  __device__ __host__
  ValueType Get(vtkm::Id index) const
  {
    return this->GetFunction(this->SrcPortal, index);
  }
  __device__ __host__
  void Set(vtkm::Id index, const ValueType &value) const
  {
    this->SetFunction(this->SrcPortal, index, value);
  }
};


__global__ void kernel(ArrayPortalVirtual<int>  ints,
                       ArrayPortalVirtual<float> floats)
{
 { printf("hello world\n"); }
 { printf("i value = %d", ints.Get(1)); }
 // { printf("f value = %d", floats.Get(1)); }
}



int main(int, char**)
{
  cudaError_t error;

  //make some cuda memory and map it to a portal
  using vtkm::exec::cuda::internal::ArrayPortalFromThrust;
  thrust::system::cuda::vector<double> data(1024);
  ArrayPortalFromThrust<double> portal(data.data(),data.data()+1023);


  //wrap that cuda memory in two virtual portals
  ArrayPortalVirtual<int> virt_int(&portal, vtkm::cont::DeviceAdapterTagCuda());
  ArrayPortalVirtual<float> virt_float(&portal, vtkm::cont::DeviceAdapterTagCuda());

  //need to launch a kernel to verify everything is working properly
  cudaDeviceSynchronize();
  error = cudaGetLastError();
  printf("CUDA state before kernel launch: %s\n", cudaGetErrorString(error));

  kernel<<<1,1>>>(virt_int, virt_float);
  cudaDeviceSynchronize();
  error = cudaGetLastError();
  printf("CUDA state after kernel launch: %s\n", cudaGetErrorString(error));

  return 0;
}