Skip to content

Commit 59e520c

Browse files
IgorOchockiIgorOchocki
authored
Fix compilation error (#1371)
1 parent 149dbb6 commit 59e520c

File tree

5 files changed

+39
-39
lines changed

5 files changed

+39
-39
lines changed

‎DirectProgramming/C++SYCL/StructuredGrids/guided_HSOpticalflow_SYCLMigration/02_sycl_dpct_migrated/src/derivativesKernel.hpp

+13-13
Original file line numberDiff line numberDiff line change
@@ -52,13 +52,13 @@ using namespace sycl;
5252

5353
void ComputeDerivativesKernel(
5454
int width, int height, int stride, float *Ix, float *Iy, float *Iz,
55-
accessor<cl::sycl::float4, 2, cl::sycl::access::mode::read,
55+
accessor<sycl::float4, 2, sycl::access::mode::read,
5656
sycl::access::target::image>
5757
texSource,
58-
accessor<cl::sycl::float4, 2, cl::sycl::access::mode::read,
58+
accessor<sycl::float4, 2, sycl::access::mode::read,
5959
sycl::access::target::image>
6060
texTarget,
61-
cl::sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
61+
sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
6262
const int ix = item_ct1.get_local_id(2) +
6363
item_ct1.get_group(2) * item_ct1.get_local_range().get(2);
6464
const int iy = item_ct1.get_local_id(1) +
@@ -132,7 +132,7 @@ static void ComputeDerivatives(const float *I0, const float *I1, int w, int h,
132132
queue q) {
133133
sycl::range<3> threads(1, 6, 32);
134134
auto max_wg_size =
135-
q.get_device().get_info<cl::sycl::info::device::max_work_group_size>();
135+
q.get_device().get_info<sycl::info::device::max_work_group_size>();
136136
if (max_wg_size < 6 * 32) {
137137
threads[0] = 1;
138138
threads[2] = 32;
@@ -165,28 +165,28 @@ static void ComputeDerivatives(const float *I0, const float *I1, int w, int h,
165165
}
166166
}
167167

168-
auto texDescr = cl::sycl::sampler(
168+
auto texDescr = sycl::sampler(
169169
sycl::coordinate_normalization_mode::unnormalized,
170170
sycl::addressing_mode::clamp_to_edge, sycl::filtering_mode::nearest);
171171

172172
auto texSource =
173-
cl::sycl::image<2>(I0_p, cl::sycl::image_channel_order::rgba,
174-
cl::sycl::image_channel_type::fp32, range<2>(w, h),
173+
sycl::image<2>(I0_p, sycl::image_channel_order::rgba,
174+
sycl::image_channel_type::fp32, range<2>(w, h),
175175
range<1>(s * sizeof(sycl::float4)));
176176

177177
auto texTarget =
178-
cl::sycl::image<2>(I1_p, cl::sycl::image_channel_order::rgba,
179-
cl::sycl::image_channel_type::fp32, range<2>(w, h),
178+
sycl::image<2>(I1_p, sycl::image_channel_order::rgba,
179+
sycl::image_channel_type::fp32, range<2>(w, h),
180180
range<1>(s * sizeof(sycl::float4)));
181181

182182
dpct::get_default_queue()
183183
.submit([&](sycl::handler &cgh) {
184184
auto texSource_acc =
185-
texSource.template get_access<cl::sycl::float4,
186-
cl::sycl::access::mode::read>(cgh);
185+
texSource.template get_access<sycl::float4,
186+
sycl::access::mode::read>(cgh);
187187
auto texTarget_acc =
188-
texTarget.template get_access<cl::sycl::float4,
189-
cl::sycl::access::mode::read>(cgh);
188+
texTarget.template get_access<sycl::float4,
189+
sycl::access::mode::read>(cgh);
190190

191191
cgh.parallel_for(sycl::nd_range<3>(blocks * threads, threads),
192192
[=](sycl::nd_item<3> item_ct1) {

‎DirectProgramming/C++SYCL/StructuredGrids/guided_HSOpticalflow_SYCLMigration/02_sycl_dpct_migrated/src/downscaleKernel.hpp

+8-8
Original file line numberDiff line numberDiff line change
@@ -49,10 +49,10 @@ using namespace sycl;
4949
/// \param[out] out result
5050
///////////////////////////////////////////////////////////////////////////////
5151
void DownscaleKernel(int width, int height, int stride, float *out,
52-
accessor<cl::sycl::float4, 2, cl::sycl::access::mode::read,
52+
accessor<sycl::float4, 2, sycl::access::mode::read,
5353
sycl::access::target::image>
5454
tex_acc,
55-
cl::sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
55+
sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
5656
const int ix = item_ct1.get_local_id(2) +
5757
item_ct1.get_group(2) * item_ct1.get_local_range(2);
5858
const int iy = item_ct1.get_local_id(1) +
@@ -89,7 +89,7 @@ static void Downscale(const float *src, int width, int height, int stride,
8989
queue q) {
9090
sycl::range<3> threads(1, 8, 32);
9191
auto max_wg_size =
92-
q.get_device().get_info<cl::sycl::info::device::max_work_group_size>();
92+
q.get_device().get_info<sycl::info::device::max_work_group_size>();
9393
if (max_wg_size < 8 * 32) {
9494
threads[0] = 1;
9595
threads[2] = 32;
@@ -114,20 +114,20 @@ static void Downscale(const float *src, int width, int height, int stride,
114114
}
115115
}
116116

117-
auto texDescr = cl::sycl::sampler(
117+
auto texDescr = sycl::sampler(
118118
sycl::coordinate_normalization_mode::unnormalized,
119119
sycl::addressing_mode::clamp_to_edge, sycl::filtering_mode::nearest);
120120

121-
auto texFine = cl::sycl::image<2>(src_p, cl::sycl::image_channel_order::rgba,
122-
cl::sycl::image_channel_type::fp32,
121+
auto texFine = sycl::image<2>(src_p, sycl::image_channel_order::rgba,
122+
sycl::image_channel_type::fp32,
123123
range<2>(width, height),
124124
range<1>(stride * sizeof(sycl::float4)));
125125

126126
dpct::get_default_queue()
127127
.submit([&](sycl::handler &cgh) {
128128
auto tex_acc =
129-
texFine.template get_access<cl::sycl::float4,
130-
cl::sycl::access::mode::read>(cgh);
129+
texFine.template get_access<sycl::float4,
130+
sycl::access::mode::read>(cgh);
131131

132132
cgh.parallel_for(sycl::nd_range<3>(blocks * threads, threads),
133133
[=](sycl::nd_item<3> item_ct1) {

‎DirectProgramming/C++SYCL/StructuredGrids/guided_HSOpticalflow_SYCLMigration/02_sycl_dpct_migrated/src/solverKernel.hpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -184,7 +184,7 @@ static void SolveForUpdate(const float *du0, const float *dv0, const float *Ix,
184184
// CTA size
185185
sycl::range<3> threads(1, 6, 32);
186186
auto max_wg_size =
187-
q.get_device().get_info<cl::sycl::info::device::max_work_group_size>();
187+
q.get_device().get_info<sycl::info::device::max_work_group_size>();
188188
if (max_wg_size < 6 * 32) {
189189
threads[0] = 1;
190190
threads[2] = 32;

‎DirectProgramming/C++SYCL/StructuredGrids/guided_HSOpticalflow_SYCLMigration/02_sycl_dpct_migrated/src/upscaleKernel.hpp

+9-9
Original file line numberDiff line numberDiff line change
@@ -47,10 +47,10 @@ using namespace sycl;
4747
/// \param[out] out result
4848
///////////////////////////////////////////////////////////////////////////////
4949
void UpscaleKernel(int width, int height, int stride, float scale, float *out,
50-
accessor<cl::sycl::float4, 2, cl::sycl::access::mode::read,
50+
accessor<sycl::float4, 2, sycl::access::mode::read,
5151
sycl::access::target::image>
5252
texCoarse_acc,
53-
cl::sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
53+
sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
5454
const int ix = item_ct1.get_local_id(2) +
5555
item_ct1.get_group(2) * item_ct1.get_local_range().get(2);
5656
const int iy = item_ct1.get_local_id(1) +
@@ -85,7 +85,7 @@ static void Upscale(const float *src, int width, int height, int stride,
8585
float *out, queue q) {
8686
sycl::range<3> threads(1, 8, 32);
8787
auto max_wg_size =
88-
q.get_device().get_info<cl::sycl::info::device::max_work_group_size>();
88+
q.get_device().get_info<sycl::info::device::max_work_group_size>();
8989
if (max_wg_size < 8 * 32) {
9090
threads[0] = 1;
9191
threads[2] = 32;
@@ -107,20 +107,20 @@ static void Upscale(const float *src, int width, int height, int stride,
107107
src_p[index * 4 + 1] = src_p[index * 4 + 2] = src_p[index * 4 + 3] = 0.f;
108108
}
109109
}
110-
auto texDescr = cl::sycl::sampler(
110+
auto texDescr = sycl::sampler(
111111
sycl::coordinate_normalization_mode::unnormalized,
112112
sycl::addressing_mode::clamp_to_edge, sycl::filtering_mode::linear);
113113

114-
auto texCoarse = cl::sycl::image<2>(
115-
src_p, cl::sycl::image_channel_order::rgba,
116-
cl::sycl::image_channel_type::fp32, range<2>(width, height),
114+
auto texCoarse = sycl::image<2>(
115+
src_p, sycl::image_channel_order::rgba,
116+
sycl::image_channel_type::fp32, range<2>(width, height),
117117
range<1>(stride * sizeof(sycl::float4)));
118118

119119
dpct::get_default_queue()
120120
.submit([&](sycl::handler &cgh) {
121121
auto texCoarse_acc =
122-
texCoarse.template get_access<cl::sycl::float4,
123-
cl::sycl::access::mode::read>(cgh);
122+
texCoarse.template get_access<sycl::float4,
123+
sycl::access::mode::read>(cgh);
124124

125125
cgh.parallel_for(sycl::nd_range<3>(blocks * threads, threads),
126126
[=](sycl::nd_item<3> item_ct1) {

‎DirectProgramming/C++SYCL/StructuredGrids/guided_HSOpticalflow_SYCLMigration/02_sycl_dpct_migrated/src/warpingKernel.hpp

+8-8
Original file line numberDiff line numberDiff line change
@@ -49,10 +49,10 @@ using namespace sycl;
4949
///////////////////////////////////////////////////////////////////////////////
5050
void WarpingKernel(int width, int height, int stride, const float *u,
5151
const float *v, float *out,
52-
accessor<cl::sycl::float4, 2, cl::sycl::access::mode::read,
52+
accessor<sycl::float4, 2, sycl::access::mode::read,
5353
sycl::access::target::image>
5454
texToWarp,
55-
cl::sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
55+
sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
5656
const int ix = item_ct1.get_local_id(2) +
5757
item_ct1.get_group(2) * item_ct1.get_local_range().get(2);
5858
const int iy = item_ct1.get_local_id(1) +
@@ -90,7 +90,7 @@ static void WarpImage(const float *src, int w, int h, int s, const float *u,
9090
const float *v, float *out, queue q) {
9191
sycl::range<3> threads(1, 6, 32);
9292
auto max_wg_size =
93-
q.get_device().get_info<cl::sycl::info::device::max_work_group_size>();
93+
q.get_device().get_info<sycl::info::device::max_work_group_size>();
9494
if (max_wg_size < 6 * 32) {
9595
threads[0] = 1;
9696
threads[2] = 32;
@@ -110,20 +110,20 @@ static void WarpImage(const float *src, int w, int h, int s, const float *u,
110110
src_p[index * 4 + 1] = src_p[index * 4 + 2] = src_p[index * 4 + 3] = 0.f;
111111
}
112112
}
113-
auto texDescr = cl::sycl::sampler(
113+
auto texDescr = sycl::sampler(
114114
sycl::coordinate_normalization_mode::unnormalized,
115115
sycl::addressing_mode::clamp_to_edge, sycl::filtering_mode::linear);
116116

117117
auto texToWarp =
118-
cl::sycl::image<2>(src_p, cl::sycl::image_channel_order::rgba,
119-
cl::sycl::image_channel_type::fp32, range<2>(w, h),
118+
sycl::image<2>(src_p, sycl::image_channel_order::rgba,
119+
sycl::image_channel_type::fp32, range<2>(w, h),
120120
range<1>(s * sizeof(sycl::float4)));
121121

122122
dpct::get_default_queue()
123123
.submit([&](sycl::handler &cgh) {
124124
auto texToWarp_acc =
125-
texToWarp.template get_access<cl::sycl::float4,
126-
cl::sycl::access::mode::read>(cgh);
125+
texToWarp.template get_access<sycl::float4,
126+
sycl::access::mode::read>(cgh);
127127

128128
cgh.parallel_for(sycl::nd_range<3>(blocks * threads, threads),
129129
[=](sycl::nd_item<3> item_ct1) {

0 commit comments

Comments
 (0)