Finish Kokkos::Parallel in ExtrapolatedSmootherTake

[julianlitz]

Merge Request - GuideLine Checklist

Guideline to check code before resolve WIP and approval, respectively.
As many checkboxes as possible should be ticked.

Checks by code author:

Always to be checked:

• There is at least one issue associated with the pull request.
• New code adheres with the coding guidelines
• No large data files have been added to the repository. Maximum size for files should be of the order of KB not MB. In particular avoid adding of pdf, word, or other files that cannot be change-tracked correctly by git.

If functions were changed or functionality was added:

• Tests for new functionality has been added
• A local test was succesful

If new functionality was added:

• There is appropriate documentation of your work. (use doxygen style comments)

If new third party software is used:

• Did you pay attention to its license? Please remember to add it to the wiki after successful merging.

If new mathematical methods or epidemiological terms are used:

• Are new methods referenced? Did you provide further documentation?

Checks by code reviewer(s):

• Is the code clean of development artifacts e.g., unnecessary comments, prints, ...
• The ticket goals for each associated issue are reached or problems are clearly addressed (i.e., a new issue was introduced).
• There are appropriate unit tests and they pass.
• The git history is clean and linearized for the merge request. All reviewers should squash commits and write a simple and meaningful commit message.
• Coverage report for new code is acceptable.
• No large data files have been added to the repository. Maximum size for files should be of the order of KB not MB. In particular avoid adding of pdf, word, or other files that cannot be change-tracked correctly by git.

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 89.48%. Comparing base (7acade7) to head (8bc8a33).
⚠️ Report is 1 commits behind head on main.

Additional details and impacted files

@@            Coverage Diff             @@
##             main     #268      +/-   ##
==========================================
- Coverage   89.49%   89.48%   -0.02%     
==========================================
  Files          79       79              
  Lines        9223     9214       -9     
==========================================
- Hits         8254     8245       -9     
  Misses        969      969              

☔ View full report in Codecov by Sentry.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

[julianlitz]

You sure we need to remove the const&? Here all test are green with const&.

Or would this only be needed for the class lambdas due to:

[EmilyBourne]

You sure we need to remove the const&? Here all test are green with const&.

Yes. You would only see failures on GPU. All tests in the CI are CPU only (plus we hare still running the loop on OpenMP for now)

It is needed for the KOKKOS_CLASS_LAMBDA but not for the reason that you describe. The KOKKOS_CLASS_LAMBDA does not use the operator= (which would require the modification of a const member), instead it just uses the constructor.

Compiler output for operator= vs copy constructor

Compare your suggestion:

#include <utility>

class MyClass {
    public:
        MyClass() : x(0) {}
    private:
        const int x;
};

int main () {
    MyClass a;
    MyClass b;
    b = a;
    b = std::move(a);
}

which generates the following compiler errors:

tmp.cpp: In function ‘int main()’:
tmp.cpp:13:9: error: use of deleted function ‘MyClass& MyClass::operator=(const MyClass&)’
   13 |     b = a;
      |         ^
tmp.cpp:3:7: note: ‘MyClass& MyClass::operator=(const MyClass&)’ is implicitly deleted because the default definition would be ill-formed:
    3 | class MyClass {
      |       ^~~~~~~
tmp.cpp:3:7: error: non-static const member ‘const int MyClass::x’, cannot use default assignment operator
tmp.cpp:14:20: error: use of deleted function ‘MyClass& MyClass::operator=(MyClass&&)’
   14 |     b = std::move(a);
      |                    ^
tmp.cpp:3:7: note: ‘MyClass& MyClass::operator=(MyClass&&)’ is implicitly deleted because the default definition would be ill-formed:
    3 | class MyClass {
      |       ^~~~~~~
tmp.cpp:3:7: error: non-static const member ‘const int MyClass::x’, cannot use default assignment operator

with a version using the copy constructor:

#include <utility>

class MyClass {
    public:
        MyClass() : x(0) {}
    private:
        const int x;
};

int main () {
    MyClass a;
    MyClass b(a);
    MyClass c(std::move(a));
}

which compiles with no errors

For GPU the issue is:

#include <iostream>
#include <Kokkos_Core.hpp>

class DataClass {
public:
    KOKKOS_FUNCTION DataClass() : x_(0) {}
    KOKKOS_DEFAULTED_FUNCTION DataClass(const DataClass&) = default;
    KOKKOS_INLINE_FUNCTION int x() const { return x_; }
private:
    const int x_; 
};

class MyClass {
public:
    MyClass(DataClass const& data_on_cpu)
     : reference_to_data_on_mem_space_where_data_was_created(data_on_cpu) {}
    KOKKOS_INLINE_FUNCTION int x() const {
       return reference_to_data_on_mem_space_where_data_was_created.x();
    }
private:
    DataClass const& reference_to_data_on_mem_space_where_data_was_created;
};

class MyClass2 {
public:
    MyClass2(DataClass const& data_on_cpu)
     : copy_of_data_on_mem_space_where_my_class2_was_created(data_on_cpu) {}
    KOKKOS_INLINE_FUNCTION int x() const {
       return copy_of_data_on_mem_space_where_my_class2_was_created.x();
    }
private:
    DataClass copy_of_data_on_mem_space_where_my_class2_was_created;
};

int main (int argc, char* argv[]) {
    Kokkos::ScopeGuard kokkos_scope(argc, argv);

    DataClass a;
    MyClass alloc_on_cpu(a);
    MyClass2 alloc_on_cpu2(a);
    std::cout << alloc_on_cpu.x() << std::endl; // OK because memory and execution is on CPU
    std::cout << alloc_on_cpu2.x() << std::endl; // OK because memory and execution is on CPU
    Kokkos::parallel_for("Test OpenMP loop", Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(0, 1),
            KOKKOS_LAMBDA(const int i) {
                // OK because memory and execution is on CPU
                Kokkos::printf("%d\n", alloc_on_cpu.x());
                // OK because memory and execution is on CPU
                Kokkos::printf("%d\n", alloc_on_cpu2.x());
            });
    Kokkos::parallel_for("Test GPU loop", Kokkos::RangePolicy<Kokkos::DefaultExecutionSpace>(0, 1),
            KOKKOS_LAMBDA(const int i) {
                // KO because a is referenced directly. It is on CPU but execution is on GPU
                Kokkos::printf("%d\n", alloc_on_cpu.x());
                // OK because memory and execution is on GPU (thanks to copy constructors)
                Kokkos::printf("%d\n", alloc_on_cpu2.x());
            });

}

This program compiles perfectly but it crashes on execution as follows:

-bash $> ./src/reproducer 
0
0
0
0
cudaDeviceSynchronize() error( cudaErrorIllegalAddress): an illegal memory access was encountered Kokkos_Cuda_Instance.cpp:154

Commenting the line marked KO leads to correct execution

[julianlitz]

Probably best if we dont use Kokkos class lambda anywhere ( expect the one in the CSR Solver - thats fine ).

They just break everything in the code so easily.