q_kernel.hpp

Go to the documentation of this file.

 
#ifndef PLSSVM_BACKENDS_SYCL_Q_KERNEL_HPP_
#define PLSSVM_BACKENDS_SYCL_Q_KERNEL_HPP_
#pragma once
 
#include "plssvm/constants.hpp"  // plssvm::kernel_index_type
 
#include "sycl/sycl.hpp"         // sycl::nd_item, sycl::pow, sycl::exp
 
namespace plssvm::sycl::detail {
 
template <typename T>
class device_kernel_q_linear {
  public:
    using real_type = T;
 
    device_kernel_q_linear(real_type *q, const real_type *data_d, const real_type *data_last, const kernel_index_type num_rows, const kernel_index_type feature_range) :
        q_{ q }, data_d_{ data_d }, data_last_{ data_last }, num_rows_{ num_rows }, feature_range_{ feature_range } {}
 
    void operator()(::sycl::id<1> index) const {
        real_type temp{ 0.0 };
        for (kernel_index_type i = 0; i < feature_range_; ++i) {
            temp += data_d_[i * num_rows_ + index] * data_last_[i];
        }
        q_[index] = temp;
    }
 
  private:
    real_type *q_;
    const real_type *data_d_;
    const real_type *data_last_;
    const kernel_index_type num_rows_;
    const kernel_index_type feature_range_;
};
 
template <typename T>
class device_kernel_q_polynomial {
  public:
    using real_type = T;
 
    device_kernel_q_polynomial(real_type *q, const real_type *data_d, const real_type *data_last, const kernel_index_type num_rows, const kernel_index_type num_cols, const int degree, const real_type gamma, const real_type coef0) :
        q_{ q }, data_d_{ data_d }, data_last_{ data_last }, num_rows_{ num_rows }, num_cols_{ num_cols }, degree_{ degree }, gamma_{ gamma }, coef0_{ coef0 } {}
 
    void operator()(::sycl::id<1> index) const {
        real_type temp{ 0.0 };
        for (kernel_index_type i = 0; i < num_cols_; ++i) {
            temp += data_d_[i * num_rows_ + index] * data_last_[i];
        }
        q_[index] = ::sycl::pow(gamma_ * temp + coef0_, static_cast<real_type>(degree_));
    }
 
  private:
    real_type *q_;
    const real_type *data_d_;
    const real_type *data_last_;
    const kernel_index_type num_rows_;
    const kernel_index_type num_cols_;
    const int degree_;
    const real_type gamma_;
    const real_type coef0_;
};
 
template <typename T>
class device_kernel_q_rbf {
  public:
    using real_type = T;
 
    device_kernel_q_rbf(real_type *q, const real_type *data_d, const real_type *data_last, const kernel_index_type num_rows, const kernel_index_type num_cols, const real_type gamma) :
        q_{ q }, data_d_{ data_d }, data_last_{ data_last }, num_rows_{ num_rows }, num_cols_{ num_cols }, gamma_{ gamma } {}
 
    void operator()(::sycl::id<1> index) const {
        real_type temp{ 0.0 };
        for (kernel_index_type i = 0; i < num_cols_; ++i) {
            temp += (data_d_[i * num_rows_ + index] - data_last_[i]) * (data_d_[i * num_rows_ + index] - data_last_[i]);
        }
        q_[index] = ::sycl::exp(-gamma_ * temp);
    }
 
  private:
    real_type *q_;
    const real_type *data_d_;
    const real_type *data_last_;
    const kernel_index_type num_rows_;
    const kernel_index_type num_cols_;
    const real_type gamma_;
};
 
}  // namespace plssvm::sycl::detail
 
#endif  // PLSSVM_BACKENDS_SYCL_Q_KERNEL_HPP_