Skip to content

Commit

Permalink
changed sycl clusterization/seed finding kernel calls to be compatibl…
Browse files Browse the repository at this point in the history 
…e with deprecated accessors
  • Loading branch information
@guilhermeAlmeida1
guilhermeAlmeida1 committed Jan 17, 2023
1 parent 48cf4c9 commit f45c5be
Show file tree
Hide file tree
Showing 2 changed files with 57 additions and 36 deletions.
76 changes: 47 additions & 29 deletions device/sycl/src/clusterization/clusterization_algorithm.sycl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -128,14 +128,21 @@ void fast_sv_1(index_t* f, index_t* gf, unsigned char adjc, index_t adjv[8],

class ccl_kernel {
public:
ccl_kernel(const alt_cell_collection_types::const_view cells,
const cell_module_collection_types::const_view modules,
const ccl_partition_collection_types::const_view partitions,
alt_measurement_collection_types::view measurements,
unsigned int* num_measurements,
::sycl::local_accessor<index_t> father,
::sycl::local_accessor<index_t> grandfather,
::sycl::local_accessor<unsigned int> outCounter)
ccl_kernel(
const alt_cell_collection_types::const_view cells,
const cell_module_collection_types::const_view modules,
const ccl_partition_collection_types::const_view partitions,
alt_measurement_collection_types::view measurements,
unsigned int* num_measurements,
::sycl::accessor<index_t, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
father,
::sycl::accessor<index_t, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
grandfather,
::sycl::accessor<unsigned int, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
outCounter)
: cells_view(cells),
modules_view(modules),
partitions_view(partitions),
Expand Down Expand Up @@ -227,9 +234,10 @@ class ccl_kernel {
* themself assigned as a parent.
*/
if (f[tid] == tid) {
::sycl::atomic_ref<unsigned int, ::sycl::memory_order::relaxed,
::sycl::memory_scope::work_group,
::sycl::access::address_space::local_space>(outi)
::sycl::ext::oneapi::atomic_ref<
unsigned int, ::sycl::memory_order::relaxed,
::sycl::memory_scope::work_group,
::sycl::access::address_space::local_space>(outi)
.fetch_add(1);
}

Expand All @@ -244,12 +252,12 @@ class ccl_kernel {
* amount of threads per block, this has no sever implications.
*/
if (tid == 0) {
outi =
::sycl::atomic_ref<unsigned int, ::sycl::memory_order::relaxed,
::sycl::memory_scope::device,
::sycl::access::address_space::global_space>(
measurement_count)
.fetch_add(outi);
outi = ::sycl::ext::oneapi::atomic_ref<
unsigned int, ::sycl::memory_order::relaxed,
::sycl::memory_scope::device,
::sycl::access::address_space::global_space>(
measurement_count)
.fetch_add(outi);
}

item.barrier();
Expand All @@ -273,10 +281,10 @@ class ccl_kernel {
* output array which we can write to.
*/
const unsigned int id =
::sycl::atomic_ref<unsigned int, ::sycl::memory_order::relaxed,
::sycl::memory_scope::work_group,
::sycl::access::address_space::local_space>(
outi)
::sycl::ext::oneapi::atomic_ref<
unsigned int, ::sycl::memory_order::relaxed,
::sycl::memory_scope::work_group,
::sycl::access::address_space::local_space>(outi)
.fetch_add(1);

device::aggregate_cluster(cells_device, modules_device, &f[0],
Expand All @@ -291,9 +299,15 @@ class ccl_kernel {
const ccl_partition_collection_types::const_view partitions_view;
alt_measurement_collection_types::view measurements_view;
unsigned int* m_measurement_count;
::sycl::local_accessor<index_t> f;
::sycl::local_accessor<index_t> f_next;
::sycl::local_accessor<unsigned int> m_outi;
::sycl::accessor<index_t, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
f;
::sycl::accessor<index_t, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
f_next;
::sycl::accessor<unsigned int, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
m_outi;
};

} // namespace kernels
Expand Down Expand Up @@ -344,11 +358,15 @@ clusterization_algorithm::output_type clusterization_algorithm::operator()(
details::get_queue(m_queue)
.submit([&ndrange, &cells, &modules, &partitions, &measurements_view,
&num_measurements_device](::sycl::handler& h) {
::sycl::local_accessor<index_t, 1> f(
partitioning::MAX_CELLS_PER_PARTITION, h);
::sycl::local_accessor<index_t, 1> f_next(
partitioning::MAX_CELLS_PER_PARTITION, h);
::sycl::local_accessor<unsigned int, 1> outi(1, h);
::sycl::accessor<index_t, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
f(partitioning::MAX_CELLS_PER_PARTITION, h);
::sycl::accessor<index_t, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
f_next(partitioning::MAX_CELLS_PER_PARTITION, h);
::sycl::accessor<unsigned int, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
outi(1, h);

h.parallel_for<kernels::ccl_kernel>(
ndrange, kernels::ccl_kernel(
Expand Down
17 changes: 10 additions & 7 deletions device/sycl/src/seeding/seed_finding.sycl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -255,9 +255,11 @@ seed_finding::output_type seed_finding::operator()(
details::get_queue(m_queue).submit([&](::sycl::handler& h) {
// Array for temporary storage of triplets for comparing within seed
// selecting kernel
::sycl::local_accessor<scalar, 1> local_mem(
m_seedfilter_config.compatSeedLimit * weightUpdatingLocalSize,
h);
::sycl::accessor<scalar, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
local_mem(m_seedfilter_config.compatSeedLimit *
weightUpdatingLocalSize,
h);

h.parallel_for<kernels::update_triplet_weights>(
weightUpdatingRange,
Expand Down Expand Up @@ -310,10 +312,11 @@ seed_finding::output_type seed_finding::operator()(
.submit([&](::sycl::handler& h) {
// Array for temporary storage of triplets for comparing within seed
// selecting kernel
::sycl::local_accessor<triplet, 1> local_mem(
m_seedfilter_config.max_triplets_per_spM *
seedSelectingLocalSize,
h);
::sycl::accessor<triplet, 1, ::sycl::access::mode::read_write,
::sycl::access::target::local>
local_mem(m_seedfilter_config.max_triplets_per_spM *
seedSelectingLocalSize,
h);

h.parallel_for<kernels::select_seeds>(
seedSelectingRange,
Expand Down

0 comments on commit f45c5be

Please sign in to comment.