[SYCL] Don't run SYCLOptimizeBarriersPass with O0 #20037

MrSidims · 2025-09-10T11:41:20Z

No description provided.

Signed-off-by: Sidorov, Dmitry <[email protected]>

MrSidims · 2025-09-10T11:42:12Z

Not entirely sure, how to add a test for this change - with O0 SYCL code won't be having memory semantic constants be propagated in the first place.

Signed-off-by: Sidorov, Dmitry <[email protected]>

maarquitos14 · 2025-09-10T12:36:34Z

Not entirely sure, how to add a test for this change - with O0 SYCL code won't be having memory semantic constants be propagated in the first place.

Maybe with some kind of log to make sure the barrier does not run?

Fznamznon · 2025-09-10T12:42:59Z

Not entirely sure, how to add a test for this change - with O0 SYCL code won't be having memory semantic constants be propagated in the first place.

Why do we disable it then?

MrSidims · 2025-09-10T15:55:51Z

Not entirely sure, how to add a test for this change - with O0 SYCL code won't be having memory semantic constants be propagated in the first place.

Why do we disable it then?

Because of the OpenMP compiler, that inserts barriers implicitly (with the constants) and OpenMP offloading reuses SYCL offloading model, making this pass running.

Fznamznon · 2025-09-10T15:59:57Z

Not entirely sure, how to add a test for this change - with O0 SYCL code won't be having memory semantic constants be propagated in the first place.

Why do we disable it then?

Because of the OpenMP compiler, that inserts barriers implicitly (with the constants) and OpenMP offloading reuses SYCL offloading model, making this pass running.

Can we use an OpenMP test?

MrSidims · 2025-09-10T16:11:56Z

Not entirely sure, how to add a test for this change - with O0 SYCL code won't be having memory semantic constants be propagated in the first place.

Why do we disable it then?

Because of the OpenMP compiler, that inserts barriers implicitly (with the constants) and OpenMP offloading reuses SYCL offloading model, making this pass running.

Can we use an OpenMP test?

AFAIK OpenMP solution is a part of our downstream, so adding a test here won't help much :) The only other solution I see is to add an LLVM IR test with opt -O0 passes as a RUN command, but not sure if it's 'LLVM way' of creating unit tests.

MrSidims · 2025-09-10T16:13:32Z

Not entirely sure, how to add a test for this change - with O0 SYCL code won't be having memory semantic constants be propagated in the first place.

Why do we disable it then?

Because of the OpenMP compiler, that inserts barriers implicitly (with the constants) and OpenMP offloading reuses SYCL offloading model, making this pass running.

Can we use an OpenMP test?

AFAIK OpenMP solution is a part of our downstream, so adding a test here won't help much :) The only other solution I see is to add an LLVM IR test with opt -O0 passes as a RUN command, but not sure if it's 'LLVM way' of creating unit tests.

Actually, I have an idea for SYCL'ish test, let me try it.

sycl/test/check_device_code/narrow-barrier-explicit-spirv-call.cpp

Signed-off-by: Sidorov, Dmitry <[email protected]>

github-actions · 2025-09-10T19:04:19Z

@intel/llvm-gatekeepers please consider merging

maarquitos14 · 2025-09-11T07:15:21Z

sycl/test/check_device_code/narrow-barrier-explicit-spirv-call.cpp

