[SYCL][Bindless] Update and add support for SPV_INTEL_bindless_image extension new revision #13753

DBDuncan · 2024-05-10T16:58:52Z

Add support to emit instructions that convert handles to images, samplers and sampled images

… extension new revision Add support to emit instructions that convert handles to images, samplers and sampled images

DBDuncan · 2024-05-10T17:00:44Z

There will be a corresponding SPIRV-LLVM-Translator pull request soon.

DBDuncan · 2024-05-13T12:07:00Z

There will be a corresponding SPIRV-LLVM-Translator pull request soon.

Now up here: KhronosGroup/SPIRV-LLVM-Translator#2559

Needs to be merged first and then eventually pulled into DPC++. Then two tests in this PR should pass.

aelovikov-intel · 2024-05-13T14:43:49Z

sycl/include/sycl/ext/oneapi/bindless_images.hpp

+// Image types used for generating SPIR-V
+#ifdef __SYCL_DEVICE_ONLY__
+template <int NDims>
+using OCLImageTyRead =


Is this part of the specification or should it be moved into the detail namespace?

Ah, yeah it should be in the detail namespace. Moved it.

sycl/include/sycl/ext/oneapi/bindless_images.hpp

aelovikov-intel · 2024-05-13T14:47:10Z

sycl/include/sycl/ext/oneapi/bindless_images.hpp

-    return __invoke__ImageReadLod<DataT>(imageHandle.raw_handle, coords, level);
+    return __invoke__ImageReadLod<DataT>(
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle,
+                                        OCLImageTyRead<coordSize>),


Why wouldn't you put OCLImageTyRead inside CONVERT_HANDLE_TO_SAMPLED_IMAGE? Or better yet inside __invoke__ImageReadLod...

I do agree that putting OCLImageTyRead inside CONVERT_HANDLE_TO_SAMPLED_IMAGE is better. But I am not sure about putting it inside __invoke__ImageReadLod. That would require also putting sampled_opencl_image_type inside __invoke__ImageReadLod

Because these functions are used to also create the PTX instructions for the CUDA backend I feel putting it inside __invoke__ImageReadLod would complicate things. What do you think?

I have moved OCLImageTyRead into CONVERT_HANDLE_TO_SAMPLED_IMAGE for now unless something different is decided.

Technically, this file is outside SYCL RT codeownership, so the decision is ultimately yours :)

I will leave it as is for now. Potentially revisit this later.

sycl/include/sycl/ext/oneapi/bindless_images.hpp

wenju-he · 2024-05-14T01:14:58Z

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -230,6 +230,18 @@ template <typename SampledType, typename TempRetT, typename TempArgT>
 extern __DPCPP_SYCL_EXTERNAL TempRetT __spirv_ImageSampleCubemap(SampledType,
                                                                 TempArgT);

+template <typename RetT>
+extern __DPCPP_SYCL_EXTERNAL RetT
+__spirv_ConvertHandleToImageINTEL(unsigned long);


unsigned long is 32bit on windows?

can we use following code?

template <class RetT, class HandleT> extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ConvertHandleToImageINTEL(HandleT);

Good point. Added.

…side 'CONVERT_HANDLE_TO_SAMPLED_IMAGE' as sampled images only ever use read image type

sycl/include/sycl/ext/oneapi/bindless_images.hpp

aelovikov-intel · 2024-05-14T14:27:56Z

sycl/include/sycl/ext/oneapi/bindless_images.hpp

+#if defined(__NVPTX__)
+#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords)                       \
+  __invoke__ImageFetch<DataT>(raw_handle, coords)
+#elif defined(__SPIR__)
+#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords)                       \
+  __invoke__ImageRead<DataT>(raw_handle, coords)
+#else
+#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords)                       \
+  __invoke__ImageFetch<DataT>(raw_handle, coords)
+#endif


Are these all the same?

Not quite. In the CUDA backend, __invoke__ImageFetch is used for a basic unsampled image read. When this is encountered by llvm-spirv, it crashes as it tries to emit, OpImageFetch, which is not implemented in llvm-spirv. But for a basic unsampled image read, we want to emit OpImageRead instead. So __invoke__ImageRead is used when compiling to SPIR-V.

Interestingly, and slightly annoyingly, OpImageFetch must always return a vector of four components. But that is prob not a big deal and is something to think about much later when adding support for fetching data from sampled images.

Also, turns out, llvm-spirv does not have an assert to check if an instruction is fully implemented before trying to call Inst->init(); at line 185 of SPIRVInstruction.h and jumping to random memory. That was a bit annoying.

I have removed one of the redundant branches.

sycl/include/sycl/ext/oneapi/bindless_images.hpp

mdtoguchi · 2024-05-14T16:40:59Z

clang/lib/Driver/ToolChains/Clang.cpp

