unpack.hpp

Go to the documentation of this file.

#ifndef STAN_MATH_OPENCL_KERNELS_UNPACK_HPP
#define STAN_MATH_OPENCL_KERNELS_UNPACK_HPP
#ifdef STAN_OPENCL
 
#include <stan/math/opencl/kernel_cl.hpp>
#include <stan/math/opencl/buffer_types.hpp>
#include <string>
 
namespace stan {
namespace math {
namespace opencl_kernels {
// \cond
static constexpr const char* unpack_kernel_code = STRINGIFY(
    // \endcond
    __kernel void unpack(__global double* B, __global double* A,
                         unsigned int rows, unsigned int cols,
                         unsigned int view) {
      int i = get_global_id(0);
      int j = get_global_id(1);
      if (i < rows && j < cols) {
        // the packed matrices are stored in row major
        if (view == LOWER) {
          const int column_offset = j * rows - (j * (j - 1)) / 2;
          const int row_offset = (i - j);
          if (j <= i) {
            B(i, j) = A[column_offset + row_offset];
          } else {
            B(i, j) = 0.0;
          }
        } else {
          const int column_offset = j * (j + 1) / 2;
          if (j >= i) {
            B(i, j) = A[column_offset + i];
          } else {
            B(i, j) = 0.0;
          }
        }
      }
    }
    // \cond
);
// \endcond
 
const kernel_cl<out_buffer, in_buffer, int, int, matrix_cl_view> unpack(
    "unpack", {indexing_helpers, unpack_kernel_code});
 
}  // namespace opencl_kernels
}  // namespace math
}  // namespace stan
#endif
#endif