[SPIR][SimplifyCFG] Do not generate select insructions with sampler type

[Fznamznon]

In SPIR-V neither image nor sampler types cannot go through phi and
select instructions. Lowering of phi nodes and select instructions with
image/sampler result type by SPIR-V translator tool is very complicated
in common case, so the transformation is disabled for image/sampler
types in the pass which does it.

[Fznamznon]

/summary:run

[AlexeySachkov]

I was wondering if that is something which we can upstream to llvm/llvm-project instead of having here as a customization. However, I do understand that neither OpenCL nor spir specs require that, so it could be hard to motivate. Tagging @AnastasiaStulova here to get her opinion on that.

[bader]

In SPIR-V neither image nor sampler types cannot go through phi and
select instructions.

However, I do understand that neither OpenCL nor spir specs require that, so it could be hard to motivate.

I think it make sense to check with if such difference in OpenCL and SPIR-V specifications is intended. @bashbaug, @mkinsner, do know the motivation to put an addition restriction on SPIR-V format for image/sampler types?

Alternative solution would be setting simplifycfg-sink-common option to false in the driver when we compile for spir target (it impacts non -image/-sampler types as well).

[bashbaug]

Interesting, we were discussing a similar issue internally a few weeks ago and I believe any differences are unintentional.

The OpenCL C spec has a bunch of restrictions around images (see item (b) specifically):
https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#restrictions

An image type (image2d_t, image3d_t, image2d_array_t, image1d_t, image1d_buffer_t or image1d_array_t) can only be used as the type of a function argument. An image function argument cannot be modified. Elements of an image can only be accessed using the built-in image read and write functions. [...]

I think these will also disallow using images and samplers for a select or phi.

For the old SPIR spec, we have:

i32 values that represent sampler t objects, can only be passed as arguments to images built-ins. Any other operation involving these i32 values is implementation defined.

So, samplers are restricted in SPIR also. I didn't find similar text for images but I suspect this was just an omission.

[gmlueck]

Will this break conditional selection of an image accessor in SYCL? Currently, nothing in the SYCL spec disallows code like this:

using my_image_accessor = sycl::sampled_image_accessor<sycl::int4, 2>;

void my_device_function(my_image_accessor acc1, my_image_accessor acc2, bool condition) {
  my_image_accessor acc = acc1;
  if (condition) {
    acc = acc2;
  }
  sycl::int4 val = acc.read({1.0, 2.0});
  // ...
}

[Fznamznon]

Will this break conditional selection of an image accessor in SYCL? Currently, nothing in the SYCL spec disallows code like this:

No, It shouldn't break this. Such SYCL will be represented in the IR with another set of instructions that doesn't use phi and select with image/sampler result, but it still be valid and functional code.

EDIT: I looked at the code snippet a bit more carefully, and I think this is yet another case where SYCL is less restrictive than OpenCL (another example I can remember is that sampler/image struct fields are disallowed in OpenCL and basically our image accessor or SYCL sampler class just break this restriction). Unoptimized, this code is invalid in OpenCL just because image accessors contain field of image opaque type inside, in case some optimizations are enabled - this will likely also break the restriction that we are discussing here, i.e. images can be used only as arguments to builtins. This is not connected with this patch though, since with it or without, such valid SYCL code cannot be valid in OpenCL, I think.

[Fznamznon]

Alternative solution would be setting simplifycfg-sink-common option to false in the driver when we compile for spir target (it impacts non -image/-sampler types as well).

I think it may negatively affect performance for non -image/-sampler types, so I changed behavior only for restricted types.

An image type (image2d_t, image3d_t, image2d_array_t, image1d_t, image1d_buffer_t or image1d_array_t) can only be used as the type of a function argument. An image function argument cannot be modified. Elements of an image can only be accessed using the built-in image read and write functions. [...]

I think these will also disallow using images and samplers for a select or phi.

I wasn't sure about this. The original code I was looking at was valid from this point of view, the thing is that it somehow was transformed to code that is explicitly disallowed by SPIR-V spec looked for me like some kind of implementation detail which is out of scope of OpenCL spec.

For the old SPIR spec, we have:
i32 values that represent sampler t objects, can only be passed as arguments to images built-ins. Any other operation involving these i32 values is implementation defined.
So, samplers are restricted in SPIR also. I didn't find similar text for images but I suspect this was just an omission.

The phrase "Any other operation involving these i32 values is implementation defined." looks a bit suspicious here. "Implementation defined" doesn't sound a bit different from "disallowed".

I case I'm overreacting, should I prepare review for this patch in LLORG then?

[bader]

Will this break conditional selection of an image accessor in SYCL? Currently, nothing in the SYCL spec disallows code like this:

No, It shouldn't break this. Such SYCL will be represented in the IR with another set of instructions that doesn't use phi and select with image/sampler result, but it still be valid and functional code.

