[sycl-post-link] Don't remove llvm.compiler.used for NVPTX #15224

hdelan · 2024-08-29T11:05:29Z

llvm.compiler.used is a global symbol which tells the compiler not to touch some other global symbols until backend lowering. The llvm.compiler.used symbol itself is thus removed automatically during lowering for NVPTX and AMDGCN. Removal at sycl-post-link, prior to lowering, was causing issues for these backends, where symbols protected by llvm.compiler.used were getting removed after sycl-post-link and before lowering.

We retain the current behaviour for SPIR-V, as SPIR-V generation is handled in llvm-spirv anyway, not in the LLVM SPIR-V backend.

Also adds tests to make sure static device_globals are handled properly for NVPTX/AMDGCN.

aelovikov-intel · 2024-08-29T16:24:18Z

sycl/test/check_device_code/device_global_static.cpp

@@ -0,0 +1,29 @@
+// Tests that the llvm.compiler.used symbol, which is used to implement static


Why sycl/test instead of sycl-post-link test?

It's a bit of a mess but the llvm.used symbol is handled in different parts of the compilation depending on backend. So it's important to make sure that the tests are working for an entire compilation pipeline, not just for a particular stage of the action.

What is the latest stage it is being removed? Why can't we have test for that component? Is the "mess" final, or do we expect pipelines to be unified?

llvm.used should survive to final linking, llvm.compiler.used is okay for intermediary tools to drop: https://llvm.org/docs/LangRef.html#the-llvm-used-global-variable

Do we need a test when programmer uses __attribute__((used)) on something to ensure that something isn't being dropped?

Device image compression will be optional (#15124) - unless you specify --offload-compress in CLI, by default, we won't compress images. So, image compression might not be an issue for this particular test.

If image compression is "opt-in" and we will not opt in in this test, then why is this an issue @aelovikov-intel ? It seems unlikely that binary formatting is going to change drastically in the next while, and if it does then I think this test will be the least of our problems.

I am using strings as it's a simpler, more portable way to find a global symbol, using the current device code formats, than using proprietary binutils that are likely to belong to proprietary toolkits. For instance cuobjdump is not necessarily distributed in every CUDA version. Relying on a tool like this may break CI with a CUDA toolkit version change. Likewise I don't want to rely on a tool like spirv-dis which often breaks and doesn't keep up to date with the latest spir-v features.

Thanks for response @uditagarwal97 :)

I've updated the test to use llvm-strings instead of strings. Is this sufficient @aelovikov-intel ?

If I just delete the test will you approve this PR? I don't think this discussion is going anywhere

Yes.

If image compression is "opt-in" and we will not opt in in this test, then why is this an issue @aelovikov-intel ?

Because that's an implementation detail of something unrelated to your test and it must not make any assumptions of it.

clang/lib/CodeGen/CodeGenModule.cpp

llvm/tools/sycl-post-link/sycl-post-link.cpp

asudarsa

Looks good overall. Few minor comments. Thanks

clang/lib/CodeGen/CodeGenModule.cpp

These are removed automatically during lowering.

This is removed at lowering for NVPTX/AMDGCN. TODO: check that this is working for SPIR-V.

sycl post link no longer removes llvm.compiler.used.

For spirv remove the llvm.compiler.used symbol at sycl-post-link. This is because we use llvm-spirv for lowering, not the SPIR-V backend.

asudarsa · 2024-09-04T15:42:24Z

llvm/tools/sycl-post-link/sycl-post-link.cpp

-  // used inside the device code after they have been removed from
-  // "llvm.compiler.used" they can be erased safely.
-  Modified |= removeDeviceGlobalFromCompilerUsed(*M.get());
+  // to keep the optimizer from wrongfully removing them. llvm.compiler.used


We intend to replace the translator with SPIR-V backend eventually. Can we remove the check at that point of time?

Thanks

Once the SPIRV backend is being used instead of llvm-spirv I think it'd be better to remove the llvm.compiler.used symbol in the backend, instead of in sycl-post-link, although some might disagree. I think it's OK to keep the current behaviour for the moment unless it becomes a problem later.

asudarsa

Changes look good to me. Just one clarification requested. Thanks

Make sure that `strings` can successfully detect the symbol in an object file.

aelovikov-intel

I'd prefer the test to be removed, but if you really want to keep it so be it.

hdelan · 2024-09-06T08:50:54Z

Thanks @aelovikov-intel I'd like to keep it in if possible. I can remove at any point in time if it becomes a problem. Ping @intel/llvm-gatekeepers this can be merged, thanks.

sarnex · 2024-09-06T14:09:33Z

@hdelan The test is failing on HIP in postcommit, can you take a look and disable the test if the fix isn't quick? Thx

https://github.com/intel/llvm/actions/runs/10735485792/job/29773246908

