[SYCL] Optional kernel features: implement split based on reqd-sub-group-size #8167
Conversation
…oup-size This patch implements device code split based on reqd-sub-group-size attribute, enables generation of "reqd_sub_group_size" property in "SYCL/device requirements" property set, and adds support of reqd_sub_group_size to sycl::is_compatible Design: https://github.com/intel/llvm/blob/sycl/sycl/doc/design/OptionalDeviceFeatures.md#changes-to-the-device-code-split-algorithm E2E tests: TBA
| ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 | ||
| ; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \ | ||
| ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 | ||
| ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M2-IR \ |
There was a problem hiding this comment.
It is ok to change this test, per-reqd-wg-size-split-2.ll, sycl-esimd-large-grf.ll and sycl-large-grf.ll below: by adding new optional kernel feature reqd-sub-group-size to internal data structures, re-hashing happened which affected name generation logic. This does not affect customers as customers don't use these temp files directly, and didn't break anything in the pipeline.
| } // namespace detail | ||
| } // __SYCL_INLINE_VER_NAMESPACE(_V1) | ||
| } // namespace sycl | ||
|
|
||
| namespace sycl { | ||
| __SYCL_INLINE_VER_NAMESPACE(_V1) { | ||
| namespace detail { |
There was a problem hiding this comment.
nit: Normally, I'd prefer those simplifications to be a separate PR.
|
@v-klochkov no-sycl-esimd-split-shared-func.ll test started to fail because this patch adds device code splitting for reqd_sub_group_size. Can you please advise, how to update this test? Or probably it doesn't make sense now? |
|
/verify with intel/llvm-test-suite#1569 |
|
/verify with intel/llvm-test-suite#1569 |
I am not sure yet if the test should stay with your fix, WITHOUT THE FIX: WITH THE FIX: b1) Not OK: SIMD_CALLEE is not removed and left as a dead code in |
Device code split based on reqd-sub-group-size is a part of optional kernel features which are part of SYCL 2020. This esimd test uses reqd-sub-group-size attribute, so this is expected that now more than one modules are generated because we do split on reqd-sub-group-size. So I guess test should be re-written or esimd feature should be revised to take into account this SYCL 2020 feature. E.g., regarding the test, removing |
The attribute "!sycl_explicit_simd" must have higher priority than "!intel_reqd_sub_group_size 1". It is Ok to separate funcs with "!intel_reqd_sub_group_size 1" to a separate module, The current version of PR generates wrong IR. The PR cloned SHARED_F 2 times instead of 1, which is unnecessary. Also, why modified compiler moved this func to test_esimd_2.ll: and this one to test_0.ll? |
) #### Intro This is a refactoring of how we perform device code split in `sycl-post-link`, which is intended to solve several existing issues with the current implementation: 1. increased peak RAM consumption by `sycl-post-link` 2. bad scaling with more and more split "dimensions" being added 3. increased tests maintenance cost due to non-deterministic order (between commits) of output files produced by `sycl-post-link` #### A bit more context about the issues above: (1) Increase peak RAM consumption is caused by the fact that we currently preserve **all** splits in-memory, even though we can process them on-by-one and discard them as soon as we stored them to a disk. This was implemented as a memory consumption optimization in #5021, but it got accidentally reverted in #7302 as an attempt to workaround (2). (2) is pretty much summarized in our source code: https://github.com/intel/llvm/blob/afebb2543ccecb89f83c84b68fba7616bbab89ac/llvm/tools/sycl-post-link/sycl-post-link.cpp#L806-L811 (3) is caused by a bad implementation decision made in #7302: because every split is now identified by a hash, every time you add a new split "dimension"/new feature to an account, it results in different hashes for existing tests. Just look how many unrelated tests had to be updated in #7512, #8056 and #8167 #### Now to the PR itself: It introduces a new infrastructure for categorizing/grouping kernel functions: instead of using hashes, we now build a string description for each kernel function and then group kernels with the same description string together. String description is built by a new entity: it accepts a set of rules, where each rule is a simple function which returns a string for passed `llvm::Function`. Results of all rules are concatenated together and rules are invoked in a stable order of their registration. There is a simple API for building those rules. It provides some predefined rules for the most popular use cases like turning a function attribute or a metadata into a string descriptor for the function. There is also a possibility to pass a custom callback there to implement more complicated logic. #### How does this PR help with issues above? (1) and (2) are fixed in conjunction: `sycl-post-link` was refactored to avoid storing more than one split module at a time and that is possible because the PR unifies per-scope and optional-kernel-features splitters into a single generic splitter. The new API for kernels categorization seems to be flexible enough to provide that infrastructure so merged splitters still look OK code-wise. (3) is caused by using string identifiers instead of hashes as well as by using a data structure which sorts identifiers. #### Any other benefits from this PR? About 50 lines of code less to support :) Extending device code split for more optional features would be even easier than it is now: instead of adding several changes to various places around `UsedOptionalFeatures` structure, it will be just adding a 1-3 lines of code. Please also note that `UsedOptionalFeatures` contains tons of inconsistencies in its implementation, which will all gone with this PR: in `operator==` we don't use hash and instead compare certain fields directly (and we do miss some of them); `generateModuleName` method skips some of optional features and ignores them. Cross-module `device_global` usages checks should now work at all split dimensions (except for ESIMD). #### Any potential downsides? With current `UsedOptionalFeatures` there is a possibility to embed various information (used aspects, `large-grf` flag, etc.) directly during device code split to avoid re-gathering that information later when we generate properties. With the suggested approach, it would be harder to do, because it doesn't seem to naturally fit to the proposed infrastructure: see changes I did around `large-grf` in this PR. However, we have never actually implemented this and re-querying some metadata from function doesn't seem like a bottleneck, so it should really be a very minor and only theoretical downside.
|
Superseded by #9928 |
This patch implements device code split based on reqd-sub-group-size attribute, enables generation of "reqd_sub_group_size" property in "SYCL/device requirements" property set, and adds support of reqd_sub_group_size to sycl::is_compatible
Design:
https://github.com/intel/llvm/blob/sycl/sycl/doc/design/OptionalDeviceFeatures.md#changes-to-the-device-code-split-algorithm
E2E tests: intel/llvm-test-suite#1569