Skip to content

Fix compilation error #1371

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Feb 22, 2023
Merged

Fix compilation error #1371

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
26 changes: 13 additions & 13 deletions ...dGrids/guided_HSOpticalflow_SYCLMigration/02_sycl_dpct_migrated/src/derivativesKernel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,13 +52,13 @@ using namespace sycl;

void ComputeDerivativesKernel(
int width, int height, int stride, float *Ix, float *Iy, float *Iz,
accessor<cl::sycl::float4, 2, cl::sycl::access::mode::read,
accessor<sycl::float4, 2, sycl::access::mode::read,
sycl::access::target::image>
texSource,
accessor<cl::sycl::float4, 2, cl::sycl::access::mode::read,
accessor<sycl::float4, 2, sycl::access::mode::read,
sycl::access::target::image>
texTarget,
cl::sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
const int ix = item_ct1.get_local_id(2) +
item_ct1.get_group(2) * item_ct1.get_local_range().get(2);
const int iy = item_ct1.get_local_id(1) +
Expand Down Expand Up @@ -132,7 +132,7 @@ static void ComputeDerivatives(const float *I0, const float *I1, int w, int h,
queue q) {
sycl::range<3> threads(1, 6, 32);
auto max_wg_size =
q.get_device().get_info<cl::sycl::info::device::max_work_group_size>();
q.get_device().get_info<sycl::info::device::max_work_group_size>();
if (max_wg_size < 6 * 32) {
threads[0] = 1;
threads[2] = 32;
Expand Down Expand Up @@ -165,28 +165,28 @@ static void ComputeDerivatives(const float *I0, const float *I1, int w, int h,
}
}

auto texDescr = cl::sycl::sampler(
auto texDescr = sycl::sampler(
sycl::coordinate_normalization_mode::unnormalized,
sycl::addressing_mode::clamp_to_edge, sycl::filtering_mode::nearest);

auto texSource =
cl::sycl::image<2>(I0_p, cl::sycl::image_channel_order::rgba,
cl::sycl::image_channel_type::fp32, range<2>(w, h),
sycl::image<2>(I0_p, sycl::image_channel_order::rgba,
sycl::image_channel_type::fp32, range<2>(w, h),
range<1>(s * sizeof(sycl::float4)));

auto texTarget =
cl::sycl::image<2>(I1_p, cl::sycl::image_channel_order::rgba,
cl::sycl::image_channel_type::fp32, range<2>(w, h),
sycl::image<2>(I1_p, sycl::image_channel_order::rgba,
sycl::image_channel_type::fp32, range<2>(w, h),
range<1>(s * sizeof(sycl::float4)));

dpct::get_default_queue()
.submit([&](sycl::handler &cgh) {
auto texSource_acc =
texSource.template get_access<cl::sycl::float4,
cl::sycl::access::mode::read>(cgh);
texSource.template get_access<sycl::float4,
sycl::access::mode::read>(cgh);
auto texTarget_acc =
texTarget.template get_access<cl::sycl::float4,
cl::sycl::access::mode::read>(cgh);
texTarget.template get_access<sycl::float4,
sycl::access::mode::read>(cgh);

cgh.parallel_for(sycl::nd_range<3>(blocks * threads, threads),
[=](sycl::nd_item<3> item_ct1) {
Expand Down
16 changes: 8 additions & 8 deletions ...redGrids/guided_HSOpticalflow_SYCLMigration/02_sycl_dpct_migrated/src/downscaleKernel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,10 +49,10 @@ using namespace sycl;
/// \param[out] out result
///////////////////////////////////////////////////////////////////////////////
void DownscaleKernel(int width, int height, int stride, float *out,
accessor<cl::sycl::float4, 2, cl::sycl::access::mode::read,
accessor<sycl::float4, 2, sycl::access::mode::read,
sycl::access::target::image>
tex_acc,
cl::sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
const int ix = item_ct1.get_local_id(2) +
item_ct1.get_group(2) * item_ct1.get_local_range(2);
const int iy = item_ct1.get_local_id(1) +
Expand Down Expand Up @@ -89,7 +89,7 @@ static void Downscale(const float *src, int width, int height, int stride,
queue q) {
sycl::range<3> threads(1, 8, 32);
auto max_wg_size =
q.get_device().get_info<cl::sycl::info::device::max_work_group_size>();
q.get_device().get_info<sycl::info::device::max_work_group_size>();
if (max_wg_size < 8 * 32) {
threads[0] = 1;
threads[2] = 32;
Expand All @@ -114,20 +114,20 @@ static void Downscale(const float *src, int width, int height, int stride,
}
}

auto texDescr = cl::sycl::sampler(
auto texDescr = sycl::sampler(
sycl::coordinate_normalization_mode::unnormalized,
sycl::addressing_mode::clamp_to_edge, sycl::filtering_mode::nearest);

auto texFine = cl::sycl::image<2>(src_p, cl::sycl::image_channel_order::rgba,
cl::sycl::image_channel_type::fp32,
auto texFine = sycl::image<2>(src_p, sycl::image_channel_order::rgba,
sycl::image_channel_type::fp32,
range<2>(width, height),
range<1>(stride * sizeof(sycl::float4)));

dpct::get_default_queue()
.submit([&](sycl::handler &cgh) {
auto tex_acc =
texFine.template get_access<cl::sycl::float4,
cl::sycl::access::mode::read>(cgh);
texFine.template get_access<sycl::float4,
sycl::access::mode::read>(cgh);

cgh.parallel_for(sycl::nd_range<3>(blocks * threads, threads),
[=](sycl::nd_item<3> item_ct1) {
Expand Down
2 changes: 1 addition & 1 deletion ...cturedGrids/guided_HSOpticalflow_SYCLMigration/02_sycl_dpct_migrated/src/solverKernel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -184,7 +184,7 @@ static void SolveForUpdate(const float *du0, const float *dv0, const float *Ix,
// CTA size
sycl::range<3> threads(1, 6, 32);
auto max_wg_size =
q.get_device().get_info<cl::sycl::info::device::max_work_group_size>();
q.get_device().get_info<sycl::info::device::max_work_group_size>();
if (max_wg_size < 6 * 32) {
threads[0] = 1;
threads[2] = 32;
Expand Down
18 changes: 9 additions & 9 deletions ...turedGrids/guided_HSOpticalflow_SYCLMigration/02_sycl_dpct_migrated/src/upscaleKernel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,10 +47,10 @@ using namespace sycl;
/// \param[out] out result
///////////////////////////////////////////////////////////////////////////////
void UpscaleKernel(int width, int height, int stride, float scale, float *out,
accessor<cl::sycl::float4, 2, cl::sycl::access::mode::read,
accessor<sycl::float4, 2, sycl::access::mode::read,
sycl::access::target::image>
texCoarse_acc,
cl::sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
const int ix = item_ct1.get_local_id(2) +
item_ct1.get_group(2) * item_ct1.get_local_range().get(2);
const int iy = item_ct1.get_local_id(1) +
Expand Down Expand Up @@ -85,7 +85,7 @@ static void Upscale(const float *src, int width, int height, int stride,
float *out, queue q) {
sycl::range<3> threads(1, 8, 32);
auto max_wg_size =
q.get_device().get_info<cl::sycl::info::device::max_work_group_size>();
q.get_device().get_info<sycl::info::device::max_work_group_size>();
if (max_wg_size < 8 * 32) {
threads[0] = 1;
threads[2] = 32;
Expand All @@ -107,20 +107,20 @@ static void Upscale(const float *src, int width, int height, int stride,
src_p[index * 4 + 1] = src_p[index * 4 + 2] = src_p[index * 4 + 3] = 0.f;
}
}
auto texDescr = cl::sycl::sampler(
auto texDescr = sycl::sampler(
sycl::coordinate_normalization_mode::unnormalized,
sycl::addressing_mode::clamp_to_edge, sycl::filtering_mode::linear);

auto texCoarse = cl::sycl::image<2>(
src_p, cl::sycl::image_channel_order::rgba,
cl::sycl::image_channel_type::fp32, range<2>(width, height),
auto texCoarse = sycl::image<2>(
src_p, sycl::image_channel_order::rgba,
sycl::image_channel_type::fp32, range<2>(width, height),
range<1>(stride * sizeof(sycl::float4)));

dpct::get_default_queue()
.submit([&](sycl::handler &cgh) {
auto texCoarse_acc =
texCoarse.template get_access<cl::sycl::float4,
cl::sycl::access::mode::read>(cgh);
texCoarse.template get_access<sycl::float4,
sycl::access::mode::read>(cgh);

cgh.parallel_for(sycl::nd_range<3>(blocks * threads, threads),
[=](sycl::nd_item<3> item_ct1) {
Expand Down
16 changes: 8 additions & 8 deletions ...turedGrids/guided_HSOpticalflow_SYCLMigration/02_sycl_dpct_migrated/src/warpingKernel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,10 +49,10 @@ using namespace sycl;
///////////////////////////////////////////////////////////////////////////////
void WarpingKernel(int width, int height, int stride, const float *u,
const float *v, float *out,
accessor<cl::sycl::float4, 2, cl::sycl::access::mode::read,
accessor<sycl::float4, 2, sycl::access::mode::read,
sycl::access::target::image>
texToWarp,
cl::sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
sycl::sampler texDesc, sycl::nd_item<3> item_ct1) {
const int ix = item_ct1.get_local_id(2) +
item_ct1.get_group(2) * item_ct1.get_local_range().get(2);
const int iy = item_ct1.get_local_id(1) +
Expand Down Expand Up @@ -90,7 +90,7 @@ static void WarpImage(const float *src, int w, int h, int s, const float *u,
const float *v, float *out, queue q) {
sycl::range<3> threads(1, 6, 32);
auto max_wg_size =
q.get_device().get_info<cl::sycl::info::device::max_work_group_size>();
q.get_device().get_info<sycl::info::device::max_work_group_size>();
if (max_wg_size < 6 * 32) {
threads[0] = 1;
threads[2] = 32;
Expand All @@ -110,20 +110,20 @@ static void WarpImage(const float *src, int w, int h, int s, const float *u,
src_p[index * 4 + 1] = src_p[index * 4 + 2] = src_p[index * 4 + 3] = 0.f;
}
}
auto texDescr = cl::sycl::sampler(
auto texDescr = sycl::sampler(
sycl::coordinate_normalization_mode::unnormalized,
sycl::addressing_mode::clamp_to_edge, sycl::filtering_mode::linear);

auto texToWarp =
cl::sycl::image<2>(src_p, cl::sycl::image_channel_order::rgba,
cl::sycl::image_channel_type::fp32, range<2>(w, h),
sycl::image<2>(src_p, sycl::image_channel_order::rgba,
sycl::image_channel_type::fp32, range<2>(w, h),
range<1>(s * sizeof(sycl::float4)));

dpct::get_default_queue()
.submit([&](sycl::handler &cgh) {
auto texToWarp_acc =
texToWarp.template get_access<cl::sycl::float4,
cl::sycl::access::mode::read>(cgh);
texToWarp.template get_access<sycl::float4,
sycl::access::mode::read>(cgh);

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