Skip to content

Commit

Permalink
run clang-format on all files
Browse files Browse the repository at this point in the history
  • Loading branch information
@MarcelKoch
MarcelKoch committed Nov 14, 2023
1 parent 033f661 commit e9131ea
Show file tree
Hide file tree
Showing 6 changed files with 112 additions and 104 deletions.
4 changes: 2 additions & 2 deletions core/test/base/deferred_factory.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -86,8 +86,8 @@ struct test_impl<gko::xstd::void_t<decltype(T(std::declval<Args>()...))>, T,

// specialization for DF2 with_factory_list
template <typename... Args>
struct test_impl<gko::xstd::void_t<decltype(DF2::param{}.with_factory_list(
std::declval<Args>()...))>,
struct test_impl<gko::xstd::void_t<decltype(
DF2::param{}.with_factory_list(std::declval<Args>()...))>,
DummyFlag, Args...> : std::true_type {};

// test the object can be constructable or not with Args.
Expand Down
59 changes: 32 additions & 27 deletions dpcpp/factorization/par_ilut_filter_kernels.hpp.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -115,12 +115,13 @@ void threshold_filter_nnz(dim3 grid, dim3 block,
remove_complex<ValueType> threshold, IndexType* nnz,
bool lower)
{
queue->parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(subgroup_size)]] {
threshold_filter_nnz<subgroup_size>(
row_ptrs, vals, num_rows, threshold, nnz, lower, item_ct1);
});
queue->parallel_for(sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1)
[[sycl::reqd_sub_group_size(subgroup_size)]] {
threshold_filter_nnz<subgroup_size>(
row_ptrs, vals, num_rows, threshold, nnz,
lower, item_ct1);
});
}


Expand Down Expand Up @@ -152,14 +153,15 @@ void threshold_filter(dim3 grid, dim3 block, size_type dynamic_shared_memory,
const IndexType* new_row_ptrs, IndexType* new_row_idxs,
IndexType* new_col_idxs, ValueType* new_vals, bool lower)
{
queue->parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(subgroup_size)]] {
threshold_filter<subgroup_size>(
old_row_ptrs, old_col_idxs, old_vals, num_rows, threshold,
new_row_ptrs, new_row_idxs, new_col_idxs, new_vals, lower,
item_ct1);
});
queue->parallel_for(sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1)
[[sycl::reqd_sub_group_size(subgroup_size)]] {
threshold_filter<subgroup_size>(
old_row_ptrs, old_col_idxs, old_vals,
num_rows, threshold, new_row_ptrs,
new_row_idxs, new_col_idxs, new_vals, lower,
item_ct1);
});
}


Expand All @@ -183,12 +185,13 @@ void bucket_filter_nnz(dim3 grid, dim3 block, size_type dynamic_shared_memory,
const BucketType* buckets, IndexType num_rows,
BucketType bucket, IndexType* nnz)
{
queue->parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(subgroup_size)]] {
bucket_filter_nnz<subgroup_size>(row_ptrs, buckets, num_rows,
bucket, nnz, item_ct1);
});
queue->parallel_for(sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1)
[[sycl::reqd_sub_group_size(subgroup_size)]] {
bucket_filter_nnz<subgroup_size>(
row_ptrs, buckets, num_rows, bucket, nnz,
item_ct1);
});
}


Expand Down Expand Up @@ -222,13 +225,15 @@ void bucket_filter(dim3 grid, dim3 block, size_type dynamic_shared_memory,
IndexType* new_row_idxs, IndexType* new_col_idxs,
ValueType* new_vals)
{
queue->parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(subgroup_size)]] {
bucket_filter<subgroup_size>(
old_row_ptrs, old_col_idxs, old_vals, buckets, num_rows, bucket,
new_row_ptrs, new_row_idxs, new_col_idxs, new_vals, item_ct1);
});
queue->parallel_for(sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1)
[[sycl::reqd_sub_group_size(subgroup_size)]] {
bucket_filter<subgroup_size>(
old_row_ptrs, old_col_idxs, old_vals,
buckets, num_rows, bucket, new_row_ptrs,
new_row_idxs, new_col_idxs, new_vals,
item_ct1);
});
}