EDIT: I looked at the code snippet a bit more carefully, and I think this is yet another case where SYCL is less restrictive than OpenCL (another example I can remember is that sampler/image struct fields are disallowed in OpenCL and basically our image accessor or SYCL sampler class just break this restriction). Unoptimized, this code is invalid in OpenCL just because image accessors contain field of image opaque type inside, in case some optimizations are enabled - this will likely also break the restriction that we are discussing here, i.e. images can be used only as arguments to builtins. This is not connected with this patch though, since with it or without, such valid SYCL code cannot be valid in OpenCL, I think.

@intel/dpcpp-specification-reviewers, do you have any ideas how image/sampler restrictions can be implemented for SYCL image/sampler classes? I don't see an easy way for C++ classes to implement such restrictions and diagnose "invalid" image/sampler usages at compile time.

[gmlueck]

@intel/dpcpp-specification-reviewers, do you have any ideas how image/sampler restrictions can be implemented for SYCL image/sampler classes? I don't see an easy way for C++ classes to implement such restrictions and diagnose "invalid" image/sampler usages at compile time.

As @bashbaug mentioned, we've been talking about this a little, though we have no resolution currently. One idea was that we could change the definition of image_accessor to contain a pointer to a SPIR-V OpTypeImage and then somehow cause the kernel to store the OpTypeImage value to memory at the beginning of the kernel. Operations on image_accessor would then no longer be on OpTypeImage values, so we avoid the restrictions. However, @bashbaug wasn't sure if it was legal to store the OpTypeImage value to memory.

Unoptimized, this code is invalid in OpenCL just because image accessors contain field of image opaque type inside, in case some optimizations are enabled - this will likely also break

This seems like something that must be fixed. The only way to pass an "image" opaque type into a SYCL kernel is via an accessor. Perhaps we need some special handling for this argument type which decomposes the image_accessor parameter and passes the underlying "image" type directly as an argument (rather than as a structure field).

[Fznamznon]

Perhaps we need some special handling for this argument type which decomposes the image_accessor parameter and passes the underlying "image" type directly as an argument (rather than as a structure field).

This already happens for kernel arguments. I'm not sure that changing of user's functions signature is a good idea though.

[gmlueck]

This already happens for kernel arguments. I'm not sure that changing of user's functions signature is a good idea though.

Are you proposing that it should be legal to capture an image accessor as a kernel argument, but it should not be legal to pass an image accessor as a parameter to an internal function inside the kernel? For example:

void foo(sycl::queue q) {
  q.submit([&](sycl::handler &cgh) {
    sycl::sampled_image_accessor<sycl::int4, 2> acc = /*...*/;
    cgh.parallel_for([=](sycl::item it) {
      sycl::int4 val = acc.read({1.0, 2.0});  // This is OK
      my_device_func(acc);                    // This is not allowed
    });
  });
}

void my_device_func(sycl::sampled_image_accessor<sycl::int4, 2> acc) {
  sycl::int4 val = acc.read({1.0, 2.0});
}

I think that would be a huge limitation to the language. Essentially, it means that kernels operating on images have to be written as one single function. That seems unrealistically limiting.

[Fznamznon]

Are you proposing that it should be legal to capture an image accessor as a kernel argument, but it should not be legal to pass an image accessor as a parameter to an internal function inside the kernel?

No, I agree that this is too restrictive. I personally think that some of OpenCL rules are too restrictive as well. Probably they were designed for old devices or smth like that, because in SYCL we break some of OpenCL rules once we started supporting image accessors and sampler objects and everything still works.
Also, now with enabled early optimizations I think most of functions that are not used from another translation units (or even only another shared library in case we enable some link time inlining) will be inlined allowing to optimize out image accessor types. That will bring some requirement for SYCL_EXTERNAL macro similar to what we have already about generic address space support. Relying on optimizations still concerns me a bit, since there is always possibility that some optimization won't work because of some heuristics that decides whether apply the change or not, so we would have to make them really aggressive. So, I'd still vote for OpenCL restrictions lifting if this is possible.

But, my original concern was also about assignment operation between two image accessors. I'm not sure there is an OpenCL/SPIRV built-in function doing something like this for images, so I think we may end up with code breaking the rules as well.

[gmlueck]

But, my original concern was also about assignment operation between two image accessors. I'm not sure there is an OpenCL/SPIRV built-in function doing something like this for images, so I think we may end up with code breaking the rules as well.

This relates to the idea I mentioned above. If the "image_accessor" type really contained a pointer to the OpenCL image, then this wouldn't be a problem. Assigning one accessor to another would just assign one pointer to another, and this would not violate any SPIRV rules. However, as I said, it's not clear whether it's valid to even have a pointer to an OpenCL image.

Relying on optimizations still concerns me a bit

Yes, I think this is a problem. We cannot get into a situation where legal SYCL programs fail when we disable optimizations. If correctness relies on some compiler transformation, we need to ensure that the transformation always happens, regardless of whether optimizations are enabled.

