gpumat.hh

// This file is part of fml which is released under the Boost Software
// License, Version 1.0. See accompanying file LICENSE or copy at
// https://www.boost.org/LICENSE_1_0.txt
 
#ifndef FML_GPU_GPUMAT_H
#define FML_GPU_GPUMAT_H
#pragma once
 
 
#include <cstdint>
 
#include "../_internals/print.hh"
#include "../_internals/rand.hh"
#include "../_internals/types.hh"
#include "../_internals/unimat.hh"
 
#include "arch/arch.hh"
 
#include "internals/gpuscalar.hh"
#include "internals/kernelfuns.hh"
#include "internals/launcher.hh"
 
#include "card.hh"
#include "gpuvec.hh"
 
 
namespace fml
{
  template <typename REAL>
  class gpumat : public fml::unimat<REAL>
  {
    public:
      gpumat();
      gpumat(std::shared_ptr<card> gpu);
      gpumat(std::shared_ptr<card> gpu, len_t nrows, len_t ncols);
      gpumat(std::shared_ptr<card> gpu, REAL *data, len_t nrows, len_t ncols, bool free_on_destruct=false);
      gpumat(const gpumat &x);
      ~gpumat();
      
      void resize(len_t nrows, len_t ncols);
      void resize(std::shared_ptr<card> gpu, len_t nrows, len_t ncols);
      void inherit(std::shared_ptr<card> gpu, REAL *data, len_t nrows, len_t ncols, bool free_on_destruct=false);
      gpumat<REAL> dupe() const;
      
      void print(uint8_t ndigits=4, bool add_final_blank=true) const;
      void info() const;
      
      void fill_zero();
      void fill_val(const REAL v);
      void fill_linspace();
      void fill_linspace(const REAL start, const REAL stop);
      void fill_eye();
      void fill_diag(const gpuvec<REAL> &v);
      void fill_runif(const uint32_t seed, const REAL min=0, const REAL max=1);
      void fill_runif(const REAL min=0, const REAL max=1);
      void fill_rnorm(const uint32_t seed, const REAL mean=0, const REAL sd=1);
      void fill_rnorm(const REAL mean=0, const REAL sd=1);
      
      void diag(gpuvec<REAL> &v);
      void antidiag(gpuvec<REAL> &v);
      void scale(const REAL s);
      void rev_rows();
      void rev_cols();
      
      bool any_inf() const;
      bool any_nan() const;
      
      REAL get(const len_t i) const;
      REAL get(const len_t i, const len_t j) const;
      void set(const len_t i, const REAL v);
      void set(const len_t i, const len_t j, const REAL v);
      void get_row(const len_t i, gpuvec<REAL> &v) const;
      void set_row(const len_t i, const gpuvec<REAL> &v);
      void get_col(const len_t j, gpuvec<REAL> &v) const;
      void set_col(const len_t j, const gpuvec<REAL> &v);
      
      bool operator==(const gpumat<REAL> &x) const;
      bool operator!=(const gpumat<REAL> &x) const;
      gpumat<REAL>& operator=(const gpumat<REAL> &x);
      
      std::shared_ptr<card> get_card() const {return c;};
      dim3 get_blockdim() const {return dim_block;};
      dim3 get_griddim() const {return dim_grid;};
      
    protected:
      std::shared_ptr<card> c;
    
    private:
      dim3 dim_block;
      dim3 dim_grid;
      
      void free();
      void check_params(len_t nrows, len_t ncols);
      void check_gpu(std::shared_ptr<card> gpu);
  };
}
 
 
 
// -----------------------------------------------------------------------------
// public
// -----------------------------------------------------------------------------
 
// constructors/destructor
 
template <typename REAL>
fml::gpumat<REAL>::gpumat()
{
  this->c=NULL;
  this->m=0;
  this->n=0;
  this->data=NULL;
  this->free_data = false;
}
 
template <typename REAL>
fml::gpumat<REAL>::gpumat(std::shared_ptr<fml::card> gpu)
{
  check_gpu(gpu);
  
  this->c = gpu;
  
  this->m = 0;
  this->n = 0;
  this->data = NULL;
  
  this->free_data = true;
}
 
 
 
template <typename REAL>
fml::gpumat<REAL>::gpumat(std::shared_ptr<fml::card> gpu, len_t nrows, len_t ncols)
{
  check_params(nrows, ncols);
  check_gpu(gpu);
  
  this->c = gpu;
  
  const size_t len = (size_t) nrows * ncols * sizeof(REAL);
  this->data = (REAL*) this->c->mem_alloc(len);
  
  this->m = nrows;
  this->n = ncols;
  
  dim_block = fml::kernel_launcher::dim_block2();
  dim_grid = fml::kernel_launcher::dim_grid(this->m, this->n);
  
  this->free_data = true;
}
 
 
 
