Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Convert PreconditionChebyshev::set_initial_guess_kernel to Kokkos #15091

Merged
merged 1 commit into from
Apr 15, 2023
Merged

Convert PreconditionChebyshev::set_initial_guess_kernel to Kokkos #15091

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
58 changes: 12 additions & 46 deletions include/deal.II/lac/precondition.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3354,28 +3354,6 @@ namespace internal
vector.add(-mean_value);
}

template <typename Number>
void
set_initial_guess(
::dealii::LinearAlgebra::distributed::Vector<Number, MemorySpace::Host>
&vector)
{
// Choose a high-frequency mode consisting of numbers between 0 and 1
// that is cheap to compute (cheaper than random numbers) but avoids
// obviously re-occurring numbers in multi-component systems by choosing
// a period of 11.
// Make initial guess robust with respect to number of processors
// by operating on the global index.
types::global_dof_index first_local_range = 0;
if (!vector.locally_owned_elements().is_empty())
first_local_range = vector.locally_owned_elements().nth_index_in_set(0);
for (unsigned int i = 0; i < vector.locally_owned_size(); ++i)
vector.local_element(i) = (i + first_local_range) % 11;

const Number mean_value = vector.mean_value();
vector.add(-mean_value);
}

template <typename Number>
void
set_initial_guess(
Expand All @@ -3385,25 +3363,10 @@ namespace internal
set_initial_guess(vector.block(block));
}


# ifdef DEAL_II_WITH_CUDA
template <typename Number>
__global__ void
set_initial_guess_kernel(const types::global_dof_index offset,
const unsigned int locally_owned_size,
Number * values)

{
const unsigned int index = threadIdx.x + blockDim.x * blockIdx.x;
if (index < locally_owned_size)
values[index] = (index + offset) % 11;
}

template <typename Number>
template <typename Number, typename MemorySpace>
void
set_initial_guess(
::dealii::LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>
&vector)
::dealii::LinearAlgebra::distributed::Vector<Number, MemorySpace> &vector)
{
// Choose a high-frequency mode consisting of numbers between 0 and 1
// that is cheap to compute (cheaper than random numbers) but avoids
Expand All @@ -3416,16 +3379,19 @@ namespace internal
first_local_range = vector.locally_owned_elements().nth_index_in_set(0);

const auto n_local_elements = vector.locally_owned_size();
const int n_blocks =
1 + (n_local_elements - 1) / CUDAWrappers::block_size;
set_initial_guess_kernel<<<n_blocks, CUDAWrappers::block_size>>>(
first_local_range, n_local_elements, vector.get_values());
AssertCudaKernel();

Number * values_ptr = vector.get_values();
Kokkos::RangePolicy<typename MemorySpace::kokkos_space::execution_space,
Kokkos::IndexType<types::global_dof_index>>
policy(0, n_local_elements);
Kokkos::parallel_for(
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I hope this won't be the new standard how to write code in deal.II...

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Would you mind elaborating?

"PreconditionChebyshev::set_initial_guess",
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It really feel strange to specify a string for a function call.

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Well, this helps Kokkos' tools for profiling.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

One wished that they had made that the last argument to the function, with a default to "". But having it is a good idea.

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

You can omit the label anyway (but that's not recommended since it makes identifying kernels in Kokkos tools, Nsight Systems, VTune, etc. much harder).

policy,
KOKKOS_LAMBDA(types::global_dof_index i) {
values_ptr[i] = (i + first_local_range) % 11;
});
const Number mean_value = vector.mean_value();
vector.add(-mean_value);
}
# endif // DEAL_II_WITH_CUDA

struct EigenvalueTracker
{
Expand Down