[SYCL] Switch to using integration footer by default by dm-vodopyanov · Pull Request #3777 · intel/llvm

dm-vodopyanov · 2021-05-18T15:49:09Z

Added spec_constant_integration.hpp header file, which is included from
integration footer.
Added get_spec_constant_symbolic_ID_wrapper which in combination
with spec_constant_integration.hpp allows us to avoid exploiting a UB
related to instantiating a template before we have seen all its
specializations.

Replaced -fsycl-use-footer compiler flag with -fno-sycl-use-footer
flag with the opposite semantics.

…e header This patch moves functions which use `__builtin_unique_stable_name` to a separate header which should only be included to integration footer. This was made to get rid of UB when the compiler sees defenition of the template specialization before seeing a declaration.

bader · 2021-05-18T16:31:59Z

This patch moves functions which use __builtin_unique_stable_name to a
separate header which should only be included to integration footer.

This description is not accurate KernelInfoData method still uses __builtin_unique_stable_name:

llvm/sycl/include/CL/sycl/detail/kernel_desc.hpp

Line 102 in 22f9009

static constexpr auto n = __builtin_unique_stable_name(T);

romanovvlad · 2021-05-19T04:49:53Z

This patch moves functions which use __builtin_unique_stable_name to a
separate header which should only be included to integration footer.

This description is not accurate KernelInfoData method still uses __builtin_unique_stable_name:

llvm/sycl/include/CL/sycl/detail/kernel_desc.hpp

Line 102 in 22f9009

static constexpr auto n = __builtin_unique_stable_name(T);

Right, it is not directly related to __builtin_unique_stable_name, we rather need to guarantee that calls of get_spec_constant_symbolic_ID appear after specializations of this functions which defined in the integration footer.

…_stable_name-to-separate-header

…bled by default for any app

erichkeane · 2021-06-01T17:37:33Z

I don't think we should be putting the last patch into the repo, we are about a week and a half away from getting the actual implementation, and this will delay it.

…e_name-to-separate-header

erichkeane

Just 2 CFE test nits.

…din_unique_stable_name-to-separate-header

The primary function/wrapper which is used by DPC++ RT is left as-is, i.e. `get_spec_constant_symbolic_ID`. The internal function and its specialziations which are generated by the compiler renamed to `get_spec_constant_symbolic_ID_impl`.

romanovvlad

LGTM

AlexeySachkov · 2021-06-28T12:31:08Z

@AGindinson, @mdtoguchi, could you please take a look at the driver part? Basically, we have enabled integration footer by default and turned existing -fsycl-use-footer into -fno-sycl-use-footer and adjusted tests. However, I'm not exactly sure that we have enough coverage for this new option.

@AaronBallman, @elizabethandrews, @premanandrao, could you please take a look at FE part? The changes are quite simple: we have added a few more includes to emitted integration footer + did renaming get_spec_constant_symbolic_ID -> get_spec_constant_symbolic_ID_impl

elizabethandrews

FE changes LGTM. Thanks!

…din_unique_stable_name-to-separate-header

…eparate-header' of github.com:dm-vodopyanov/llvm into private/dvodopya/move-__buildin_unique_stable_name-to-separate-header

elizabethandrews

FE changes LGTM

romanovvlad

LGTM

