Add common interface for simple kernels

[upsj]

What?

This PR adds the capability to implement simple kernels with a uniform syntax across all executors:

exec->run_kernel([] GKO_KERNEL(auto row, auto col, auto alpha, auto x, auto y) {
                     x(row, col) = alpha(0, col) * x(row, col) + y(row, col);
                 }, x->get_size(), alpha, x, y);

Why?

This might help simplify our solver porting efforts for SYCL, and also allow us to execute custom operations on our data in examples transparently across executors, like I am doing in #701 on the CPU.

How?

It works in two three steps:

1. a uniform kernel dispatch approach for 1D and 2D kernels that call the provided lambda for all entries of the index range
2. mapping the input parameters alpha, x, y to intermediate device representations. The mappings right now are:
   • matrix::Dense<ValueType>* -> matrix_accessor<ValueType>
   • Array<ValueType>* -> ValueType*
   • other T* -> T* (the disambiguation here is currently a bit hacky and limited to arithmetic types)
3. Mapping the device representation to kernel parameters. The mapped parameters are then passed to the kernel lambda function. This is only relevant for compact(...) and vector(...) which omit (vector -> ValueType*) or infer the stride from kernel dimensions (compact -> matrix_accessor<ValueType>)

By wrapping the input objects in wrapper types, we could implement more specific semantics like
pointwise for (row, col), colwise for (0, col), rowwise for (row, 0) or scalar for (0, 0). This could simplify the example to

exec->run_kernel([] GKO_KERNEL(auto row, auto col, auto alpha, auto x, auto y) {
                     x = alpha * x + y;
                 }, x->get_size(), colwise(alpha), pointwise(x), pointwise(y));

or even

exec->run_kernel([] GKO_KERNEL(auto alpha, auto x, auto y) { x = alpha * x + y; },
                 x->get_size(), colwise(alpha), pointwise(x), pointwise(y));

if we added optional row_index() and column_index() pseudo parameters to run_kernel

I had to re-implement a tiny accessor replacement, since the value semantics implemented by gko::range don't work well in this generic context.

Finally, GKO_KERNEL is defined to be empty for CPU and SYCL or __device__ for CUDA and HIP.

TODO

• Fix nvcc flags on Windows
• Check if this has any performance impact:
  • The last time I checked this, I couldn't see any difference in code generation, but that might have changed?
  • We now specify a different stride for each matrix parameter. Does that matter, or is constant memory fast enough anyways? This is no longer true, we use compact representations now
• Add tests
  • Kernel launch
  • Stride handling
  • Division by zero handling
• Add documentation

[upsj]

The code that triggers the compiler error on Intel:

template <typename F, typename... Args>
void run(F f, Args... args) {
    #pragma omp parallel for
    for (unsigned i = 0; i < 100; i++) {
        f(args...);
    }
}

int main() {
    int a;
    run([](auto j) {}, a);
}

[yhmtsai]

I would not say this reduce the porting effort, but it really reduce a lot code from porting. Thanks!
Originally, our porting script always generates 1 ~ 2 layers to keep the same interface and template usage as CUDA/HIP. Thus, it will gives many layer just for internal, which is somehow annoying.
I think we can leave these kernels launch args to compiler decide and it should not give any performance impact.
Do you have any idea such that we do not need to copy the same thing over different executors?
Instantiating in core directly seems to be impossible because we need different compiler to compile it.
Is putting it in generic folder like common a possible way? but it will increase the complexity to find the actual kernel code

[upsj]

@yhmtsai Yes, that was the idea, we could basically have a core/device/simple_headers.hpp file that includes the correct executor's headers etc. dependent on which executor is being compiled.

[sonarcloud]

Kudos, SonarCloud Quality Gate passed!

0 Bugs
0 Vulnerabilities
0 Security Hotspots
1 Code Smell

82.6% Coverage
7.3% Duplication

[codecov]

Codecov Report

Merging #733 (29af274) into develop (d169abf) will decrease coverage by 1.27%.
The diff coverage is 96.14%.

@@             Coverage Diff             @@
##           develop     #733      +/-   ##
===========================================
- Coverage    94.38%   93.11%   -1.28%     
===========================================
  Files          400      408       +8     
  Lines        32253    32621     +368     
===========================================
- Hits         30441    30374      -67     
- Misses        1812     2247     +435     

Impacted Files	Coverage Δ
include/ginkgo/core/base/math.hpp	100.00% <ø> (ø)
include/ginkgo/core/matrix/dense.hpp	95.12% <ø> (ø)
omp/matrix/dense_kernels.cpp	97.61% <ø> (-0.71%)	⬇️
omp/test/matrix/dense_kernels.cpp	99.72% <ø> (-0.10%)	⬇️
test/utils/executor.hpp	14.28% <14.28%> (ø)
core/test/matrix/dense.cpp	73.52% <75.00%> (+0.06%)	⬆️
omp/base/kernel_launch.hpp	89.58% <89.58%> (ø)
common/base/kernel_launch_solver.hpp	90.90% <90.90%> (ø)
common/solver/fcg_kernels.cpp	93.33% <93.33%> (ø)
test/solver/ir_kernels.cpp	96.39% <93.75%> (ø)
... and 39 more

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update d169abf...29af274. Read the comment docs.