Dpcpp port ELL and the porting script

[yhmtsai]

This PR adds the dpcpp ported ELL and the porting script
For review before merging dense, the https://github.com/ginkgo-project/ginkgo/pull/779/files/7ac6d0909f006b52c5516e63b50cf0dff6b18012..3fc8bea0d5d6f81ef4be661aade1254ad1d54355 changes only contains the ell part

TODO:

• clear some unused code in script
• add ConfigSelection into ELL (maybe not)
• merge dense

[yhmtsai]

format!

[tcojean]

For now, I only look into the oneAPI porting script to get that out of the
way. My main concern is to create a proper documentation of this as well as a
few fixes. I do a few tentative comments already in line in the files. I think
in addition to that, having a main README.md file which explains the
workflow and how the different parts are organized and interact with each
other would be useful.

[tcojean]

Here you mean that it doesn't work with x = y = 3 sort of cases?

[yhmtsai]

Yes, but I think the template will not have this case?

[tcojean]

continue?

[yhmtsai]

no, I need to reset the temp to empty

[tcojean]

Since the function is also in the other file, you could share this.

[tcojean]

This can be shared as well

[yhmtsai]

the situation is complex, and we do not add the config to cuda.
Thus, I will delete the adding Config part first.

[tcojean]

Extra comment? # Replace everything except begin or end characters, resp. (<[ and )>]

[tcojean]

Extra comment? # Check that there are as many begin as end characters

[tcojean]

Possible comment:
# Converts a CUDA kernel call to the DPC++ equivalent. Also takes care of the config selection calls and DPC++ queue manipulation if needed.

[tcojean]

Tentative other name:
convert_cuda_kernel_call

[tcojean]

Maybe a comment for the regex?

# [1]: kernel name
# [2]: template parameters if it exists
# [3]: CUDA parameters
# [4]: the rest of the function call (parameters, ...)

As a note, couldn't the template parameters call another templated struct? In
that case, maybe the check [^<>]* isn't warranted for the template portion?

[yhmtsai]

you are right, it can not handle call<t<1>, ..><<<>>>()
should we consider it? it maybe is possible to handle up to one level but arbitrary template is hard or impossible?

[tcojean]

Continued

[tcojean]

local content="${BASH_REMATCH[3]}"

[tcojean]

I think a comment is warranted here:
# Do not fall into the trap of counting parameters in a sub-expression, such as a function call inside the CUDA kernel parameters

[tcojean]

If I'm correct, all you do in the previous loop is to count the number of
variables in the list, and you reconstruct it all in $var? In that case,
isn't it possible to use three regex on $content to cleanup the annoying
parts before counting? Like:

