Skip to content

Commit 1443e99

Browse files
committed
[SYCL] Fix pointer issue
Signed-off-by: Hu, Peisen <[email protected]>
1 parent 55df793 commit 1443e99

File tree

1 file changed

+16
-14
lines changed

1 file changed

+16
-14
lines changed

sycl/test-e2e/GroupAlgorithm/root_group.cpp

Lines changed: 16 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -69,7 +69,7 @@ template <typename T> class TestKernel1 {
6969
volatile float X = 1.0f;
7070
volatile float Y = 1.0f;
7171
auto root = it.ext_oneapi_get_root_group();
72-
*m_data[root.get_local_id()] = root.get_local_id();
72+
(*m_data)[root.get_local_id()] = root.get_local_id();
7373
sycl::group_barrier(root);
7474
// Delay half of the workgroups with extra work to check that the barrier
7575
// synchronizes the whole device.
@@ -78,10 +78,10 @@ template <typename T> class TestKernel1 {
7878
Y += sycl::cos(Y);
7979
}
8080
root = sycl::ext::oneapi::experimental::this_work_item::get_root_group<1>();
81-
int sum = *m_data[root.get_local_id()] +
82-
*m_data[root.get_local_range() - root.get_local_id() - 1];
81+
int sum = (*m_data)[root.get_local_id()] +
82+
(*m_data)[root.get_local_range() - root.get_local_id() - 1];
8383
sycl::group_barrier(root);
84-
*m_data[root.get_local_id()] = sum;
84+
(*m_data)[root.get_local_id()] = sum;
8585
}
8686
auto get(sycl::ext::oneapi::experimental::properties_tag) {
8787
return sycl::ext::oneapi::experimental::properties{
@@ -121,17 +121,19 @@ template <typename T> class TestKernel2 {
121121
void operator()(sycl::nd_item<1> it) const {
122122
const auto root = it.ext_oneapi_get_root_group();
123123
if (root.leader() || root.get_local_id() == 3) {
124-
*m_testResults[0] = root.get_group_id() == sycl::id<1>(0);
125-
*m_testResults[1] = root.leader() ? root.get_local_id() == sycl::id<1>(0)
126-
: root.get_local_id() == sycl::id<1>(3);
127-
*m_testResults[2] = root.get_group_range() == sycl::range<1>(1);
128-
*m_testResults[3] = root.get_local_range() == it.get_global_range();
129-
*m_testResults[4] = root.get_max_local_range() == root.get_local_range();
130-
*m_testResults[5] = root.get_group_linear_id() == 0;
131-
*m_testResults[6] =
124+
(*m_testResults)[0] = root.get_group_id() == sycl::id<1>(0);
125+
(*m_testResults)[1] = root.leader()
126+
? root.get_local_id() == sycl::id<1>(0)
127+
: root.get_local_id() == sycl::id<1>(3);
128+
(*m_testResults)[2] = root.get_group_range() == sycl::range<1>(1);
129+
(*m_testResults)[3] = root.get_local_range() == it.get_global_range();
130+
(*m_testResults)[4] =
131+
root.get_max_local_range() == root.get_local_range();
132+
(*m_testResults)[5] = root.get_group_linear_id() == 0;
133+
(*m_testResults)[6] =
132134
root.get_local_linear_id() == root.get_local_id().get(0);
133-
*m_testResults[7] = root.get_group_linear_range() == 1;
134-
*m_testResults[8] =
135+
(*m_testResults)[7] = root.get_group_linear_range() == 1;
136+
(*m_testResults)[8] =
135137
root.get_local_linear_range() == root.get_local_range().size();
136138
}
137139
}

0 commit comments

Comments
 (0)