+
+int main(int argc, char *argv[]) {
+  sycl::queue q{sycl::property::queue::enable_profiling{}};
+  float *sum = sycl::malloc_shared<float>(NUMBER_OF_WORK_GROUPS, q);


This has to be freed.

the test is removed

sycl/test/check_device_code/narrow-barrier-explicit-spirv-call.cpp

Fznamznon · 2025-09-11T08:27:59Z

sycl/test/check_device_code/narrow-barrier-explicit-spirv-call.cpp

@@ -0,0 +1,68 @@
+// RUN: %clangxx -fsycl-device-only -fsycl-unnamed-lambda -S -Xclang -emit-llvm -Xclang -no-enable-noundef-analysis -O2 %s -o - | FileCheck %s --check-prefix=CHECK-O2


Sorry, I'm not really happy with the test.
The problem I see here is that the tests that reside in sycl/test/check_device_code/ are intended to check that SYCL library is using front-end provided hooks/attributes correctly, this test is explicitly using SPIR-V builtin which should not happen in these tests.
If we simply want to check that the builtin is treated properly by the optimizations, it is better to add a simpler smaller test to clang/CodeGenSYCL. This way the test will also be much faster, since inclusion of the whole sycl.hpp won't be needed. If we want to see that SYCL application as user would write it generates the right code, we can leave the test here but we should not use the builtin explicitly.

Strongly disagree, the test shouldn't be placed in clang/CodeGenSYCL . There should be frontend tests only, they must not test optimizations happening later.

this test is explicitly using SPIR-V builtin which should not happen in these tests

using sycl::group_builtin won't expose missing O0 check in the backend utils, as the SYCL API is using several wrapper functions to get scope, memory semantics etc before calling spirv builtin.

Strongly disagree, the test shouldn't be placed in clang/CodeGenSYCL . There should be frontend tests only, they must not test optimizations happening later.

The patch makes sure that clang doesn't run a pass with O0. I don't think a test for that patch actually checks the optimization's behavior.

Didn't know, that there is a test for the pipeline check, moved the check there, thanks!

Strongly disagree, the test shouldn't be placed in clang/CodeGenSYCL . There should be frontend tests only, they must not test optimizations happening later.

Never heard of this limitation before. I see 200 tests in clang/test/CodeGen use -O2. Most of these tests check LLVM IR after optimizations, not the LLVM IR emitted by the front-end.

Signed-off-by: Sidorov, Dmitry <[email protected]>

clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp

Co-authored-by: Marcos Maronas <[email protected]>

[SYCL] Don't run SYCLOptimizeBarriersPass with O0

b78605b

Signed-off-by: Sidorov, Dmitry <[email protected]>

MrSidims requested review from a team as code owners September 10, 2025 11:41

MrSidims temporarily deployed to WindowsCILock September 10, 2025 11:41 — with GitHub Actions Inactive

MrSidims had a problem deploying to WindowsCILock September 10, 2025 12:17 — with GitHub Actions Error

MrSidims temporarily deployed to WindowsCILock September 10, 2025 12:17 — with GitHub Actions Inactive

format

96a7a7a

Signed-off-by: Sidorov, Dmitry <[email protected]>

MrSidims temporarily deployed to WindowsCILock September 10, 2025 12:27 — with GitHub Actions Inactive

maarquitos14 approved these changes Sep 10, 2025

View reviewed changes

MrSidims temporarily deployed to WindowsCILock September 10, 2025 13:12 — with GitHub Actions Inactive

MrSidims had a problem deploying to WindowsCILock September 10, 2025 13:12 — with GitHub Actions Failure

add test

1aeedf7

MrSidims requested a review from a team as a code owner September 10, 2025 17:09

MrSidims requested a review from steffenlarsen September 10, 2025 17:09

MrSidims had a problem deploying to WindowsCILock September 10, 2025 17:09 — with GitHub Actions Error

MrSidims commented Sep 10, 2025

View reviewed changes

sycl/test/check_device_code/narrow-barrier-explicit-spirv-call.cpp Outdated Show resolved Hide resolved

MrSidims had a problem deploying to WindowsCILock September 10, 2025 17:15 — with GitHub Actions Error

adjust test

68a84f0

Signed-off-by: Sidorov, Dmitry <[email protected]>

MrSidims force-pushed the run-barriers-with-o0 branch from 38937e9 to 68a84f0 Compare September 10, 2025 17:16

MrSidims temporarily deployed to WindowsCILock September 10, 2025 17:16 — with GitHub Actions Inactive

MrSidims temporarily deployed to WindowsCILock September 10, 2025 17:49 — with GitHub Actions Inactive

steffenlarsen approved these changes Sep 11, 2025

View reviewed changes

maarquitos14 reviewed Sep 11, 2025

View reviewed changes

Fznamznon reviewed Sep 11, 2025

View reviewed changes

move the test

73266d4

Signed-off-by: Sidorov, Dmitry <[email protected]>

MrSidims temporarily deployed to WindowsCILock September 11, 2025 13:41 — with GitHub Actions Inactive

MrSidims requested a review from Fznamznon September 11, 2025 14:27

MrSidims temporarily deployed to WindowsCILock September 11, 2025 14:31 — with GitHub Actions Inactive

maarquitos14 reviewed Sep 11, 2025

View reviewed changes

clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp Outdated Show resolved Hide resolved

Update clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp

df7e2a8

Co-authored-by: Marcos Maronas <[email protected]>

MrSidims had a problem deploying to WindowsCILock September 11, 2025 14:56 — with GitHub Actions Failure

Fznamznon approved these changes Sep 11, 2025

View reviewed changes

MrSidims temporarily deployed to WindowsCILock September 11, 2025 15:37 — with GitHub Actions Inactive

		@@ -0,0 +1,68 @@
		// RUN: %clangxx -fsycl-device-only -fsycl-unnamed-lambda -S -Xclang -emit-llvm -Xclang -no-enable-noundef-analysis -O2 %s -o - \| FileCheck %s --check-prefix=CHECK-O2

[SYCL] Don't run SYCLOptimizeBarriersPass with O0 #20037

Are you sure you want to change the base?

[SYCL] Don't run SYCLOptimizeBarriersPass with O0 #20037

Conversation

MrSidims commented Sep 10, 2025

Uh oh!

MrSidims commented Sep 10, 2025

Uh oh!

maarquitos14 commented Sep 10, 2025

Uh oh!

Fznamznon commented Sep 10, 2025

Uh oh!

MrSidims commented Sep 10, 2025

Uh oh!

Fznamznon commented Sep 10, 2025

Uh oh!

MrSidims commented Sep 10, 2025

Uh oh!

MrSidims commented Sep 10, 2025

Uh oh!

Uh oh!

github-actions bot commented Sep 10, 2025

Uh oh!

maarquitos14 Sep 11, 2025

Choose a reason for hiding this comment

Uh oh!

MrSidims Sep 11, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Fznamznon Sep 11, 2025

Choose a reason for hiding this comment

Uh oh!

MrSidims Sep 11, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

MrSidims Sep 11, 2025

Choose a reason for hiding this comment

Uh oh!

Fznamznon Sep 11, 2025

Choose a reason for hiding this comment

Uh oh!

MrSidims Sep 11, 2025

Choose a reason for hiding this comment

Uh oh!

bader Sep 11, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

MrSidims Sep 11, 2025 •

edited

Loading