[SYCL] Option to disable alloca address space for sret arguments #17976

Fznamznon · 2025-04-11T13:23:17Z

Recent community change llvm/llvm-project#114062 enabled the use of
alloca address space for sret arguments. This causes several issues for sycl, particularly for the SPIR
target where this leads to invalid address-space-casts from the local address space.

The new option -foffload-use-alloca-addrspace-for-srets is TRUE by default (and produces the current
community behavior) and is set to FALSE in sycl device compilation modes (where the prior behavior is
retained). The commit also reverts some test changes made to reflect the current community behavior.

This reverts commit afe4c53, reversing changes made to bd2d590.

This reverts commit 3c3ad27.

aelovikov-intel

SYCL RT LGTM, but this PR needs a better description.

Fznamznon

Thank you for picking this up!
See also #17976 (comment)

Fznamznon · 2025-04-22T08:54:04Z

sycl/test/check_device_code/extensions/address_cast.cpp

@@ -13,63 +13,63 @@ using namespace sycl::ext::oneapi::experimental;

 namespace static_as_cast {
 // CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast19to_global_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
-// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 captures(none) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {{.*}}{
+// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {{.*}}{


It would be great if we could check both modes.

dm-vodopyanov · 2025-04-22T12:13:51Z

Should this file be also updated? (these lines should be removed)

llvm/sycl/cts_exclude_filter/compfails

Lines 3 to 4 in 46da71a

    
           # CMPLRLLVM-66370 
        
           hierarchical

Fznamznon · 2025-04-22T12:29:26Z

Should this file be also updated? (these lines should be removed)

llvm/sycl/cts_exclude_filter/compfails

Lines 3 to 4 in 46da71a

# CMPLRLLVM-66370

hierarchical

Yes, I would expect this problem to be fixed/hidden by the option until we figure out how to deal with alloca for srets properly.

mdtoguchi · 2025-04-22T18:05:32Z

clang/lib/Driver/ToolChains/Clang.cpp

+    // Disable this option for SPIR targets.
+    // TODO:  This needs to be re-enabled once we have a real fix.
+    CmdArgs.push_back("-fno-offload-use-alloca-addrspace-for-srets");
+  }


Updated change LGTM

aelovikov-intel

sycl/cts_exclude_filter/compfails LGTM

premanandrao · 2025-04-22T20:17:56Z

@intel/llvm-gatekeepers, can this please be merged?

