[SYCL] Fix zero dimension accessors on FPGA in AOT mode

[cperkinsintel]

Zero dimension accessors aren't working on FPGA with atomic data. When aot compiling for FPGA we use address space global_device_space instead of just global_space. This seems to be confusing the specialization for the zero dimension accessor. Not overspecifying the address space fixes the problem with no other change in functionality.

[cperkinsintel]

ping to reviewers.

[romanovvlad]

1. Could you please add a test?

Could you please add more details on why it helps and why it's correct?
Currently we have accessor<..., global_device_space> -> atomic<..., global_device_space> which looks correct.
With the patch we will have: accessor<..., global_device_space> -> atomic<..., global_space>, this looses "special" address space.

Maybe instead of converting everything to global_space we could have to methods:

1. One works with every address space and does conversion to AS
2. Another works with global_device_space and global_host_space and does conversion to accessor<..., global_space> by firstly converting to atomic<..., global_device_space> using the first method, then converting to atomic<..., global_space> using conversion provided by atomic itself on https://github.com/intel/llvm/blob/sycl/sycl/include/CL/sycl/atomic.hpp#L203

[cperkinsintel]

@romanovvlad - it's in support of some new testing coming to SYCL_CTS. But I can add a test.

As to why it is working, I honestly don't know. I agree with you that the original code seems fine. But it is not fine. When compiling with the global_device_space address space in the return type of these atomic values, the enable_if_t SFINAE fails and the specialization is then not available, leading to compile failures. Without the AS included in the specialization, the enable_f_t passes and the specialization is generated. And the actual values that are created (for the multi_ptr and the owning atomic) DO have the desired global_device_space address space.

[cperkinsintel]

@GarveyJoe , I believe you implemented the original code for the USM address spaces, can you advise? We have a remote case where zero dimension accessors with access::atomic aren't getting the correct template specialization. It seems that when global_device_space is used as the address space in the return type, that the SFINAE fails and then the specialization isn't generated leading to problems when using the FPGA aot compiler.

I frankly don't see what the root of the problem is, it looks like it should work to me. But it isn't. My simple solution here is to simply not use the address space when returning the type for atomic from the enable_if_t SFINAE, but keep it when we actually instantiate the type. That way the value with global_device_space gets to pun in place for global_space, which I now believe is the right intention. But I'm not sure. Should I instead write another conversion, this one from global space to something? (like global_device_space or global_host_space). That doesn't seem correct. Anyway, your insight would be very welcome.

[MrSidims]

Correction, not Joe but me implemented the device/host address spaces. I'll take a look at the issue/patch when I refresh my memory. I remember, there were difficulties to make atomics working.

[MrSidims]

/summary:run

[MrSidims]

A test would be helpful to understand the issue you are facing. Is it some sort of substituted type causes a side-effect problem?

[cperkinsintel]

@MrSidims - I added a test, something I meant to do last week before I got pulled into something else.

[cperkinsintel]

@MrSidims any chance to review this yet? The test should demonstrate the problem pretty well. It's definitely in a far corner.

[MrSidims]

// note: candidate function not viable: no known conversion from 'const sycl::accessor<int, 0, sycl::access::mode::atomic, sycl::access::target::global_buffer>' to 'atomic_t' (aka 'atomic<int>') for 1st argument

So it does look like, that existence of implicit conversion confuses SFINAE. Okay, since summary run was successful the patch is LGTM.
Not sure whether it's a good suggestion, but may be it worth to guard this change with ENABLE_USM macro.

[cperkinsintel]

ping to reviewers - this should be ready for review.

@MrSidims - guarding behind the ENABLE_USM_SPACE macro is a good idea. Done.

[MrSidims]

LGTM, yet it should be approved by RT team member