FAIL: SYCL :: DeviceGlobal/device_global_static.cpp (413 of 2168)
******************** TEST 'SYCL :: DeviceGlobal/device_global_static.cpp' FAILED ********************
Exit Code: -6

Command Output (stdout):
--
# RUN: at line 1
/__w/llvm/llvm/toolchain/bin//clang++  -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa  /__w/llvm/llvm/llvm/sycl/test-e2e/DeviceGlobal/device_global_static.cpp -o /__w/llvm/llvm/build-e2e/DeviceGlobal/Output/device_global_static.cpp.tmp.out
# executed command: /__w/llvm/llvm/toolchain/bin//clang++ -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa /__w/llvm/llvm/llvm/sycl/test-e2e/DeviceGlobal/device_global_static.cpp -o /__w/llvm/llvm/build-e2e/DeviceGlobal/Output/device_global_static.cpp.tmp.out
# note: command had no output on stdout or stderr
# RUN: at line 2
env ONEAPI_DEVICE_SELECTOR=hip:gpu  /__w/llvm/llvm/build-e2e/DeviceGlobal/Output/device_global_static.cpp.tmp.out
# executed command: env ONEAPI_DEVICE_SELECTOR=hip:gpu /__w/llvm/llvm/build-e2e/DeviceGlobal/Output/device_global_static.cpp.tmp.out
# .---command stderr------------
# | <HIP>[ERROR]: 
# | UR HIP ERROR:
# | 	Value:           500
# | 	Name:            hipErrorNotFound
# | 	Description:     named symbol not found
# | 	Function:        getGlobalVariablePointer
# | 	Source Location: /__w/llvm/llvm/build/_deps/unified-runtime-src/source/adapters/hip/program.cpp:254
# | 
# | <HIP>[ERROR]: 
# | UR HIP ERROR:
# | 	Value:           UR_RESULT_ERROR_UNKNOWN
# | 	Function:        deviceGlobalCopyHelper
# | 	Source Location: /__w/llvm/llvm/build/_deps/unified-runtime-src/source/adapters/hip/enqueue.cpp:1691
# | 
# | terminate called after throwing an instance of 'sycl::_V1::exception'
# |   what():  Native API failed. Native API returns: 2147483646 (UR_RESULT_ERROR_UNKNOWN)
# `-----------------------------
# error: command failed with exit status: -6

hdelan · 2024-09-09T09:20:30Z

Thanks @sarnex . Have disabled test here

hdelan requested review from a team as code owners August 29, 2024 11:05

hdelan requested a review from aelovikov-intel August 29, 2024 11:05

hdelan mentioned this pull request Aug 29, 2024

[clang] Modify linkage and register initialization of device_global #15148

Closed

hdelan had a problem deploying to WindowsCILock August 29, 2024 11:06 — with GitHub Actions Error

hdelan force-pushed the dont-remove-llvm-compiler-used branch from 3cfff78 to 8ca0294 Compare August 29, 2024 11:14

hdelan had a problem deploying to WindowsCILock August 29, 2024 11:15 — with GitHub Actions Error

hdelan force-pushed the dont-remove-llvm-compiler-used branch from 8ca0294 to 9784bb5 Compare August 29, 2024 11:25

hdelan had a problem deploying to WindowsCILock August 29, 2024 11:25 — with GitHub Actions Failure

hdelan temporarily deployed to WindowsCILock August 29, 2024 12:10 — with GitHub Actions Inactive

hdelan requested a review from a team as a code owner August 29, 2024 15:22

hdelan had a problem deploying to WindowsCILock August 29, 2024 15:22 — with GitHub Actions Error

hdelan changed the title ~~[AMDGCN][NVPTX] Don't remove llvm.compiler.used for NVPTX/AMDGCN~~ [sycl-post-link] Don't remove llvm.compiler.used for NVPTX Aug 29, 2024

aelovikov-intel reviewed Aug 29, 2024

View reviewed changes

asudarsa reviewed Aug 29, 2024

View reviewed changes

clang/lib/CodeGen/CodeGenModule.cpp Outdated Show resolved Hide resolved

asudarsa reviewed Aug 29, 2024

View reviewed changes

llvm/tools/sycl-post-link/sycl-post-link.cpp Outdated Show resolved Hide resolved

asudarsa reviewed Aug 29, 2024

View reviewed changes

elizabethandrews reviewed Aug 30, 2024

View reviewed changes

clang/lib/CodeGen/CodeGenModule.cpp Outdated Show resolved Hide resolved

hdelan added 2 commits September 2, 2024 11:49

Don't remove llvm.compiler.used for NVPTX/AMDGCN

72a6032

These are removed automatically during lowering.

Don't remove llvm.compiler.used at sycl post link

74fa7ba

This is removed at lowering for NVPTX/AMDGCN. TODO: check that this is working for SPIR-V.

