Skip to content

fix the SYCL DCT sample #2693

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

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
Open

fix the SYCL DCT sample #2693

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
33 changes: 17 additions & 16 deletions DirectProgramming/C++SYCL/SpectralMethods/DiscreteCosineTransform/src/DCT.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,8 +47,8 @@ void MatrixTranspose(float x[block_size], float xinv[block_size]) {
}

// Multiply two matrices x and y and write output to xy
SYCL_EXTERNAL void MatrixMultiply(multi_ptr<const float, access::address_space::global_space, (sycl::access::decorated)2> x,
multi_ptr<const float, access::address_space::global_space, (sycl::access::decorated)2> y,
SYCL_EXTERNAL void MatrixMultiply(const float x[block_size],
const float y[block_size],
float xy[block_size]) {
for (int i = 0; i < block_dims; ++i) {
for (int j = 0; j < block_dims; ++j) {
Expand All @@ -61,7 +61,8 @@ SYCL_EXTERNAL void MatrixMultiply(multi_ptr<const float, access::address_space::
}

// Processes an individual 8x8 subset of image data
SYCL_EXTERNAL void ProcessBlock(multi_ptr<const rgb, access::address_space::global_space, sycl::access::decorated::no> indataset, rgb* outdataset,
SYCL_EXTERNAL void ProcessBlock(multi_ptr<const rgb, access::address_space::global_space, sycl::access::decorated::no> indataset,
multi_ptr<rgb, access::address_space::global_space, sycl::access::decorated::no> outdataset,
multi_ptr<const float, access::address_space::global_space, sycl::access::decorated::no> dct,
multi_ptr<const float, access::address_space::global_space, sycl::access::decorated::no> dctinv,
int start_index, int width) {
Expand Down Expand Up @@ -111,8 +112,8 @@ SYCL_EXTERNAL void ProcessBlock(multi_ptr<const rgb, access::address_space::glob

// Computation of the discrete cosine transform of the image section of size
// 8x8 for red values
MatrixMultiply(dct, red_input, temp);
MatrixMultiply(temp, dctinv, interim);
MatrixMultiply(dct.get(), red_input, temp);
MatrixMultiply(temp, dctinv.get(), interim);

// Computation of quantization phase using the quantization matrix
for (int i = 0; i < block_size; ++i)
Expand All @@ -123,8 +124,8 @@ SYCL_EXTERNAL void ProcessBlock(multi_ptr<const rgb, access::address_space::glob
interim[i] = sycl::floor((interim[i] * quant[i]) + 0.5f);

// Computation of Inverse Discrete Cosine Transform (IDCT)
MatrixMultiply(dctinv, interim, temp);
MatrixMultiply(temp, dct, product);
MatrixMultiply(dctinv.get(), interim, temp);
MatrixMultiply(temp, dct.get(), product);

// Translating the pixels values from [-128, 127] range to [0, 255] range
// and writing to output image data
Expand All @@ -146,8 +147,8 @@ SYCL_EXTERNAL void ProcessBlock(multi_ptr<const rgb, access::address_space::glob

// Computation of the discrete cosine transform of the image section of size
// 8x8 for blue values
MatrixMultiply(dct, blue_input, temp);
MatrixMultiply(temp, dctinv, interim);
MatrixMultiply(dct.get(), blue_input, temp);
MatrixMultiply(temp, dctinv.get(), interim);

// Computation of quantization phase using the quantization matrix
for (int i = 0; i < block_size; ++i)
Expand All @@ -158,8 +159,8 @@ SYCL_EXTERNAL void ProcessBlock(multi_ptr<const rgb, access::address_space::glob
interim[i] = sycl::floor((interim[i] * quant[i]) + 0.5f);

// Computation of Inverse Discrete Cosine Transform (IDCT)
MatrixMultiply(dctinv, interim, temp);
MatrixMultiply(temp, dct, product);
MatrixMultiply(dctinv.get(), interim, temp);
MatrixMultiply(temp, dct.get(), product);

// Translating the pixels values from [-128, 127] range to [0, 255] range
// and writing to output image data
Expand All @@ -181,8 +182,8 @@ SYCL_EXTERNAL void ProcessBlock(multi_ptr<const rgb, access::address_space::glob

// Computation of the discrete cosine transform of the image section of size
// 8x8 for green values
MatrixMultiply(dct, green_input, temp);
MatrixMultiply(temp, dctinv, interim);
MatrixMultiply(dct.get(), green_input, temp);
MatrixMultiply(temp, dctinv.get(), interim);

// Computation of quantization phase using the quantization matrix
for (int i = 0; i < block_size; ++i)
Expand All @@ -193,8 +194,8 @@ SYCL_EXTERNAL void ProcessBlock(multi_ptr<const rgb, access::address_space::glob
interim[i] = sycl::floor((interim[i] * quant[i]) + 0.5f);

// Computation of Inverse Discrete Cosine Transform (IDCT)
MatrixMultiply(dctinv, interim, temp);
MatrixMultiply(temp, dct, product);
MatrixMultiply(dctinv.get(), interim, temp);
MatrixMultiply(temp, dct.get(), product);

// Translating the pixels values from [-128, 127] range to [0, 255] range
// and writing to output image data
Expand Down Expand Up @@ -235,7 +236,7 @@ void ProcessImage(rgb* indataset, rgb* outdataset, int width, int height) {
h.parallel_for(
range<2>(width / block_dims, height / block_dims), [=](auto idx) {
int start_index = idx[0] * block_dims + idx[1] * block_dims * width;
ProcessBlock(i_acc.get_multi_ptr<sycl::access::decorated::no>(), o_acc.get_multi_ptr<sycl::access::decorated::no>().get(),
ProcessBlock(i_acc.get_multi_ptr<sycl::access::decorated::no>(), o_acc.get_multi_ptr<sycl::access::decorated::no>(),
d_acc.get_multi_ptr<sycl::access::decorated::no>(), di_acc.get_multi_ptr<sycl::access::decorated::no>(), start_index,
width);
});
Expand Down