Expand Down
50 changes: 25 additions & 25 deletions dpcpp/factorization/par_ilut_select_kernels.hpp.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -68,13 +68,13 @@ void build_searchtree(dim3 grid, dim3 block, size_type dynamic_shared_memory,
sycl::access::target::local>
sh_samples_acc_ct1(sycl::range<1>(1024 /*sample_size*/), cgh);

cgh.parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(
config::warp_size)]] {
build_searchtree(input, size, tree_output, item_ct1,
sh_samples_acc_ct1.get_pointer());
});
cgh.parallel_for(sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1)
[[sycl::reqd_sub_group_size(config::warp_size)]] {
build_searchtree(
input, size, tree_output, item_ct1,
sh_samples_acc_ct1.get_pointer());
});
});
}

Expand Down Expand Up @@ -256,12 +256,13 @@ void block_prefix_sum(dim3 grid, dim3 block, size_type dynamic_shared_memory,
cgh);

cgh.parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(
config::warp_size)]] {
block_prefix_sum(counters, totals, num_blocks, item_ct1,
(IndexType*)warp_sums_acc_ct1.get_pointer());
});
sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1)
[[sycl::reqd_sub_group_size(config::warp_size)]] {
block_prefix_sum(
counters, totals, num_blocks, item_ct1,
(IndexType*)warp_sums_acc_ct1.get_pointer());
});
});
}

Expand Down Expand Up @@ -363,12 +364,12 @@ void basecase_select(dim3 grid, dim3 block, size_type dynamic_shared_memory,
sh_local_acc_ct1(sycl::range<1>(1024 /*basecase_size*/), cgh);

cgh.parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(
config::warp_size)]] {
basecase_select(input, size, rank, out, item_ct1,
(ValueType*)sh_local_acc_ct1.get_pointer());
});
sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1)
[[sycl::reqd_sub_group_size(config::warp_size)]] {
basecase_select(input, size, rank, out, item_ct1,
(ValueType*)sh_local_acc_ct1.get_pointer());
});
});
}

Expand Down Expand Up @@ -403,12 +404,11 @@ template <typename IndexType>
void find_bucket(dim3 grid, dim3 block, size_type dynamic_shared_memory,
sycl::queue* queue, IndexType* prefix_sum, IndexType rank)
{
queue->parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(
config::warp_size)]] {
find_bucket(prefix_sum, rank, item_ct1);
});
queue->parallel_for(sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1)
[[sycl::reqd_sub_group_size(config::warp_size)]] {
find_bucket(prefix_sum, rank, item_ct1);
});
}


