ocl_kernels_hyb.hpp

Go to the documentation of this file.

#ifndef PARALUTION_OCL_KERNELS_HYB_HPP_
#define PARALUTION_OCL_KERNELS_HYB_HPP_

namespace paralution {

const char *ocl_kernels_hyb =
        "__kernel void kernel_ell_nnz_coo(const int nrow, const int max_row,\n"
        "                                 __global const int *row_offset, __global int *nnz_coo) {\n"
        "\n"
        "  int gid = get_global_id(0);\n"
        "\n"
        "  if (gid < nrow) {\n"
        "\n"
        "    nnz_coo[gid] = 0;\n"
        "    int nnz_per_row = row_offset[gid+1] - row_offset[gid];\n"
        "\n"
        "    if (nnz_per_row > max_row)\n"
        "      nnz_coo[gid] = nnz_per_row - max_row;\n"
        "\n"
        "  }\n"
        "\n"
        "}\n"
        "\n"
        "__kernel void kernel_ell_fill_ell(const int nrow, const int max_row,\n"
        "                                  __global const int *row_offset, __global const int *col,\n"
        "                                  __global const ValueType *val, __global int *ELL_col,\n"
        "                                  __global ValueType *ELL_val, __global int *nnz_ell) {\n"
        "\n"
        "  int gid = get_global_id(0);\n"
        "\n"
        "  if (gid < nrow) {\n"
        "\n"
        "    int n = 0;\n"
        "\n"
        "    for (int i=row_offset[gid]; i<row_offset[gid+1]; ++i) {\n"
        "\n"
        "      if (n >= max_row) break;\n"
        "\n"
        "      int idx = n * nrow + gid;\n"
        "\n"
        "      ELL_col[idx] = col[i];\n"
        "      ELL_val[idx] = val[i];\n"
        "\n"
        "      ++n;\n"
        "\n"
        "    }\n"
        "\n"
        "    nnz_ell[gid] = n;\n"
        "\n"
        "  }\n"
        "\n"
        "}\n"
        "\n"
        "__kernel void kernel_ell_fill_coo(const int nrow, __global const int *row_offset,\n"
        "                                  __global const int *col, __global const ValueType *val,\n"
        "                                  __global const int *nnz_coo, __global const int *nnz_ell,\n"
        "                                  __global int *COO_row, __global int *COO_col, __global ValueType *COO_val) {\n"
        "\n"
        "  int gid = get_global_id(0);\n"
        "\n"
        "  if (gid < nrow) {\n"
        "\n"
        "    int row_ptr = row_offset[gid+1];\n"
        "\n"
        "    for (int i=row_ptr - nnz_coo[gid]; i<row_ptr; ++i) {\n"
        "\n"
        "      int idx = i - nnz_ell[gid];\n"
        "\n"
        "      COO_row[idx] = gid;\n"
        "      COO_col[idx] = col[i];\n"
        "      COO_val[idx] = val[i];\n"
        "\n"
        "    }\n"
        "\n"
        "  }\n"
        "\n"
        "}\n"
        "\n"
;
}

#endif // PARALUTION_OCL_KERNELS_HYB_HPP_