* upstream/sycl: (649 commits) [SYCL][Driver][NFC] Update integration footer test for 32-bit host (intel#4039) [SYCL][L0] Initialize descriptor .stype and .pNext (intel#4032) [SYCL] Add sycl::kernel::get_kernel_bundle method (intel#3855) [SYCL] Add support for device UUID as a SYCL extension. (intel#3696) [SYCL][Matrix] Add spec document for the matrix extension interface and its first implementation for AMX (intel#3551) Fix debug build mangler test after PR#3992 (8f38045). (intel#4033) [Driver][SYCL] Restrict user -include file in final integration footer step (intel#4036) [SYCL] [Tests] Do not copy device binary image mocks (intel#4023) [SYCL][Doc] Update docs to reflect new compiler features (intel#4030) [SYCL][CUDA] cl_khr_fp16 extension connected to cuda PI. (intel#4029) [SYCL][NFC] Refactor RT unit tests (intel#4021) [SYCL] Switch to using integration footer by default (intel#3777) [SYCL][CUDA] Add the Use Default Stream property (intel#4004) Uplift GPU RT version for Linux to 21.24.20098 (intel#4003) [SYCL][CUDA] atomic_ref.fetch_add used for fp64 reduction if device.has(atomic64) (intel#3950) [Driver][SYCL] Differentiate host dependency link from regular host link (intel#4002) [SYCL][ESIMD] Support device half type in intrinsics. (intel#4024) [SYCL] Allow fpga_reg only for PODs and Trivially-copyable structs (intel#3643) [SYCL][FPGA] Restore legacy debug info version for the hardware (intel#3991) [SYCL][PI][L0] Force reset of memcpy command-list. (intel#4001) ...

For a vector type, add `SPV_EXT_long_vector` extension and prioritize it over `SPV_INTEL_vector_compute` whenever the vector type doesn't have a standard component size (2, 3, 4, 8, or 16). Original commit: KhronosGroup/SPIRV-LLVM-Translator@7d4b128df7e6ff3

#3829) #3777 added `SPV_EXT_long_vector` and gave it priority over `SPV_INTEL_vector_compute`: when both extensions are enabled and a vector has a non-standard component count, the translator now emits `CapabilityLongVectorEXT` instead of `CapabilityVectorAnyINTEL`. This regresses `VectorCompute` modules. The offloading flow (producing both `VectorCompute` and non-`VectorCompute` modules) applies a single, fixed extension list uniformly to every device image, so a `VectorCompute` module is translated with both `SPV_INTEL_vector_compute` and `SPV_EXT_long_vector` enabled. Before #3777 such a module emitted `VectorAnyINTEL` for its non-standard vectors; after #3777 it emits `LongVectorEXT`. `VectorCompute` consumers require `VectorAnyINTEL`, so this is a functional regression. Hence, the enabled extension list does not identify a module as `VectorCompute` - the `VectorCompute` metadata (`VCFunction` /`VCGlobalVariable`) does. `IsVectorComputeModule` records the presence of that metadata before type translation. `VectorCompute` modules keep using `VectorAnyINTEL`, restoring the pre-#3777 behavior; all other modules still prefer the multi-vendor `LongVectorEXT` when its extension is enabled. --------- Co-authored-by: Ye, Qi <qi.ye@intel.com> Original commit: KhronosGroup/SPIRV-LLVM-Translator@3238c2903cd177e

dm-vodopyanov requested a review from a team as a code owner May 18, 2021 15:49

dm-vodopyanov requested a review from romanovvlad May 18, 2021 15:49

dm-vodopyanov added 2 commits May 18, 2021 18:50

Add missing EOL

dd66b44

Remove extra endif

22f9009

bader changed the title ~~[SYCL] Move __buildin_unique_stable_name dependent funcs to a separate header~~ [SYCL] Move __builtin_unique_stable_name dependent funcs to a separate header May 18, 2021

dm-vodopyanov changed the title ~~[SYCL] Move __builtin_unique_stable_name dependent funcs to a separate header~~ [SYCL] Move get_spec_constant_symbolic_ID to a separate header May 19, 2021

dm-vodopyanov added 3 commits May 20, 2021 20:03

Merge branch 'intel:sycl' into private/dvodopya/move-__buildin_unique…

3db201f

…_stable_name-to-separate-header

Merge branch 'intel:sycl' into private/dvodopya/move-__buildin_unique…

68187a5

…_stable_name-to-separate-header

Add -fsycl-use-footer to tests and include missing header to int footer

1bd1e94

dm-vodopyanov requested review from AaronBallman, elizabethandrews and premanandrao as code owners May 31, 2021 13:13

AlexeySachkov reviewed Jun 1, 2021

View reviewed changes

Comment thread clang/lib/Sema/SemaSYCL.cpp Outdated

elizabethandrews requested a review from erichkeane June 1, 2021 15:23

Enable __builtin_unique_stable_name in FE; Make -fsycl-use-footer ena…

1bd83e0

…bled by default for any app

dm-vodopyanov requested review from AGindinson and mdtoguchi as code owners June 1, 2021 17:33

mdtoguchi reviewed Jun 1, 2021

View reviewed changes

Comment thread clang/test/Driver/sycl-preprocess.cpp

dm-vodopyanov added 5 commits June 9, 2021 16:16

Merge branch 'sycl' into private/dvodopya/move-__buildin_unique_stabl…

0c9e8e2

…e_name-to-separate-header

Fix some test failures

0d3e000

Merge branch 'sycl' into private/dvodopya/move-__buildin_unique_stabl…

224eb1a

…e_name-to-separate-header

Update PR to use new builtin

ed1686d

Fix SYCL :: basic_tests/SYCL-2020*** tests

feb54c4

erichkeane reviewed Jun 15, 2021

View reviewed changes

Comment thread clang/test/CodeGenSYCL/anonymous_integration_footer.cpp Outdated

Comment thread clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp Outdated

AlexeySachkov added 2 commits June 16, 2021 16:17

Fix check-clang

e56ee41

Merge remote-tracking branch 'sycl' into private/dvodopya/move-__buil…

30ce3d0

…din_unique_stable_name-to-separate-header

AlexeySachkov dismissed stale reviews from mdtoguchi and AaronBallman via 3ef8887 June 25, 2021 13:27

Merge remote-tracking branch 'sycl' into private/dvodopya/move-__buil…

77e7988

…din_unique_stable_name-to-separate-header

AlexeySachkov removed the request for review from kbobrovs June 25, 2021 13:31

AlexeySachkov changed the title ~~[SYCL] Move get_spec_constant_symbolic_ID to a separate header~~ [SYCL] Switch to using integration footer by default Jun 25, 2021

AlexeySachkov added 4 commits June 28, 2021 10:10

Adjust sycl-offload-amdgcn.cpp test

68c1f1a

Merge remote-tracking branch 'sycl' into private/dvodopya/move-__buil…

30de609

…din_unique_stable_name-to-separate-header

Add a test for -fno-sycl-use-footer

37be540

Rename get_spec_constant_symbolic_ID

48e59c9

The primary function/wrapper which is used by DPC++ RT is left as-is, i.e. `get_spec_constant_symbolic_ID`. The internal function and its specialziations which are generated by the compiler renamed to `get_spec_constant_symbolic_ID_impl`.

romanovvlad reviewed Jun 28, 2021

View reviewed changes

Fix comment and clang-format

b73a756

elizabethandrews previously approved these changes Jun 28, 2021

View reviewed changes

mdtoguchi reviewed Jun 28, 2021

View reviewed changes

Comment thread clang/test/Driver/sycl-int-footer.cpp

AlexeySachkov added 3 commits June 29, 2021 12:38

Merge remote-tracking branch 'sycl' into private/dvodopya/move-__buil…

d5765f4

…din_unique_stable_name-to-separate-header

Add test for compiler phases w/ and w/o integration footer

8e58b19

Merge branch 'private/dvodopya/move-__buildin_unique_stable_name-to-s…

d5e10a2

…eparate-header' of github.com:dm-vodopyanov/llvm into private/dvodopya/move-__buildin_unique_stable_name-to-separate-header

AlexeySachkov dismissed elizabethandrews’s stale review via d5e10a2 June 29, 2021 11:17

An attempt to make sycl-int-footer.cpp OS-independent

9476403

AlexeySachkov requested review from elizabethandrews and mdtoguchi June 30, 2021 07:01

mdtoguchi approved these changes Jun 30, 2021

View reviewed changes

elizabethandrews approved these changes Jun 30, 2021

View reviewed changes

romanovvlad approved these changes Jun 30, 2021

View reviewed changes

romanovvlad merged commit 4a6d21f into intel:sycl Jun 30, 2021

dm-vodopyanov deleted the private/dvodopya/move-__buildin_unique_stable_name-to-separate-header branch February 10, 2022 14:06

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[SYCL] Switch to using integration footer by default#3777

[SYCL] Switch to using integration footer by default#3777
romanovvlad merged 34 commits into
intel:syclfrom
dm-vodopyanov:private/dvodopya/move-__buildin_unique_stable_name-to-separate-header

dm-vodopyanov commented May 18, 2021 •

edited by AlexeySachkov

Loading

Uh oh!

bader commented May 18, 2021

Uh oh!

romanovvlad commented May 19, 2021

Uh oh!

Uh oh!

erichkeane commented Jun 1, 2021

Uh oh!

Uh oh!

erichkeane left a comment

Uh oh!

Uh oh!

Uh oh!

romanovvlad left a comment

Uh oh!

AlexeySachkov commented Jun 28, 2021

Uh oh!

elizabethandrews left a comment

Uh oh!

Uh oh!

elizabethandrews left a comment

Uh oh!

romanovvlad left a comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

9 participants

Uh oh!

Conversation

dm-vodopyanov commented May 18, 2021 • edited by AlexeySachkov Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bader commented May 18, 2021

Uh oh!

romanovvlad commented May 19, 2021

Uh oh!

Uh oh!

erichkeane commented Jun 1, 2021

Uh oh!

Uh oh!

erichkeane left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

romanovvlad left a comment

Choose a reason for hiding this comment

Uh oh!

AlexeySachkov commented Jun 28, 2021

Uh oh!

elizabethandrews left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

elizabethandrews left a comment

Choose a reason for hiding this comment

Uh oh!

romanovvlad left a comment

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

9 participants

dm-vodopyanov commented May 18, 2021 •

edited by AlexeySachkov

Loading