@@ -10394,7 +10394,8 @@ static void getOtherSPIRVTransOpts(Compilation &C,
      ",+SPV_INTEL_fpga_argument_interfaces"
      ",+SPV_INTEL_fpga_invocation_pipelining_attributes"
      ",+SPV_INTEL_fpga_latency_control"
-      ",+SPV_INTEL_task_sequence";
+      ",+SPV_INTEL_task_sequence"
+      ",+SPV_INTEL_bindless_images";


Could you also update clang/test/Driver/sycl-spirv-ext.c to cover your new entry?

Oh, I missed that. Will do.

DBDuncan · 2024-05-22T12:40:12Z

Friendly ping @aelovikov-intel

aelovikov-intel

Comments about tests could be resolved in a different PR.

aelovikov-intel · 2024-05-22T14:49:23Z

sycl/include/sycl/ext/oneapi/bindless_images.hpp

+
+// Macros are required because it is not legal for a function to return
+// a variable of type 'opencl_image_type'.
+#if defined(__SPIR__)


nit: I'd personally merge the #ifs into a single one providing three macros at once. Can be ignored.

Yeah, actually after combining them they did look better. Done.

aelovikov-intel · 2024-05-22T14:53:33Z

sycl/test/extensions/bindless_images_SPIRV_inst.cpp

+int main() {
+
+  sycl::device dev;
+  sycl::queue q(dev);
+  auto ctxt = q.get_context();


I don't think we need any of this for this test. One can use SYCL_EXTERNAL and only run device compilation to produce much compact IR. See https://github.com/intel/llvm/blob/sycl/sycl/test/check_device_code/vector/vector_as.cpp for an example.

aelovikov-intel · 2024-05-22T14:54:23Z

sycl/test/extensions/bindless_images_SPIRV_inst.cpp

+  try {
+    q.submit([&](sycl::handler &cgh) {
+      cgh.parallel_for<image_addition>(1, [=](sycl::id<1> id) {
+#ifdef __SYCL_DEVICE_ONLY__


That can become troublesome if customers have to structure their code like this. In the worst case it can lead to different lambda captures between host/device.

Customers do not need to structure their code like this. I was just intending to invoke the new __spirv_ builtins manually to verify they worked. The guards are needed as opencl image types do not compile on host.

There is the other test, sycl/test/extensions/bindless_images.cpp which does also check that the SPIR-V instructions are emitted correctly, but it does not check that __spirv_ConvertHandleToSamplerINTEL works correctly as it is currently not used in the SYCL bindless extension. So I wanted to keep this to ensure all of the new builtins were covered.

DBDuncan · 2024-05-22T17:21:43Z

Comments about tests could be resolved in a different PR.

Good idea. I will double check with my team. Should be fine.

DBDuncan · 2024-05-23T13:15:34Z

Comments about tests could be resolved in a different PR.

Good idea. I will double check with my team. Should be fine.

@aelovikov-intel On your comment that I can use SYCL_EXTERNAL like in vector_as.cpp, we think it is best to leave it as is as we don't think it is worth it at the moment to update to use SYCL_EXTERNAL and all the filecheck checks.

If ok for you, can you approve and I will move forward to getting this merged.

The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753 (cherry picked from commit e9a9a3f)

DBDuncan · 2024-05-24T11:48:59Z

@intel/llvm-gatekeepers Should be ready to merge. There are some failing tests but all of them look unrelated to the changes here.

steffenlarsen · 2024-05-24T12:36:22Z

The issue causing no_sycl_hpp_in_e2e_tests.cpp have been addressed by #13889.

…ion (KhronosGroup#2559) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753 Original PR KhronosGroup#2559

…ion (KhronosGroup#2559) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

…ion (#2559) (#3002) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

…ion (#2559) (#3008) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

…#2559) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

…on (KhronosGroup#2559) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

…ion (#2559) (#3009) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

…#2559) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

…on (KhronosGroup#2559) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

…ion (#2559) (#3016) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

[SYCL][Bindless] Update and add support for SPV_INTEL_bindless_images…

dd48b8e

… extension new revision Add support to emit instructions that convert handles to images, samplers and sampled images

DBDuncan mentioned this pull request May 10, 2024

Add updated SPV_INTEL_bindless_images preview extension KhronosGroup/SPIRV-LLVM-Translator#2559

Merged

DBDuncan had a problem deploying to WindowsCILock May 10, 2024 17:13 — with GitHub Actions Failure

DBDuncan temporarily deployed to WindowsCILock May 10, 2024 18:17 — with GitHub Actions Inactive

DBDuncan marked this pull request as ready for review May 13, 2024 12:31

DBDuncan requested review from a team as code owners May 13, 2024 12:31

DBDuncan requested a review from aelovikov-intel May 13, 2024 12:31

merge branch 'origin/sycl' into duncan/spirv-bindless-images-patch2

7fd7b37

DBDuncan had a problem deploying to WindowsCILock May 13, 2024 12:58 — with GitHub Actions Failure

DBDuncan temporarily deployed to WindowsCILock May 13, 2024 14:28 — with GitHub Actions Inactive

aelovikov-intel reviewed May 13, 2024

View reviewed changes

wenju-he reviewed May 14, 2024

View reviewed changes

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

wenju-he reviewed May 14, 2024

View reviewed changes

Move opencl types to detail namespace and move 'OCLImageTyRead' to in…

7fd6ce0

…side 'CONVERT_HANDLE_TO_SAMPLED_IMAGE' as sampled images only ever use read image type

DBDuncan had a problem deploying to WindowsCILock May 14, 2024 14:22 — with GitHub Actions Failure

aelovikov-intel reviewed May 14, 2024

View reviewed changes

DBDuncan temporarily deployed to WindowsCILock May 14, 2024 15:09 — with GitHub Actions Inactive

mdtoguchi reviewed May 14, 2024

View reviewed changes

DBDuncan added 2 commits May 14, 2024 17:46

merge macro branches

9fb5c7d

Template '__spirv_ConvertHandleTo*' builtin handle type

f674ba9

DBDuncan had a problem deploying to WindowsCILock May 14, 2024 16:48 — with GitHub Actions Failure

DBDuncan temporarily deployed to WindowsCILock May 14, 2024 17:28 — with GitHub Actions Inactive

Add 'SPV_INTEL_bindless_images' to 'sycl-spirv-ext.c'

b564671

DBDuncan had a problem deploying to WindowsCILock May 15, 2024 12:41 — with GitHub Actions Failure

DBDuncan temporarily deployed to WindowsCILock May 15, 2024 13:24 — with GitHub Actions Inactive

aelovikov-intel approved these changes May 22, 2024

View reviewed changes

Combine '#ifs' into single '#if'

f8f391f

DBDuncan temporarily deployed to WindowsCILock May 23, 2024 14:10 — with GitHub Actions Inactive

DBDuncan temporarily deployed to WindowsCILock May 23, 2024 14:51 — with GitHub Actions Inactive

aelovikov-intel approved these changes May 23, 2024

View reviewed changes

steffenlarsen merged commit 83db85f into intel:sycl May 24, 2024
11 of 14 checks passed

ProGTX mentioned this pull request Jul 18, 2024

[SYCL][SPV_INTEL_bindless_images] Add OpConvertHandleToSampledImageINTEL back to the extension #13392

Closed

KanclerzPiotr mentioned this pull request Feb 7, 2025

[Backport to 15] Add updated SPV_INTEL_bindless_images preview extension (#2559) KhronosGroup/SPIRV-LLVM-Translator#3002

Merged

KanclerzPiotr mentioned this pull request Feb 10, 2025

[Backport to 16] Add updated SPV_INTEL_bindless_images preview extension (#2559) KhronosGroup/SPIRV-LLVM-Translator#3008

Merged

KanclerzPiotr pushed a commit to KanclerzPiotr/SPIRV-LLVM-Translator that referenced this pull request Feb 11, 2025

Add updated SPV_INTEL_bindless_images preview extension (KhronosGroup…

2628594

…#2559) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

KanclerzPiotr mentioned this pull request Feb 11, 2025

[Backport to 18] Add updated SPV_INTEL_bindless_images preview extension (#2559) KhronosGroup/SPIRV-LLVM-Translator#3009

Merged

This was referenced Feb 12, 2025

[Backport to 15] Add updated SPV_INTEL_bindless_images preview extension (#2559) KhronosGroup/SPIRV-LLVM-Translator#3012

Closed

[Backport to 17] Add updated SPV_INTEL_bindless_images preview extension (#2559) KhronosGroup/SPIRV-LLVM-Translator#3016

Merged

KanclerzPiotr pushed a commit to KanclerzPiotr/SPIRV-LLVM-Translator that referenced this pull request Feb 13, 2025

Add updated SPV_INTEL_bindless_images preview extension (KhronosGroup…

c974cea

…#2559) The updated SPV_INTEL_bindless_images spec can be seen in this PR: intel/llvm#13753

[SYCL][Bindless] Update and add support for SPV_INTEL_bindless_image extension new revision #13753

[SYCL][Bindless] Update and add support for SPV_INTEL_bindless_image extension new revision #13753

Uh oh!

Conversation

DBDuncan commented May 10, 2024

Uh oh!

DBDuncan commented May 10, 2024

Uh oh!

DBDuncan commented May 13, 2024

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

DBDuncan commented May 22, 2024

Uh oh!

aelovikov-intel left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

DBDuncan commented May 22, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

DBDuncan commented May 23, 2024

Uh oh!

DBDuncan commented May 24, 2024

Uh oh!

steffenlarsen commented May 24, 2024

Uh oh!

Uh oh!

Uh oh!

DBDuncan commented May 22, 2024 •

edited

Loading