struct indirect_pointer_t {
  void *ptr;
  size_t value_size;

  indirect_pointer_t(cccl_iterator_t &it) : ptr(&it.state), value_size(it.value_type.size) {
     assert( it.type == cccl_iterator_kind_t::CCCL_POINTER );
  }

  indirect_pointer_t& operator+=(uint64_t offset) {
    char **ptr_ptr = reinterpret_cast<char **>(ptr);
    *ptr_ptr += (offset * value_size);
    return *this;
  }

  void* operator&() const
  {
    return ptr;
  }

};

struct indirect_iterator_t {
    // -----
    // type definitions
    // ----
    void *ptr;
    difference_type *offset_ptr;

    indirect_iterator_t(cccl_iterator_t &it) : ptr() {
      size_t offset_offset = align_up(it.size, sizeof(difference_type));
      size_t combined_nbytes = offset_offset + sizeof(difference_type); 
      // allocate memory for user-defined state followed by the offset 
      ptr = calloc(combined_nbytes);
      // copy content of state from `cccl_iterator_t` to allocation
      ::memcpy(ptr, it.state, it.size);
      // set offset to zero
      offset_ptr = ptr + offset_offset;
    }

    ~indirect_iterator_t () noexcept {
      // deallocate memory 
      free(ptr);
    }

    indirect_iterator_t& operator+=(difference_type offset) {
      *offset_ptr += offset;
      return *this;
    }
}

struct __align__(ALIGNMENT) input_iterator_t {
  // type definitions

  __device__ inline value_type operator*() const { 
       const input_iterator_t &it = (*this + *offset_ptr);
       return DEREF(it.data);
  }
  __device__ inline input_iterator_t& operator+=(difference_type diff) {
      ADVANCE(data, diff);
      return *this;
  }
  __device__ inline value_type operator[](difference_type diff) const {
      return *(*this + diff);
  }
  __device__ inline input_iterator_t operator+(difference_type diff) const {
      input_iterator_t result = *this;
      result += diff;
      return result;
  }

   char data[ STATE_SIZE ];
   int64_t offset;
};