template <typename REAL>
fml::gpumat<REAL>::gpumat(std::shared_ptr<fml::card> gpu, REAL *data_, len_t nrows, len_t ncols, bool free_on_destruct)
{
  check_params(nrows, ncols);
  check_gpu(gpu);
  
  this->c = gpu;
  
  this->m = nrows;
  this->n = ncols;
  this->data = data_;
  
  dim_block = fml::kernel_launcher::dim_block2();
  dim_grid = fml::kernel_launcher::dim_grid(this->m, this->n);
  
  this->free_data = free_on_destruct;
}
 
 
 
template <typename REAL>
fml::gpumat<REAL>::gpumat(const gpumat<REAL> &x)
{
  this->m = x.nrows();
  this->n = x.ncols();
  this->data = x.data_ptr();
  
  dim_block = fml::kernel_launcher::dim_block2();
  dim_grid = fml::kernel_launcher::dim_grid(this->m, this->n);
  
  this->c = x.get_card();
  
  this->free_data = false;
}
 
 
 
template <typename REAL>
fml::gpumat<REAL>::~gpumat()
{
  this->free();
  c = NULL;
}
 
 
 
// memory management
 
template <typename REAL>
void fml::gpumat<REAL>::resize(len_t nrows, len_t ncols)
{
  check_params(nrows, ncols);
  
  const size_t len = (size_t) nrows * ncols * sizeof(REAL);
  const size_t oldlen = (size_t) this->m * this->n * sizeof(REAL);
  
  if (len == oldlen)
  {
    this->m = nrows;
    this->n = ncols;
    return;
  }
  
  REAL *realloc_ptr;
  realloc_ptr = (REAL*) this->c->mem_alloc(len);
  
  const size_t copylen = std::min(len, oldlen);
  this->c->mem_gpu2gpu(realloc_ptr, this->data, copylen);
  this->c->mem_free(this->data);
  this->data = realloc_ptr;
  
  this->m = nrows;
  this->n = ncols;
  
  dim_block = fml::kernel_launcher::dim_block2();
  dim_grid = fml::kernel_launcher::dim_grid(this->m, this->n);
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::resize(std::shared_ptr<fml::card> gpu, len_t nrows, len_t ncols)
{
  check_gpu(gpu);
  
  this->c = gpu;
  this->resize(nrows, ncols);
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::inherit(std::shared_ptr<fml::card> gpu, REAL *data, len_t nrows, len_t ncols, bool free_on_destruct)
{
  check_params(nrows, ncols);
  check_gpu(gpu);
  
  this->free();
  
  this->c = gpu;
  
  this->m = nrows;
  this->n = ncols;
  this->data = data;
  
  dim_block = fml::kernel_launcher::dim_block2();
  dim_grid = fml::kernel_launcher::dim_grid(this->m, this->n);
  
  this->free_data = free_on_destruct;
}
 
 
 
template <typename REAL>
fml::gpumat<REAL> fml::gpumat<REAL>::dupe() const
{
  fml::gpumat<REAL> cpy(this->c, this->m, this->n);
  
  const size_t len = (size_t) this->m * this->n * sizeof(REAL);
  this->c->mem_gpu2gpu(cpy.data_ptr(), this->data, len);
  
  return cpy;
}
 
 
 
// printers
 
template <typename REAL>
void fml::gpumat<REAL>::print(uint8_t ndigits, bool add_final_blank) const
{
  for (int i=0; i<this->m; i++)
  {
    for (int j=0; j<this->n; j++)
    {
      REAL tmp;
      this->c->mem_gpu2cpu(&tmp, this->data + (i + this->m*j), sizeof(REAL));
      this->printval(tmp, ndigits);
    }
  
    fml::print::putchar('\n');
  }
  
  if (add_final_blank)
    fml::print::putchar('\n');
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::info() const
{
  fml::print::printf("# gpumat ");
  fml::print::printf("%dx%d ", this->m, this->n);
  fml::print::printf("type=%s ", typeid(REAL).name());
  fml::print::printf("\n");
}
 
 
 
// fillers
 
template <typename REAL>
void fml::gpumat<REAL>::fill_zero()
{
  const size_t len = (size_t) this->m * this->n * sizeof(REAL);
  this->c->mem_set(this->data, 0, len);
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::fill_val(const REAL v)
{
  fml::kernelfuns::kernel_fill_val<<<dim_grid, dim_block>>>(v, this->m, this->n, this->data);
  this->c->check();
}
 
 
 
template <typename T>
void fml::gpumat<T>::fill_linspace()
{
  T start = 1;
  T stop = (T) (this->m * this->n);
  this->fill_linspace(start, stop);
}
 
template <typename REAL>
void fml::gpumat<REAL>::fill_linspace(REAL start, REAL stop)
{
  // if (start == stop)
  //   this->fill_val(start);
  // else
  // {
    fml::kernelfuns::kernel_fill_linspace<<<dim_grid, dim_block>>>(start, stop, this->m, this->n, this->data);
    this->c->check();
  // }
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::fill_eye()
{
  fml::kernelfuns::kernel_fill_eye<<<dim_grid, dim_block>>>(this->m, this->n, this->data);
  this->c->check();
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::fill_diag(const gpuvec<REAL> &v)
{
  fml::kernelfuns::kernel_fill_diag<<<dim_grid, dim_block>>>(v.size(), v.data_ptr(), this->m, this->n, this->data);
  this->c->check();
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::fill_runif(const uint32_t seed, const REAL min, const REAL max)
{
  const size_t len = (size_t) this->m * this->n;
  gpurand::gen_runif(seed, len, this->data);
  fml::kernelfuns::kernel_fill_runif_update<<<dim_grid, dim_block>>>(min, max, this->m, this->n, this->data);
  this->c->check();
}
 
template <typename REAL>
void fml::gpumat<REAL>::fill_runif(const REAL min, const REAL max)
{
  uint32_t seed = fml::rand::get_seed();
  this->fill_runif(seed, min, max);
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::fill_rnorm(const uint32_t seed, const REAL mean, const REAL sd)
{
  const size_t len = (size_t) this->m * this->n;
  gpurand::gen_rnorm(seed, mean, sd, len, this->data);
}
 
template <typename REAL>
void fml::gpumat<REAL>::fill_rnorm(const REAL mean, const REAL sd)
{
  uint32_t seed = fml::rand::get_seed();
  this->fill_rnorm(mean, sd);
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::diag(gpuvec<REAL> &v)
{
  const len_t minmn = std::min(this->m, this->n);
  v.resize(minmn);
  
  fml::kernelfuns::kernel_diag<<<dim_grid, dim_block>>>(this->m, this->n, this->data, v.data_ptr());
  this->c->check();
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::antidiag(gpuvec<REAL> &v)
{
  const len_t minmn = std::min(this->m, this->n);
  v.resize(minmn);
  
  fml::kernelfuns::kernel_antidiag<<<dim_grid, dim_block>>>(this->m, this->n, this->data, v.data_ptr());
  this->c->check();
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::scale(const REAL s)
{
  fml::kernelfuns::kernel_scale<<<dim_grid, dim_block>>>(s, this->m, this->n, this->data);
  this->c->check();
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::rev_rows()
{
  fml::kernelfuns::kernel_rev_rows<<<dim_grid, dim_block>>>(this->m, this->n, this->data);
  this->c->check();
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::rev_cols()
{
  fml::kernelfuns::kernel_rev_cols<<<dim_grid, dim_block>>>(this->m, this->n, this->data);
  this->c->check();
}
 
 
 
template <typename REAL>
bool fml::gpumat<REAL>::any_inf() const
{
  int has_inf = 0;
  gpuscalar<int> has_inf_gpu(c, has_inf);
  
  fml::kernelfuns::kernel_any_inf<<<dim_grid, dim_block>>>(this->m, this->n, this->data, has_inf_gpu.data_ptr());
  
  has_inf_gpu.get_val(&has_inf);
  this->c->check();
  
  return (bool) has_inf;
}
 
 
 
template <typename REAL>
bool fml::gpumat<REAL>::any_nan() const
{
  int has_nan = 0;
  gpuscalar<int> has_nan_gpu(c, has_nan);
  
  fml::kernelfuns::kernel_any_nan<<<dim_grid, dim_block>>>(this->m, this->n, this->data, has_nan_gpu.data_ptr());
  
  has_nan_gpu.get_val(&has_nan);
  this->c->check();
  
  return (bool) has_nan;
}
 
 
 
// operators
 
template <typename REAL>
REAL fml::gpumat<REAL>::get(const len_t i) const
{
  this->check_index(i);
 
  REAL ret;
  this->c->mem_gpu2cpu(&ret, this->data + i, sizeof(REAL));
  return ret;
}
 
template <typename REAL>
REAL fml::gpumat<REAL>::get(const len_t i, const len_t j) const
{
  this->check_index(i, j);
 
  REAL ret;
  this->c->mem_gpu2cpu(&ret, this->data + (i + this->m*j), sizeof(REAL));
  return ret;
}
 
template <typename REAL>
void fml::gpumat<REAL>::set(const len_t i, const REAL v)
{
  this->check_index(i);
  this->c->mem_cpu2gpu(this->data + i, &v, sizeof(REAL));
}
 
template <typename REAL>
void fml::gpumat<REAL>::set(const len_t i, const len_t j, const REAL v)
{
  this->check_index(i, j);
  this->c->mem_cpu2gpu(this->data + (i + this->m*j), &v, sizeof(REAL));
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::get_row(const len_t i, fml::gpuvec<REAL> &v) const
{
  if (i < 0 || i >= this->m)
    throw std::logic_error("invalid matrix row");
  
  v.resize(this->n);
  
  fml::kernelfuns::kernel_get_row<<<dim_grid, dim_block>>>(i, this->m, this->n, this->data, v.data_ptr());
  this->c->check();
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::set_row(const len_t i, const fml::gpuvec<REAL> &v)
{
  if (i < 0 || i >= this->m)
    throw std::logic_error("invalid matrix row");
  if (v.size() != this->n)
    throw std::runtime_error("non-conformable arguments");
  
  fml::kernelfuns::kernel_set_row<<<dim_grid, dim_block>>>(i, this->m, this->n, this->data, v.data_ptr());
  this->c->check();
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::get_col(const len_t j, fml::gpuvec<REAL> &v) const
{
  if (j < 0 || j >= this->n)
    throw std::logic_error("invalid matrix column");
  
  v.resize(this->m);
  
  fml::kernelfuns::kernel_get_col<<<dim_grid, dim_block>>>(j, this->m, this->n, this->data, v.data_ptr());
  this->c->check();
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::set_col(const len_t i, const fml::gpuvec<REAL> &v)
{
  if (i < 0 || i >= this->n)
    throw std::logic_error("invalid matrix row");
  if (v.size() != this->m)
    throw std::runtime_error("non-conformable arguments");
  
  fml::kernelfuns::kernel_set_col<<<dim_grid, dim_block>>>(i, this->m, this->n, this->data, v.data_ptr());
  this->c->check();
}
 
 
 
template <typename T>
bool fml::gpumat<T>::operator==(const fml::gpumat<T> &x) const
{
  if (this->m != x.nrows() || this->n != x.ncols())
    return false;
  else if (this->c->get_id() != x.get_card()->get_id())
    return false;
  else if (this->data == x.data_ptr())
    return true;
  
  int all_eq = 1;
  fml::gpuscalar<int> all_eq_gpu(c, all_eq);
  
  fml::kernelfuns::kernel_all_eq<<<dim_grid, dim_block>>>(this->m, this->n, this->data, x.data_ptr(), all_eq_gpu.data_ptr());
  
  all_eq_gpu.get_val(&all_eq);
  this->c->check();
  
  return (bool) all_eq;
}
 
template <typename T>
bool fml::gpumat<T>::operator!=(const fml::gpumat<T> &x) const
{
  return !(*this == x);
}
 
 
 
template <typename REAL>
fml::gpumat<REAL>& fml::gpumat<REAL>::operator=(const fml::gpumat<REAL> &x)
{
  this->c = x.get_card();
  
  this->m = x.nrows();
  this->n = x.ncols();
  this->data = x.data_ptr();
  
  this->free_data = false;
  return *this;
}
 
 
 
// -----------------------------------------------------------------------------
// private
// -----------------------------------------------------------------------------
 
template <typename REAL>
void fml::gpumat<REAL>::free()
{
  if (this->free_data && this->data)
  {
    this->c->mem_free(this->data);
    this->data = NULL;
  }
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::check_params(len_t nrows, len_t ncols)
{
  if (nrows < 0 || ncols < 0)
    throw std::runtime_error("invalid dimensions");
}
 
 
 
template <typename REAL>
void fml::gpumat<REAL>::check_gpu(std::shared_ptr<fml::card> gpu)
{
  if (!gpu->valid_card())
    throw std::runtime_error("GPU card object is invalid");
}
 
 
#endif