[ESIMD] Only dword types are supported by slm_load/store #4168
Conversation
| @@ -654,15 +654,17 @@ SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void slm_init(uint32_t size); | |||
| /// Only allow simd-16 and simd-32. | |||
| template <typename T, int n> | |||
| ESIMD_INLINE ESIMD_NODEBUG | |||
| typename sycl::detail::enable_if_t<(n == 16 || n == 32), simd<T, n>> | |||
| typename sycl::detail::enable_if_t<(n == 16 || n == 32) && (sizeof(T) == 4), | |||
There was a problem hiding this comment.
This message may support up to 4-byte (e.g., sizeof(T) == 1 or sizeof(T) == 2 can be also supported with proper fix).
There was a problem hiding this comment.
Thanks, @kychendev, does the same apply to slm_load4/store4 ? In the current implementation, they only support dword types only. I can change that as well.
There was a problem hiding this comment.
@kychendev, I wrote a test to make sure slm_load and slm_store support BYTE and WORD operand sizes and it doesn't work correctly: intel/llvm-test-suite#380. Can you please check the test to see if we have a bug somewhere, or maybe I'm missing something?
There was a problem hiding this comment.
I think the 1-byte problem is a BE glitch. From internal JIRA (by @kvladimi)
WA is to generate svm.gather that reads 2 16-bit elements from channel and then rdregion with stride 2 to ignore upper parts like this:
%res64.i.i = call <16 x i16> @llvm.genx.svm.gather.v16i16.v8i1.v8i64(<8 x i1> <i1 true, i1 true, i1 true, i1 true, i1 true, i1 true, i1 true, i1 true>, i32 1, <8 x i64> %new_offsets.i, <8 x i16> undef) %res.i.i = call <8 x i16> @llvm.genx.rdregioni.v8i16.v16i16.i16(<16 x i16> %res64.i.i, i32 0, i32 8, i32 2, i16 0, i32 undef)Requirements for SLM are exactly the same, I believe. SLM is just scatter.scaled which shall be processed in exactly the same way.
So, @DenisBakhvalov - I think some more code is needed for 1-byte case.
DenisBakhvalov
force-pushed
the
esimd_slm_double
branch
from
668ed90 to
5ead0cb
Compare
No description provided.