Expand Down
99 changes: 50 additions & 49 deletions dpcpp/solver/common_gmres_kernels.dp.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,13 +6,14 @@
// Must be called with at least `max(stride_b * num_rows, krylov_dim *
// num_cols)` threads in total.
template <size_type block_size, typename ValueType>
void initialize_kernel(
size_type num_rows, size_type num_cols, size_type krylov_dim,
const ValueType *__restrict__ b, size_type stride_b,
ValueType *__restrict__ residual, size_type stride_residual,
ValueType *__restrict__ givens_sin, size_type stride_sin,
ValueType *__restrict__ givens_cos, size_type stride_cos,
stopping_status *__restrict__ stop_status, sycl::nd_item<3> item_ct1)
void initialize_kernel(size_type num_rows, size_type num_cols,
size_type krylov_dim, const ValueType* __restrict__ b,
size_type stride_b, ValueType* __restrict__ residual,
size_type stride_residual,
ValueType* __restrict__ givens_sin, size_type stride_sin,
ValueType* __restrict__ givens_cos, size_type stride_cos,
stopping_status* __restrict__ stop_status,
sycl::nd_item<3> item_ct1)
{
const auto global_id = thread::get_thread_id_flat(item_ct1);

Expand All @@ -39,15 +40,15 @@ void initialize_kernel(

template <size_type block_size, typename ValueType>
void initialize_kernel(dim3 grid, dim3 block, size_type dynamic_shared_memory,
sycl::queue *queue, size_type num_rows,
size_type num_cols, size_type krylov_dim,
const ValueType *b, size_type stride_b,
ValueType *residual, size_type stride_residual,
ValueType *givens_sin, size_type stride_sin,
ValueType *givens_cos, size_type stride_cos,
stopping_status *stop_status)
sycl::queue* queue, size_type num_rows,
size_type num_cols, size_type krylov_dim,
const ValueType* b, size_type stride_b,
ValueType* residual, size_type stride_residual,
ValueType* givens_sin, size_type stride_sin,
ValueType* givens_cos, size_type stride_cos,
stopping_status* stop_status)
{
queue->submit([&](sycl::handler &cgh) {
queue->submit([&](sycl::handler& cgh) {
cgh.parallel_for(
sycl_nd_range(grid, block), [=](sycl::nd_item<3> item_ct1) {
initialize_kernel<block_size>(
Expand All @@ -61,12 +62,12 @@ void initialize_kernel(dim3 grid, dim3 block, size_type dynamic_shared_memory,

template <typename ValueType>
void calculate_sin_and_cos_kernel(size_type col_idx, size_type num_cols,
size_type iter, const ValueType &this_hess,
const ValueType &next_hess,
ValueType *givens_sin, size_type stride_sin,
ValueType *givens_cos, size_type stride_cos,
ValueType &register_sin,
ValueType &register_cos)
size_type iter, const ValueType& this_hess,
const ValueType& next_hess,
ValueType* givens_sin, size_type stride_sin,
ValueType* givens_cos, size_type stride_cos,
ValueType& register_sin,
ValueType& register_cos)
{
if (is_zero(this_hess)) {
register_cos = zero<ValueType>();
Expand All @@ -89,10 +90,10 @@ void calculate_sin_and_cos_kernel(size_type col_idx, size_type num_cols,
template <typename ValueType>
void calculate_residual_norm_kernel(size_type col_idx, size_type num_cols,
size_type iter,
const ValueType &register_sin,
const ValueType &register_cos,
remove_complex<ValueType> *residual_norm,
ValueType *residual_norm_collection,
const ValueType& register_sin,
const ValueType& register_cos,
remove_complex<ValueType>* residual_norm,
ValueType* residual_norm_collection,
size_type stride_residual_norm_collection)
{
const auto this_rnc =
Expand All @@ -112,13 +113,13 @@ void calculate_residual_norm_kernel(size_type col_idx, size_type num_cols,
template <size_type block_size, typename ValueType>
void givens_rotation_kernel(
size_type num_rows, size_type num_cols, size_type iter,
ValueType *__restrict__ hessenberg_iter, size_type stride_hessenberg,
ValueType *__restrict__ givens_sin, size_type stride_sin,
ValueType *__restrict__ givens_cos, size_type stride_cos,
remove_complex<ValueType> *__restrict__ residual_norm,
ValueType *__restrict__ residual_norm_collection,
ValueType* __restrict__ hessenberg_iter, size_type stride_hessenberg,
ValueType* __restrict__ givens_sin, size_type stride_sin,
ValueType* __restrict__ givens_cos, size_type stride_cos,
remove_complex<ValueType>* __restrict__ residual_norm,
ValueType* __restrict__ residual_norm_collection,
size_type stride_residual_norm_collection,
const stopping_status *__restrict__ stop_status, sycl::nd_item<3> item_ct1)
const stopping_status* __restrict__ stop_status, sycl::nd_item<3> item_ct1)
{
const auto col_idx = thread::get_thread_id_flat(item_ct1);

Expand Down Expand Up @@ -167,18 +168,18 @@ void givens_rotation_kernel(

template <size_type block_size, typename ValueType>
void givens_rotation_kernel(dim3 grid, dim3 block,
size_type dynamic_shared_memory, sycl::queue *queue,
size_type dynamic_shared_memory, sycl::queue* queue,
size_type num_rows, size_type num_cols,
size_type iter, ValueType *hessenberg_iter,
size_type stride_hessenberg, ValueType *givens_sin,
size_type stride_sin, ValueType *givens_cos,
size_type iter, ValueType* hessenberg_iter,
size_type stride_hessenberg, ValueType* givens_sin,
size_type stride_sin, ValueType* givens_cos,
size_type stride_cos,
remove_complex<ValueType> *residual_norm,
ValueType *residual_norm_collection,
remove_complex<ValueType>* residual_norm,
ValueType* residual_norm_collection,
size_type stride_residual_norm_collection,
const stopping_status *stop_status)
const stopping_status* stop_status)
{
queue->submit([&](sycl::handler &cgh) {
queue->submit([&](sycl::handler& cgh) {
cgh.parallel_for(
sycl_nd_range(grid, block), [=](sycl::nd_item<3> item_ct1) {
givens_rotation_kernel<block_size>(
Expand All @@ -195,11 +196,11 @@ void givens_rotation_kernel(dim3 grid, dim3 block,
template <size_type block_size, typename ValueType>
void solve_upper_triangular_kernel(
size_type num_cols, size_type num_rhs,
const ValueType *__restrict__ residual_norm_collection,
const ValueType* __restrict__ residual_norm_collection,
size_type stride_residual_norm_collection,
const ValueType *__restrict__ hessenberg, size_type stride_hessenberg,
ValueType *__restrict__ y, size_type stride_y,
const size_type *__restrict__ final_iter_nums, sycl::nd_item<3> item_ct1)
const ValueType* __restrict__ hessenberg, size_type stride_hessenberg,
ValueType* __restrict__ y, size_type stride_y,
const size_type* __restrict__ final_iter_nums, sycl::nd_item<3> item_ct1)
{
const auto col_idx = thread::get_thread_id_flat(item_ct1);

Expand All @@ -225,14 +226,14 @@ void solve_upper_triangular_kernel(

template <size_type block_size, typename ValueType>
void solve_upper_triangular_kernel(
dim3 grid, dim3 block, size_type dynamic_shared_memory, sycl::queue *queue,
dim3 grid, dim3 block, size_type dynamic_shared_memory, sycl::queue* queue,
size_type num_cols, size_type num_rhs,
const ValueType *residual_norm_collection,
size_type stride_residual_norm_collection, const ValueType *hessenberg,
size_type stride_hessenberg, ValueType *y, size_type stride_y,
const size_type *final_iter_nums)
const ValueType* residual_norm_collection,
size_type stride_residual_norm_collection, const ValueType* hessenberg,
size_type stride_hessenberg, ValueType* y, size_type stride_y,
const size_type* final_iter_nums)
{
queue->submit([&](sycl::handler &cgh) {
queue->submit([&](sycl::handler& cgh) {
cgh.parallel_for(
sycl_nd_range(grid, block), [=](sycl::nd_item<3> item_ct1) {
solve_upper_triangular_kernel<block_size>(
Expand Down
2 changes: 1 addition & 1 deletion examples/kokkos_assembly/kokkos_assembly.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -8,8 +8,8 @@


#include <omp.h>
#include <Kokkos_Core.hpp>
#include <ginkgo/ginkgo.hpp>
#include <Kokkos_Core.hpp>


// Creates a stencil matrix in CSR format for the given number of discretization
Expand Down
2 changes: 2 additions & 0 deletions include/ginkgo/core/log/papi.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,8 @@
#include <iostream>
#include <map>
#include <mutex>


#include <sde_lib.h>


Expand Down

0 comments on commit e9131ea

Please sign in to comment.