Skip to content

[SYCL][NFC] Increase test coverage for inline asm feature #2124

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.

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

Conversation

fadeeval
Copy link
Contributor

@fadeeval fadeeval commented Jul 15, 2020

This PR adds inline_asm feature tests with branching and loop.

Signed-off-by: Aleksander Fadeev [email protected]

@fadeeval fadeeval force-pushed the private/fadeeval/Improve_test_coverage_for_inline_asm_feature branch 2 times, most recently from a6ba188 to 73874ae Compare July 20, 2020 10:19
@fadeeval fadeeval force-pushed the private/fadeeval/Improve_test_coverage_for_inline_asm_feature branch 2 times, most recently from 7320741 to 54b9173 Compare August 20, 2020 14:26
@fadeeval fadeeval marked this pull request as ready for review August 20, 2020 14:30
@fadeeval fadeeval requested a review from a team as a code owner August 20, 2020 14:30
@fadeeval fadeeval requested a review from alexbatashev August 20, 2020 14:30
@fadeeval fadeeval force-pushed the private/fadeeval/Improve_test_coverage_for_inline_asm_feature branch 2 times, most recently from 24bd023 to e99570a Compare August 20, 2020 15:41
Copy link
Contributor

@alexbatashev alexbatashev left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Just a few stylistic nits to better match overall project code style. Please, apply my comments to other places in code when appropriate.

cgh.parallel_for<KernelFunctor<T>>(
cl::sycl::range<1>{this->getOutputBufferSize()},
[=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
bool switch_field = false;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
bool switch_field = false;
bool SwitchField = false;

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can we make this variable a kernel argument? Otherwise, it is too easy for compiler to optimize everything to just one single store, I guess

@alexbatashev alexbatashev changed the title Increase test coverage for inline asm feature [SYCL][NFC] Increase test coverage for inline asm feature Aug 20, 2020
@fadeeval
Copy link
Contributor Author

I don't like the format offered by clang-format in one position, so that the clang-format-check does not pass.

@alexbatashev
Copy link
Contributor

I don't like the format offered by clang-format in one position, so that the clang-format-check does not pass.

You can wrap this place in

// clang-format off
some code
// clang-format on

Signed-off-by: Aleksander Fadeev <[email protected]>
Add following tests:
inline asm for one branch of if
inline asm for loop
two tests of inline asm for big block of code

Signed-off-by: Aleksander Fadeev <[email protected]>
Signed-off-by: Aleksander Fadeev <[email protected]>
Signed-off-by: Aleksander Fadeev <[email protected]>
Signed-off-by: Aleksander Fadeev <[email protected]>
@fadeeval fadeeval force-pushed the private/fadeeval/Improve_test_coverage_for_inline_asm_feature branch from 7397a24 to d776183 Compare August 27, 2020 12:05
Signed-off-by: Aleksander Fadeev <[email protected]>
Copy link
Contributor

@alexbatashev alexbatashev left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@AlexeySotkin
Copy link
Contributor

Please squash the commits and fix the typo in the commit message

@alexbatashev
Copy link
Contributor

Please squash the commits and fix the typo in the commit message

I think, @bader can squash commits from GitHub GUI and use PR description as a commit message, which is kinda what he always does. There's no need to squash the commits manually, restart tests and waste CPU cycles.

@romanovvlad romanovvlad merged commit ac0e692 into intel:sycl Aug 27, 2020
jsji pushed a commit that referenced this pull request Jan 4, 2024
…" (#2276)

Per discussion in Khronos and the spec itself, OpMatrixTimesScalar should not be used in OpenCL environment, as long as the same can be achieved either via an elementwise multiplication with a cooperative matrix object created from the scalar using OpConstantComposite or by iterating over the elements of the cooperative matrix to multiply each element by the scalar.

Spec: https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/KHR/SPV_KHR_cooperative_matrix.asciidoc

This reverts commit 6522c4ac4764cc1f251edd4266d06cba64549928.

Original commit:
KhronosGroup/SPIRV-LLVM-Translator@3b6ff3e
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

5 participants