[SYCL] Implement SYCL part of sycl_ext_oneapi_prefetch #11458

KornevNikita · 2023-10-06T16:15:02Z

Spec: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_prefetch.asciidoc

The joint_prefetch functions will be done is a separate patch.

This implementation also requires a patch for llvm-spirv translator. SPIRVWriter should handle these new annotations and create the appropriate decorations in spv representation.

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

sycl/include/sycl/ext/oneapi/prefetch.hpp

sycl/include/sycl/ext/oneapi/properties/properties.hpp

aelovikov-intel · 2023-10-06T16:55:14Z

sycl/test/extensions/prefetch.cpp

+
+char data[] = {0, 1, 2, 3};
+
+// CHECK: [[PREFETCH_STR:@.*]] = private unnamed_addr addrspace(1) constant [19 x i8] c"sycl-prefetch-hint\00", section "llvm.metadata"


Are these manually written? If so, do we have an utility script to generate them automatically? If yes, one might need to disable instrumentation/use -O1 to get a more readable IR.

aelovikov-intel · 2023-10-06T16:56:32Z

sycl/test/extensions/prefetch.cpp

+  namespace syclex = sycl::ext::oneapi::experimental;
+  sycl::queue q;
+  void *dataPtr = &data;
+  q.parallel_for(1, [=](sycl::id<1> idx) {


I think this can be single_task. I'd also want to see an E2E test with this used in non-uniform control flow (I don't think the spec prohibits that).

Moved to single task - 7afe1c9. I'm going to add E2E tests a bit later when the whole feature (incl. llvm-spirv translator part) is done. Not sure what do you mean by "non-uniform control". Could you please explain a bit?

if (id.get_global_id(0) % 3 == 0)) syclex::prefetch(p);

Now that you've added joint_prefetch that just delegates to the per-WI one, I'm even more concerned about non-uniform control flow scenario...

aelovikov-intel · 2023-10-09T15:49:46Z

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

