Skip to content

Commit 19e1f8f

Browse files
authored
Merge pull request #425 from beomki-yeo/update-vecmem
Use `vecmem::sycl::local_accessor`
2 parents 48ad5a3 + ba0c6cc commit 19e1f8f

File tree

2 files changed

+8
-4
lines changed

2 files changed

+8
-4
lines changed

device/sycl/src/clusterization/clusterization_algorithm.sycl

+3-2
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,7 @@
2020
// Vecmem include(s).
2121
#include <vecmem/memory/device_atomic_ref.hpp>
2222
#include <vecmem/utils/sycl/copy.hpp>
23+
#include <vecmem/utils/sycl/local_accessor.hpp>
2324

2425
// System include(s).
2526
#include <algorithm>
@@ -117,8 +118,8 @@ clusterization_algorithm::output_type clusterization_algorithm::operator()(
117118
// Run ccl kernel
118119
details::get_queue(m_queue)
119120
.submit([&](::sycl::handler& h) {
120-
::sycl::local_accessor<unsigned int> shared_uint(3, h);
121-
::sycl::local_accessor<index_t> shared_idx(
121+
vecmem::sycl::local_accessor<unsigned int> shared_uint(3, h);
122+
vecmem::sycl::local_accessor<index_t> shared_idx(
122123
2 * max_cells_per_partition, h);
123124

124125
h.parallel_for<kernels::ccl_kernel>(

device/sycl/src/seeding/seed_finding.sycl

+5-2
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,9 @@
3131
#include "traccc/seeding/device/select_seeds.hpp"
3232
#include "traccc/seeding/device/update_triplet_weights.hpp"
3333

34+
// VecMem include(s).
35+
#include <vecmem/utils/sycl/local_accessor.hpp>
36+
3437
namespace traccc::sycl {
3538
namespace kernels {
3639

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

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

0 commit comments

Comments
 (0)