temp=${content//\[*\]/}
temp=${temp//\(*\)/}
temp=${temp//<*>/}
num=$((${#temp//[^,]} + 1))
# You can also do checks for when num == 4

[tcojean]

This seems to come from somewhere else, maybe put it as a global variable or
function variable?

[tcojean]

local suffix_matches=$(echo "${suffix}" | wc -l)

[tcojean]

IFS=';' read -ra individual_device <<< "${DEVICE_FILE}"

[tcojean]

for variable in "${individual_device[@]}"; do

[tcojean]

I fail to see for now, what is this for?

[yhmtsai]

it is optional and off.
it will move the common code to common/dpcpp_code if enabling it.

[tcojean]

Maybe a comment?
# Integrate the new DPC++ file into the correct place

[tcojean]

target_file=$(echo "${input}" | sed 's:cuda/:dpcpp/:g;s:\.cuh:\.dp\.hpp:g;s:\.cu:\.dp\.cpp:g')

[codecov]

Codecov Report

Merging #779 (7dd6ef1) into develop (47ba37d) will increase coverage by 0.54%.
The diff coverage is 100.00%.

@@             Coverage Diff             @@
##           develop     #779      +/-   ##
===========================================
+ Coverage    93.71%   94.25%   +0.54%     
===========================================
  Files          408      408              
  Lines        32758    32756       -2     
===========================================
+ Hits         30698    30874     +176     
+ Misses        2060     1882     -178     

Impacted Files	Coverage Δ
accessor/range.hpp	100.00% <100.00%> (ø)
reference/factorization/par_ilut_kernels.cpp	100.00% <0.00%> (+0.49%)	⬆️
omp/factorization/par_ilut_kernels.cpp	100.00% <0.00%> (+100.00%)	⬆️

---

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 47ba37d...7dd6ef1. Read the comment docs.

[pratikvn]

This is a bit out of my depth, but I just had a cursory look. :)

[pratikvn]

This change would make the include/ginkgo/core/base/range.hpp diverge ? Maybe we should either note this somewhere and in 2.0 fix it, or adapt the include/ginkgo/core/base/range.hpp as well.

[yhmtsai]

what do you mean diverge here? the usage of constructor keeps the same due to the accessor always require more than 1 args

[pratikvn]

I believe this is a duplication of the include/ginkgo/core/base/range.hpp, as it was needed to be separate for the Accessor. I meant now the typically the same class (range) has two different versions and we should probably unify them. Also I think you are changing one of the copy constructors below ?

[yhmtsai]

Yes, it uses enable_if_t to avoid compiler use this constructor as copy constuctor. By doing so, compiler can auto generate the copy/move constructor, so we do not need to manually point copy constructor to default behavior.
I can also change the part in include/ginkgo/core/base/range.hpp

[pratikvn]

Yeah, I guess we need to be sure that it does not break interface, I would also maybe ping @thoasm for his comments here.

[upsj]

Can you give a bit of background why we need a copy constructor and a mutable lvalue ref constructor?

[yhmtsai]

disabling Outer(const Outer&)

error: no matching constructor for initialization of 'Outer'
       run([&]() { [=]() { func(a); }; });
                                ^
candidate constructor not viable: 1st argument ('const Outer') would lose const qualifier

disabling Outer(Outer&)

note: in instantiation of function template specialization 'Outer::Outer<Outer &>' requested here
  run([&]() { [=]() { func(a); }; });

this is simple version for emulating the dpcpp submit the job.
As far as I remember, two default impl out of class is okay in C++ code, but dpcpp can not allow it.

[upsj]

Can you post a longer snippet that contains the entire context? Why do you need to modify the element you are copying in a copy constructor?

[thoasm]

Right, we talked about something like that previously.
If I remember correctly, the problem is that you need both const range & and range & for them to be used, otherwise, for a range &, the forwarding constructor is taken. Additionally, for some reason, the visual studio compiler does not accept both copy constructors range(const range &) = default and range(range &) = default (apparently, implementing them instead of =default does not help as well).
Was that about right, @yhmtsai?
I would prefer to find an alternative solution and will try something tomorrow, otherwise, we can add the SFINAE work-around to only exclude single range parameter for this constructor.

[yhmtsai]

this is the small example

#include <CL/sycl.hpp>
#include <utility>
class Inner {
public:
    explicit Inner(int, int) {}
};

class Outer {
public:
    template <typename... T>
    constexpr explicit Outer(T &&... inner) : inner_{std::forward<T>(inner)...}
    {}
    Outer(Outer &&) = default;
    // putting default inside class works
    // Outer(Outer &) = default;
    // Outer(Outer const&) = default;
    Outer(Outer &);
    Outer(Outer const &);
private:
    Inner inner_;
};

// putting outside fails
Outer::Outer(Outer &) = default;
Outer::Outer(Outer const &) = default;

void func(Outer a) {}

template <typename Lambda>
void run(Lambda f)
{
    f();
}

int main()
{
    auto a = Outer(3, 4);
    // simulate the submit
    run([&]() { [=]() { func(a); }; });
    sycl::queue q_ct1{sycl::property::queue::in_order{}};
    // real submit
    q_ct1.submit([&](sycl::handler &cgh) {
        cgh.parallel_for(sycl::range<1>{1},
                         [=](sycl::id<1> idx_id) { func(a); });
    });
    q_ct1.wait_and_throw();
    return 0;
}

the error in MSVC is shown in https://github.com/ginkgo-project/ginkgo/runs/3008759379?check_suite_focus=true
error C2580 ... multiple versions of a defaulted special member functions are not allowed

[pratikvn]

Some more minor comments

[fritzgoebel]

The ELL kernels look good to me. On the rest I think others are better qualified to review

[fritzgoebel]

Suggested change

	* The number of threads within a CUDA warp.
	* The number of threads within a DPCPP warp.

[fritzgoebel]

Suggested change

	* The maximal number of threads allowed in a CUDA warp.
	* The maximal number of threads allowed in a DPCPP warp.

[thoasm]

I would prefer to move the check_if_valid helper structure.

[upsj]

LGTM! Does it make sense to keep the conversion scripts in-tree in the long run? We're not maintaining them, and they will not be used after we finished porting.

[upsj]

Maybe instead note that c needs to be scaled beforehand, otherwise this might lead to confusion. Also:

Suggested change

	// Because the atomic operation changes the values of c during computation,
	// it can not do the right alpha * a * b + beta * c operation.
	// Thus, the dpcpp kernel only computes alpha * a * b when it uses atomic
	// operation.
	if (atomic) {
	// Because the atomic operation changes the values of c during computation,
	// it can not do the right alpha * a * b + beta * c operation.
	// Thus, the dpcpp kernel only adds alpha * a * b when it uses atomic
	// operation.
	if (atomic) {

[yhmtsai]

I will move this part into the atomic scope and the following is what I use (including c be scaled)

// Because the atomic operation changes the values of c during
// computation, it can not directly do alpha * a * b + beta * c
// operation. The beta * c needs to be done before calling this kernel.
// Then, this kernel only adds alpha * a * b when it uses atomic
// operation.

[pratikvn]

LGTM! Thanks for adding the documentation.

Changes were addressed

[sonarcloud]

Kudos, SonarCloud Quality Gate passed!   

0 Bugs
0 Vulnerabilities
0 Security Hotspots
1 Code Smell

100.0% Coverage
5.8% Duplication

[yhmtsai]

@upsj I think we put it here like cuda2hip. it will be hard to use on adaptive changes.
implement cuda -> apply cuda2dpcpp -> delete those porting on unchanged function -> fix the porting
maybe some script to only extract the changed function or selection help this situtation