Skip to content

Commit

Permalink
Merge pull request #425 from beomki-yeo/update-vecmem
Browse files Browse the repository at this point in the history
Use `vecmem::sycl::local_accessor`
  • Loading branch information
beomki-yeo committed Jun 22, 2023
2 parents 48ad5a3 + ba0c6cc commit 19e1f8f
Show file tree
Hide file tree
Showing 2 changed files with 8 additions and 4 deletions.
5 changes: 3 additions & 2 deletions device/sycl/src/clusterization/clusterization_algorithm.sycl
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,7 @@
// Vecmem include(s).
#include <vecmem/memory/device_atomic_ref.hpp>
#include <vecmem/utils/sycl/copy.hpp>
#include <vecmem/utils/sycl/local_accessor.hpp>

// System include(s).
#include <algorithm>
Expand Down Expand Up @@ -117,8 +118,8 @@ clusterization_algorithm::output_type clusterization_algorithm::operator()(
// Run ccl kernel
details::get_queue(m_queue)
.submit([&](::sycl::handler& h) {
::sycl::local_accessor<unsigned int> shared_uint(3, h);
::sycl::local_accessor<index_t> shared_idx(
vecmem::sycl::local_accessor<unsigned int> shared_uint(3, h);
vecmem::sycl::local_accessor<index_t> shared_idx(
2 * max_cells_per_partition, h);

h.parallel_for<kernels::ccl_kernel>(
Expand Down
7 changes: 5 additions & 2 deletions device/sycl/src/seeding/seed_finding.sycl
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,9 @@
#include "traccc/seeding/device/select_seeds.hpp"
#include "traccc/seeding/device/update_triplet_weights.hpp"

// VecMem include(s).
#include <vecmem/utils/sycl/local_accessor.hpp>

namespace traccc::sycl {
namespace kernels {

Expand Down Expand Up @@ -294,7 +297,7 @@ seed_finding::output_type seed_finding::operator()(
details::get_queue(m_queue).submit([&](::sycl::handler& h) {
// Array for temporary storage of triplet weights for comparing
// within kernel
::sycl::local_accessor<scalar> local_mem(
vecmem::sycl::local_accessor<scalar> local_mem(
m_seedfilter_config.compatSeedLimit * weightUpdatingLocalSize,
h);

Expand Down Expand Up @@ -341,7 +344,7 @@ seed_finding::output_type seed_finding::operator()(
.submit([&](::sycl::handler& h) {
// Array for temporary storage of triplets for comparing within
// kernel
::sycl::local_accessor<triplet> local_mem(
vecmem::sycl::local_accessor<triplet> local_mem(
m_seedfilter_config.max_triplets_per_spM *
seedSelectingLocalSize,
h);
Expand Down

0 comments on commit 19e1f8f

Please sign in to comment.