hdelan force-pushed the dont-remove-llvm-compiler-used branch from 18fe1c6 to 74fa7ba Compare September 3, 2024 15:24

hdelan had a problem deploying to WindowsCILock September 3, 2024 15:25 — with GitHub Actions Failure

hdelan had a problem deploying to WindowsCILock September 3, 2024 16:20 — with GitHub Actions Error

hdelan added 2 commits September 3, 2024 17:29

Update sycl post link tests

3f9ae01

sycl post link no longer removes llvm.compiler.used.

Reinstate llvm.compiler.used removal at SPL

1b8c225

For spirv remove the llvm.compiler.used symbol at sycl-post-link. This is because we use llvm-spirv for lowering, not the SPIR-V backend.

hdelan had a problem deploying to WindowsCILock September 3, 2024 16:48 — with GitHub Actions Error

Undo sycl-post-link tests

4b06c76

hdelan force-pushed the dont-remove-llvm-compiler-used branch from eeff007 to 4b06c76 Compare September 3, 2024 16:51

hdelan temporarily deployed to WindowsCILock September 3, 2024 16:53 — with GitHub Actions Inactive

hdelan temporarily deployed to WindowsCILock September 3, 2024 18:08 — with GitHub Actions Inactive

hdelan requested review from aelovikov-intel and asudarsa September 4, 2024 14:24

asudarsa reviewed Sep 4, 2024

View reviewed changes

asudarsa approved these changes Sep 4, 2024

View reviewed changes

Check that strings can find the symbol

d3c9529

Make sure that `strings` can successfully detect the symbol in an object file.

hdelan temporarily deployed to WindowsCILock September 4, 2024 15:49 — with GitHub Actions Inactive

hdelan temporarily deployed to WindowsCILock September 4, 2024 18:11 — with GitHub Actions Inactive

Change strings usage to llvm-strings

e852cbd

hdelan temporarily deployed to WindowsCILock September 5, 2024 16:33 — with GitHub Actions Inactive

aelovikov-intel approved these changes Sep 5, 2024

View reviewed changes

hdelan temporarily deployed to WindowsCILock September 5, 2024 18:15 — with GitHub Actions Inactive

martygrant merged commit a8c60db into intel:sycl Sep 6, 2024
13 checks passed

hdelan mentioned this pull request Sep 9, 2024

Static Device Global Doesn't Work for HIP #15329

Open

aelovikov-intel mentioned this pull request Sep 11, 2024

[SYCL] Remove device_global_static.cpp test #15362

Merged

ldrumm mentioned this pull request Dec 23, 2024

[HIP] getGlobalVariablePointer error oneapi-src/unified-runtime#2475

Closed

sherholz-intel mentioned this pull request Feb 27, 2025

[HIP] device globals amdgcn-link error #17193

Open

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[sycl-post-link] Don't remove llvm.compiler.used for NVPTX #15224

[sycl-post-link] Don't remove llvm.compiler.used for NVPTX #15224

hdelan commented Aug 29, 2024 •

edited

Loading

aelovikov-intel Aug 29, 2024

hdelan Aug 29, 2024

aelovikov-intel Aug 29, 2024 •

edited

Loading

AaronBallman Aug 29, 2024

aelovikov-intel Aug 29, 2024

hdelan Sep 5, 2024 •

edited

Loading

hdelan Sep 5, 2024

hdelan Sep 5, 2024

aelovikov-intel Sep 5, 2024

aelovikov-intel Sep 5, 2024

asudarsa left a comment

asudarsa Sep 4, 2024

hdelan Sep 4, 2024 •

edited

Loading

asudarsa left a comment

aelovikov-intel left a comment

hdelan commented Sep 6, 2024

sarnex commented Sep 6, 2024

hdelan commented Sep 9, 2024

		@@ -0,0 +1,29 @@
		// Tests that the llvm.compiler.used symbol, which is used to implement static

[sycl-post-link] Don't remove llvm.compiler.used for NVPTX #15224

[sycl-post-link] Don't remove llvm.compiler.used for NVPTX #15224

Conversation

hdelan commented Aug 29, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

aelovikov-intel Aug 29, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

hdelan Sep 5, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

asudarsa left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

hdelan Sep 4, 2024 • edited Loading

Choose a reason for hiding this comment

asudarsa left a comment

Choose a reason for hiding this comment

aelovikov-intel left a comment

Choose a reason for hiding this comment

hdelan commented Sep 6, 2024

sarnex commented Sep 6, 2024

hdelan commented Sep 9, 2024

hdelan commented Aug 29, 2024 •

edited

Loading

aelovikov-intel Aug 29, 2024 •

edited

Loading

hdelan Sep 5, 2024 •

edited

Loading

hdelan Sep 4, 2024 •

edited

Loading