-    auto DecorIt = SpirvDecorMap.find(*Property.first);
+    // Leave these annotations as is. They will be processed by SPIRVWriter.
+    if (first == "sycl-prefetch-hint" || first == "sycl-prefetch-hint-nt") {
+      return false;


This doesn't feel right to me, but it's outside of SYCL RT...

Looks like currently it's the only way to transform these properties into spirv decorations for this pointer. I tried to create spirv.Decorations metadata instead, but the compiler eliminates them with optimization flags.
Annotation also may be eliminated, but it's much less likely.

sycl/include/sycl/ext/oneapi/prefetch.hpp

aelovikov-intel · 2023-10-10T15:57:56Z

sycl/include/sycl/ext/oneapi/prefetch.hpp

+prefetch(
+    accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
+    size_t offset, size_t count, Properties properties = {}) {
+  prefetch((void *)&acc[offset], count * sizeof(DataT), properties);


Do we know for sure that count elements are consecutive in memory?

Good point. @Pennycook I'm not sure how it's intended to work in case of N-dim offset. Should we call the __spirv_ocl_prefetch spirv instruction several times for different memory segments in such case or there should be some constraint?

We struggled a bit with multi-dimensional prefetches (see the issues).

Where we landed (for now) is that the block being prefetched is assumed to be contiguous, and only the offset itself is multi-dimensional. It's effectively a shorthand to avoid computing the linear offset from the start of the buffer. Note that the specification says for the multi-dimensional cases:

Effects: Equivalent to prefetch((void*) &acc[offset], sizeof(DataT), properties).
Effects: Equivalent to prefetch((void*) &acc[offset], count, properties).

If somebody requests a multi-dimensional prefetch later, we can describe it with a range parameter in place of a size_t count, and implement it the way you suggested (by calling the instruction multiple times).

bader · 2023-10-12T14:29:38Z

@intel/dpcpp-tools-reviewers, please, take a look. The compiler change is non-functional, so should be a no brainer.

sycl/include/sycl/ext/oneapi/prefetch.hpp

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

asudarsa

SYCLLowerIR changes look good to me, as they are NFC. Just a quick request about variable name. Also, I just noticed there is no test added.

Thanks

AlexeySachkov · 2023-10-16T13:10:47Z

Pre-commit failure doesn't seem related, it is covered by #11549

The first part: intel#11458 Adjust the CompileTimePropertiesPass so it can convert new properties into spirv decorations.

SYCL part: intel/llvm#11458 intel/llvm#11597 Handle new properties and decorate prefetch's arg.

…11597) The first part: #11458 Adjust the CompileTimePropertiesPass so it can convert new properties into spirv decorations.

SYCL part: #11458 #11597 Handle new properties and decorate prefetch's arg. Original commit: KhronosGroup/SPIRV-LLVM-Translator@a76f24e

Initial

6ab79b1

KornevNikita requested review from a team as code owners October 6, 2023 16:15

KornevNikita requested a review from aelovikov-intel October 6, 2023 16:15

aelovikov-intel reviewed Oct 6, 2023

View reviewed changes

KornevNikita added 10 commits October 9, 2023 01:43

clang format

c4123e9

Add joint_prefetch

549f9e8

Fix comment

19e19a8

Use structured bindings

30e531a

Remove redundant is_valid_property

42de79f

Add explicit enum values

2d5a360

Ignore unused vars on host

c40e4bd

Use enable_if as ret type

5de267d

Use single_task in test

7afe1c9

clang-format

886f2ab

KornevNikita temporarily deployed to WindowsCILock October 9, 2023 14:44 — with GitHub Actions Inactive

KornevNikita temporarily deployed to WindowsCILock October 9, 2023 15:33 — with GitHub Actions Inactive

aelovikov-intel reviewed Oct 9, 2023

View reviewed changes

KornevNikita added 5 commits October 10, 2023 02:01

Rename first and second

a8e907b

Fix warning

01a8b5a

Delete comments

154660a

Fix

82badcc

Merge remote-tracking branch 'intel_llvm/sycl' into prefetch

7ecdbb1

KornevNikita temporarily deployed to WindowsCILock October 10, 2023 09:16 — with GitHub Actions Inactive

KornevNikita temporarily deployed to WindowsCILock October 10, 2023 10:10 — with GitHub Actions Inactive

Add helper

ebc3b23

KornevNikita temporarily deployed to WindowsCILock October 10, 2023 15:01 — with GitHub Actions Inactive

KornevNikita temporarily deployed to WindowsCILock October 10, 2023 15:48 — with GitHub Actions Inactive

aelovikov-intel reviewed Oct 10, 2023

View reviewed changes

sycl/include/sycl/ext/oneapi/prefetch.hpp Outdated Show resolved Hide resolved

aelovikov-intel reviewed Oct 10, 2023

View reviewed changes

KornevNikita temporarily deployed to WindowsCILock October 11, 2023 08:49 — with GitHub Actions Inactive

KornevNikita temporarily deployed to WindowsCILock October 11, 2023 09:32 — with GitHub Actions Inactive

aelovikov-intel approved these changes Oct 11, 2023

View reviewed changes

Remove sycl-post-link & test in favor on a follow-up patch

024c901

KornevNikita temporarily deployed to WindowsCILock October 12, 2023 12:18 — with GitHub Actions Inactive

KornevNikita temporarily deployed to WindowsCILock October 12, 2023 13:11 — with GitHub Actions Inactive

bader reviewed Oct 12, 2023

View reviewed changes

sycl/include/sycl/ext/oneapi/prefetch.hpp Outdated Show resolved Hide resolved

sarnex reviewed Oct 12, 2023

View reviewed changes

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp Outdated Show resolved Hide resolved

bader changed the title ~~[SYCL] Implement sycl part of sycl_ext_oneapi_prefetch~~ [SYCL] Implement SYCL part of sycl_ext_oneapi_prefetch Oct 12, 2023

asudarsa reviewed Oct 12, 2023

View reviewed changes

KornevNikita added 4 commits October 13, 2023 05:56

use capital initial letters

c796b9e

Move prefetch.hpp to experimental/ dir

bf3d6c9

Improve impl

6d9de5a

Add new test

9ff283e

AlexeySachkov approved these changes Oct 13, 2023

View reviewed changes

KornevNikita had a problem deploying to WindowsCILock October 13, 2023 15:17 — with GitHub Actions Failure

KornevNikita temporarily deployed to WindowsCILock October 13, 2023 16:07 — with GitHub Actions Inactive

small fix

9b2ecc8

KornevNikita temporarily deployed to WindowsCILock October 13, 2023 16:24 — with GitHub Actions Inactive

KornevNikita temporarily deployed to WindowsCILock October 13, 2023 17:18 — with GitHub Actions Inactive

AlexeySachkov merged commit e7139b0 into intel:sycl Oct 16, 2023

KornevNikita mentioned this pull request Oct 19, 2023

[SYCL] Adjust CompileTimePropertiesPass for sycl_ext_oneapi_prefetch #11597

Merged

KornevNikita added a commit to KornevNikita/SPIRV-LLVM-Translator that referenced this pull request Oct 19, 2023

Support sycl_ext_oneapi_prefetch

7358d3c

SYCL part: intel/llvm#11458 intel/llvm#11597 Handle new properties and decorate prefetch's arg.

KornevNikita mentioned this pull request Oct 19, 2023

Support sycl_ext_oneapi_prefetch KhronosGroup/SPIRV-LLVM-Translator#2189

Merged

MrSidims pushed a commit to KhronosGroup/SPIRV-LLVM-Translator that referenced this pull request Oct 26, 2023

Support sycl_ext_oneapi_prefetch (#2189)

a76f24e

SYCL part: intel/llvm#11458 intel/llvm#11597 Handle new properties and decorate prefetch's arg.

againull pushed a commit that referenced this pull request Oct 27, 2023

[SYCL] Adjust CompileTimePropertiesPass for sycl_ext_oneapi_prefetch (#…

0229456

…11597) The first part: #11458 Adjust the CompileTimePropertiesPass so it can convert new properties into spirv decorations.

jsji pushed a commit that referenced this pull request Nov 2, 2023

Support sycl_ext_oneapi_prefetch (#2189)

b359eca

SYCL part: #11458 #11597 Handle new properties and decorate prefetch's arg. Original commit: KhronosGroup/SPIRV-LLVM-Translator@a76f24e

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL] Implement SYCL part of sycl_ext_oneapi_prefetch #11458

[SYCL] Implement SYCL part of sycl_ext_oneapi_prefetch #11458

KornevNikita commented Oct 6, 2023

aelovikov-intel Oct 6, 2023

aelovikov-intel Oct 6, 2023

KornevNikita Oct 9, 2023

aelovikov-intel Oct 9, 2023

aelovikov-intel Oct 10, 2023

aelovikov-intel Oct 9, 2023

KornevNikita Oct 10, 2023

aelovikov-intel Oct 10, 2023

KornevNikita Oct 10, 2023

Pennycook Oct 10, 2023

bader commented Oct 12, 2023

asudarsa left a comment

AlexeySachkov commented Oct 16, 2023


		char data[] = {0, 1, 2, 3};

		// CHECK: [[PREFETCH_STR:@.*]] = private unnamed_addr addrspace(1) constant [19 x i8] c"sycl-prefetch-hint\00", section "llvm.metadata"

[SYCL] Implement SYCL part of sycl_ext_oneapi_prefetch #11458

[SYCL] Implement SYCL part of sycl_ext_oneapi_prefetch #11458

Conversation

KornevNikita commented Oct 6, 2023

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

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

bader commented Oct 12, 2023

asudarsa left a comment

Choose a reason for hiding this comment

AlexeySachkov commented Oct 16, 2023