[Fznamznon]

But, my original concern was also about assignment operation between two image accessors. I'm not sure there is an OpenCL/SPIRV built-in function doing something like this for images, so I think we may end up with code breaking the rules as well.

This relates to the idea I mentioned above. If the "image_accessor" type really contained a pointer to the OpenCL image, then this wouldn't be a problem. Assigning one accessor to another would just assign one pointer to another, and this would not violate any SPIRV rules. However, as I said, it's not clear whether it's valid to even have a pointer to an OpenCL image.

It seems pointers to OpenCL images are not allowed (even clang complains about it - https://godbolt.org/z/KefjKMzqf ). OpenCL C spec says:

An image type cannot be used to declare a variable, a structure or union field, an array of images, a pointer to an image, or the return type of a function.

Same thing is mentioned for samplers.
I've failed to find the same thing in SPIR-V spec though.

[gmlueck]

I've failed to find the same thing in SPIR-V spec though.

And this is what's really relevant for SYCL/DPC++. The OpenCL restrictions aren't directly relevant because SYCL code isn't compiled to OpenCL source, it's compiled to SPIR-V.

[bader]

The OpenCL restrictions aren't directly relevant because SYCL code isn't compiled to OpenCL source, it's compiled to SPIR-V.

I expect The OpenCL™ SPIR-V Environment Specification to apply OpenCL restrictions to SPIR-V.

[Fznamznon]

I expect The OpenCL™ SPIR-V Environment Specification to apply OpenCL restrictions to SPIR-V.

Being honest, I've failed to find some explicit restriction on having pointers to image/sampler (like I saw in OpenCL spec) in this spec either. But I suspect this is a bug in the OpenCL™ SPIR-V Environment Specification, or we will face some restrictions like "value of this type cannot be produced by this instruction.

In addition, I think even if we will use a pointer to image inside image accessors, I don't think that OpenCL will allow us kernel argument that is a pointer to an image, so likely we will receive an image object, then wrap it with a pointer and here probably face some SPIR-V limitation. It sounds weird for me if this is somehow allowed, because it looks like sampler/image cannot go through phi/select instructions because this way it is not traceable, but it is allowed to go be referenced through a pointer.

BTW, I somehow triggered this discussion, but I think it moves away from the purpose of this patch. The way how we store sampler/image objects in SYCL classes implementation doesn't lift SPIR-V restriction on having this types go through phi/select instructions. I even used OpenCL code to produce the test (because IR turns out much simpler) - https://godbolt.org/z/e1WGhEvYs , that means the problem exists there as well. So, can we proceed with this patch?

[bashbaug]

Being honest, I've failed to find some explicit restriction on having pointers to image/sampler (like I saw in OpenCL spec) in this spec either. But I suspect this is a bug in the OpenCL™ SPIR-V Environment Specification, or we will face some restrictions like "value of this type cannot be produced by this instruction.

Summarizing an offline discussion with @gmlueck : The SPIR-V spec (note the SPIR-V spec, not the OpenCL SPIR-V environment spec) is a bit vague what can and cannot be done with images and samplers. Images and samplers are defined to be "opaque types" in SPIR-V and have limitations that do not apply to "concrete types" like integers and floats and pointers but it's not easy to untangle exactly what the limitations are.

If we convince ourselves that the restrictions from the OpenCL C spec do not exist in the SPIR-V spec then I agree they should be documented in the OpenCL SPIR-V environment spec. I believe the intent was to not add new image capablities going through SPIR-V than were available via OpenCL C.

Please note that these restrictions can be relaxed via extensions if this functionality is needed for SYCL 2020.

[bader]

Can we allow using images/sampler types in phi/select instructions when running SPIR-V in Level Zero execution environment?
@bashbaug, do you think GPU compiler will be able to handle such SPIR-V files?

[bashbaug]

do you think GPU compiler will be able to handle such SPIR-V files?

No, I don't think it will. There are ways we could make this work but they are not enabled currently.

Can we allow using images/sampler types in phi/select instructions when running SPIR-V in Level Zero execution environment?

I consider this to be another variant of the extension caveat described above. Because this will not work today it would either need a Level Zero extension or a new Level Zero version that explicitly allows it.

This all assumes that these operations are allowed in SPIR-V, otherwise we need a SPIR-V extension as well.

[bader]

@Fznamznon, FYI - this PR can be superseded by #5283.

[maksimsab]

Issue of this PR might get back if we return simplifycfg in the future.

[bader]

Issue of this PR might get back if we return simplifycfg in the future.

That's what are tests for.

[Fznamznon]

I don't think that trying to fix the problem with samplers by fixing or disabling LLVM passes is a good idea. It is quick and simple, but who knows maybe other passes will expose the same problem. IMO, It is better to fix it by adding some requirements for SYCL spec and analyzing which SPIR-V spec requirements we can relax/clarify for modern hardware.
So, I wouldn't chase two rabbits with #5283 .