Skip to content

Commit

Permalink
Format files
Browse files Browse the repository at this point in the history
Co-authored-by: Pratik Nayak <[email protected]>
  • Loading branch information
ginkgo-bot and pratikvn committed Nov 6, 2023
1 parent 65d4ac9 commit 9f9cab4
Show file tree
Hide file tree
Showing 4 changed files with 35 additions and 33 deletions.
56 changes: 29 additions & 27 deletions dpcpp/matrix/batch_csr_kernels.dp.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -97,17 +97,18 @@ void simple_apply(std::shared_ptr<const DefaultExecutor> exec,
// Launch a kernel that has nbatches blocks, each block has max group size
exec->get_queue()->submit([&](sycl::handler& cgh) {
cgh.parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(
config::warp_size)]] {
auto group = item_ct1.get_group();
auto group_id = group.get_group_linear_id();
const auto mat_b =
batch::matrix::extract_batch_item(mat_ub, group_id);
const auto b_b = batch::extract_batch_item(b_ub, group_id);
const auto x_b = batch::extract_batch_item(x_ub, group_id);
simple_apply_kernel(mat_b, b_b.values, x_b.values, item_ct1);
});
sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1)
[[sycl::reqd_sub_group_size(config::warp_size)]] {
auto group = item_ct1.get_group();
auto group_id = group.get_group_linear_id();
const auto mat_b =
batch::matrix::extract_batch_item(mat_ub, group_id);
const auto b_b = batch::extract_batch_item(b_ub, group_id);
const auto x_b = batch::extract_batch_item(x_ub, group_id);
simple_apply_kernel(mat_b, b_b.values, x_b.values,
item_ct1);
});
});
}

Expand Down Expand Up @@ -145,22 +146,23 @@ void advanced_apply(std::shared_ptr<const DefaultExecutor> exec,
// Launch a kernel that has nbatches blocks, each block has max group size
exec->get_queue()->submit([&](sycl::handler& cgh) {
cgh.parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(
config::warp_size)]] {
auto group = item_ct1.get_group();
auto group_id = group.get_group_linear_id();
const auto mat_b =
batch::matrix::extract_batch_item(mat_ub, group_id);
const auto b_b = batch::extract_batch_item(b_ub, group_id);
const auto x_b = batch::extract_batch_item(x_ub, group_id);
const auto alpha_b =
batch::extract_batch_item(alpha_ub, group_id);
const auto beta_b =
batch::extract_batch_item(beta_ub, group_id);
advanced_apply_kernel(alpha_b.values[0], mat_b, b_b.values,
beta_b.values[0], x_b.values, item_ct1);
});
sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1)
[[sycl::reqd_sub_group_size(config::warp_size)]] {
auto group = item_ct1.get_group();
auto group_id = group.get_group_linear_id();
const auto mat_b =
batch::matrix::extract_batch_item(mat_ub, group_id);
const auto b_b = batch::extract_batch_item(b_ub, group_id);
const auto x_b = batch::extract_batch_item(x_ub, group_id);
const auto alpha_b =
batch::extract_batch_item(alpha_ub, group_id);
const auto beta_b =
batch::extract_batch_item(beta_ub, group_id);
advanced_apply_kernel(alpha_b.values[0], mat_b, b_b.values,
beta_b.values[0], x_b.values,
item_ct1);
});
});
}

Expand Down
7 changes: 3 additions & 4 deletions dpcpp/solver/batch_bicgstab_kernels.dp.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -119,10 +119,9 @@ class KernelCaller {
slm_values(sycl::range<1>(shared_size), cgh);

cgh.parallel_for(
sycl_nd_range(grid, block), [=
](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(
subgroup_size)]] [
[intel::kernel_args_restrict]] {
sycl_nd_range(grid, block),
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(
subgroup_size)]] [[intel::kernel_args_restrict]] {
auto batch_id = item_ct1.get_group_linear_id();
const auto mat_global_entry =
gko::batch::matrix::extract_batch_item(mat, batch_id);
Expand Down
4 changes: 2 additions & 2 deletions include/ginkgo/core/matrix/batch_csr.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -221,8 +221,8 @@ class Csr final
* significantly more memory efficient than the non-constant version,
* so always prefer this version.
*/
const value_type* get_const_values_for_item(size_type batch_id) const
noexcept
const value_type* get_const_values_for_item(
size_type batch_id) const noexcept
{
GKO_ASSERT(batch_id < this->get_num_batch_items());
GKO_ASSERT(values_.get_num_elems() >=
Expand Down
1 change: 1 addition & 0 deletions include/ginkgo/ginkgo.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -109,6 +109,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include <ginkgo/core/log/record.hpp>
#include <ginkgo/core/log/stream.hpp>

#include <ginkgo/core/matrix/batch_csr.hpp>
#include <ginkgo/core/matrix/batch_dense.hpp>
#include <ginkgo/core/matrix/batch_ell.hpp>
#include <ginkgo/core/matrix/batch_identity.hpp>
Expand Down

0 comments on commit 9f9cab4

Please sign in to comment.