@tahonermann reviewed and approved this for CFE. (Since I am running this PR, I don't want to self-approve this.)
@Fznamznon has reviewed it, and her review comments have been addressed.

premanandrao · 2025-04-22T20:18:43Z

The code-formatting error is about a test case, not source code. Not sure why it is being flagged.

aelovikov-intel · 2025-04-22T20:30:17Z

The code-formatting error is about a test case, not source code. Not sure why it is being flagged.

Tests should be properly formatted too... Also note

llvm/sycl/test/.clang-format

Line 2 in a53dacf

    
           CommentPragmas: "RUN|FAIL|REQUIRES|UNSUPPORTED|CHECK|expected-|update_cc_test_checks.py"

Maybe renaming the FileCheck's label to contain CHECK would silence the clang-format.

premanandrao · 2025-04-22T20:34:35Z

Tests should be properly formatted too...

Sure, but the changes it is suggesting didn't make sense to me.

Maybe renaming the FileCheck's label to contain CHECK would silence the clang-format.

This may be it. Thanks, let me try.

premanandrao · 2025-04-22T20:48:01Z

Thank you @aelovikov-intel, that was it.

premanandrao · 2025-04-22T22:37:37Z

Thank you @aelovikov-intel

Fznamznon added 2 commits April 10, 2025 05:53

Initial change for testing

dbe18de

Fix crashes related to lifetimes emission

53d4d29

Fznamznon requested review from a team as code owners April 11, 2025 13:23

Fznamznon had a problem deploying to WindowsCILock April 11, 2025 13:23 — with GitHub Actions Failure

Fznamznon marked this pull request as draft April 11, 2025 13:23

Fznamznon temporarily deployed to WindowsCILock April 11, 2025 13:45 — with GitHub Actions Inactive

Fznamznon temporarily deployed to WindowsCILock April 11, 2025 14:08 — with GitHub Actions Inactive

Additional changes to make the option work in non-SYCL modes

5e569a7

premanandrao had a problem deploying to WindowsCILock April 20, 2025 03:02 — with GitHub Actions Error

Fix clang-format issues

9d98cde

premanandrao had a problem deploying to WindowsCILock April 20, 2025 03:12 — with GitHub Actions Failure

premanandrao temporarily deployed to WindowsCILock April 20, 2025 04:22 — with GitHub Actions Inactive

premanandrao temporarily deployed to WindowsCILock April 20, 2025 04:45 — with GitHub Actions Inactive

premanandrao had a problem deploying to WindowsCILock April 20, 2025 04:45 — with GitHub Actions Failure

premanandrao added 4 commits April 20, 2025 19:15

Revert "Merge commit '3a353b1faa83' into llvmspirv_pulldown"

3c3ad27

This reverts commit afe4c53, reversing changes made to bd2d590.

Revert "Revert "Merge commit '3a353b1faa83' into llvmspirv_pulldown""

4dee04d

This reverts commit 3c3ad27.

Fix some tests that were changed earlier

408135b

One more test

c790e52

premanandrao temporarily deployed to WindowsCILock April 21, 2025 03:50 — with GitHub Actions Inactive

premanandrao temporarily deployed to WindowsCILock April 21, 2025 04:49 — with GitHub Actions Inactive

premanandrao temporarily deployed to WindowsCILock April 21, 2025 05:01 — with GitHub Actions Inactive

premanandrao marked this pull request as ready for review April 21, 2025 14:15

premanandrao requested review from a team as code owners April 21, 2025 14:15

premanandrao requested a review from aelovikov-intel April 21, 2025 14:15

premanandrao changed the title ~~[TESTING] Test option to disable alloca address space for sret arguments~~ [SYCL] Option to disable alloca address space for sret arguments Apr 21, 2025

aelovikov-intel approved these changes Apr 21, 2025

View reviewed changes

premanandrao temporarily deployed to WindowsCILock April 22, 2025 03:42 — with GitHub Actions Inactive

Fznamznon commented Apr 22, 2025

View reviewed changes

Disable the option when target is SPIR

0c4b489

premanandrao temporarily deployed to WindowsCILock April 22, 2025 15:49 — with GitHub Actions Inactive

premanandrao temporarily deployed to WindowsCILock April 22, 2025 16:24 — with GitHub Actions Inactive

premanandrao temporarily deployed to WindowsCILock April 22, 2025 17:11 — with GitHub Actions Inactive

mdtoguchi reviewed Apr 22, 2025

View reviewed changes

Remove test exclusion; modify test to check other mode

6760adb

premanandrao requested a review from a team as a code owner April 22, 2025 18:30

premanandrao temporarily deployed to WindowsCILock April 22, 2025 18:30 — with GitHub Actions Inactive

aelovikov-intel approved these changes Apr 22, 2025

View reviewed changes

premanandrao temporarily deployed to WindowsCILock April 22, 2025 19:36 — with GitHub Actions Inactive

premanandrao temporarily deployed to WindowsCILock April 22, 2025 19:45 — with GitHub Actions Inactive

Rename FileCheck string to please clang-format

8edcba2

premanandrao temporarily deployed to WindowsCILock April 22, 2025 20:42 — with GitHub Actions Inactive

premanandrao had a problem deploying to WindowsCILock April 22, 2025 21:43 — with GitHub Actions Error

premanandrao temporarily deployed to WindowsCILock April 22, 2025 21:47 — with GitHub Actions Inactive

premanandrao temporarily deployed to WindowsCILock April 22, 2025 21:59 — with GitHub Actions Inactive

aelovikov-intel merged commit 7175536 into intel:sycl Apr 22, 2025
50 of 54 checks passed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL] Option to disable alloca address space for sret arguments #17976

[SYCL] Option to disable alloca address space for sret arguments #17976

Fznamznon commented Apr 11, 2025 •

edited by premanandrao

Loading

aelovikov-intel left a comment

Fznamznon left a comment

Fznamznon Apr 22, 2025

dm-vodopyanov commented Apr 22, 2025

Fznamznon commented Apr 22, 2025

mdtoguchi Apr 22, 2025

aelovikov-intel left a comment

premanandrao commented Apr 22, 2025

premanandrao commented Apr 22, 2025

aelovikov-intel commented Apr 22, 2025 •

edited

Loading

premanandrao commented Apr 22, 2025

premanandrao commented Apr 22, 2025

premanandrao commented Apr 22, 2025

[SYCL] Option to disable alloca address space for sret arguments #17976

[SYCL] Option to disable alloca address space for sret arguments #17976

Conversation

Fznamznon commented Apr 11, 2025 • edited by premanandrao Loading

aelovikov-intel left a comment

Choose a reason for hiding this comment

Fznamznon left a comment

Choose a reason for hiding this comment

Fznamznon Apr 22, 2025

Choose a reason for hiding this comment

dm-vodopyanov commented Apr 22, 2025

Fznamznon commented Apr 22, 2025

mdtoguchi Apr 22, 2025

Choose a reason for hiding this comment

aelovikov-intel left a comment

Choose a reason for hiding this comment

premanandrao commented Apr 22, 2025

premanandrao commented Apr 22, 2025

aelovikov-intel commented Apr 22, 2025 • edited Loading

premanandrao commented Apr 22, 2025

premanandrao commented Apr 22, 2025

premanandrao commented Apr 22, 2025

Fznamznon commented Apr 11, 2025 •

edited by premanandrao

Loading

aelovikov-intel commented Apr 22, 2025 •

edited

Loading