[SYCL] Add element size argument to piKernelSetArg #5104
Conversation
This patch comes from an attempt to fix intel#5007. The issue there is that for local kernel argument the CUDA plugin uses CUDA dynamic shared memory, which gives us a single chunk of shared memory to work with. The CUDA plugin then lays out all the local kernel arguments consecutively in this single chunk of memory. And this can cause issues because simply laying the arguments out one after the other can result in misaligned arguments. In intel#5007 for example there is an `int` argument followed by a `double4` argument, so the `double4` argument ends up with the wrong alignment, only being aligned on a 4 bytes boundary following from the `int`. It is possible to adjust this and fixup the alignment when laying out the local kernel arguments in the CUDA plugin, however before this patch the only information in the plugin would be the total size of local memory required for the given arguments, which doesn't tell us anything about the required alignment. So this patch propagates the size of the elements inside of the local accessor all the way down to the PI plugin through `piKernelSetArg`, and tweaks the local argument layout in the CUDA plugin to use the type size as alignment for local kernel arguments.
|
|
||
| cl::sycl::detail::kernel_param_kind_t MType; | ||
| void *MPtr; | ||
| int MSize; | ||
| int MIndex; | ||
| int MElemSize; |
There was a problem hiding this comment.
The patch should break ABI since this structure crosses library boundaries. Breaking ABI is not allowed right now.
Can we use the strictest required alignment for vector operations as a default alignment, which I believe it should be |
Yeah, after looking at this a bit more I think you're right, we could use the largest vector size for this, I'll close this PR and open a separate one with a change to that effect. |
The issue there is that for local kernel argument the CUDA plugin uses CUDA dynamic shared memory, which gives us a single chunk of shared memory to work with. The CUDA plugin then lays out all the local kernel arguments consecutively in this single chunk of memory. And this can cause issues because simply laying the arguments out one after the other can result in misaligned arguments. So this patch is changing the argument layout to align them to the maximum necessary alignment which is the size of the largest vector type. Additionally if there is a local buffer smaller than this maximum alignment, the size of that buffer is simply used for alignment. This fixes the issue in intel#5007. See also the discussion on intel#5104 for alternative solution, that may be more efficient but would require a more intrusive ABI changing patch.
The issue there is that for local kernel argument the CUDA plugin uses CUDA dynamic shared memory, which gives us a single chunk of shared memory to work with. The CUDA plugin then lays out all the local kernel arguments consecutively in this single chunk of memory. And this can cause issues because simply laying the arguments out one after the other can result in misaligned arguments. So this patch is changing the argument layout to align them to the maximum necessary alignment which is the size of the largest vector type. Additionally if there is a local buffer smaller than this maximum alignment, the size of that buffer is simply used for alignment. This fixes the issue in #5007. See also the discussion on #5104 for alternative solution, that may be more efficient but would require a more intrusive ABI changing patch.
This patch comes from an attempt to fix #5007.
The issue there is that for local kernel argument the CUDA plugin uses
CUDA dynamic shared memory, which gives us a single chunk of shared
memory to work with.
The CUDA plugin then lays out all the local kernel arguments
consecutively in this single chunk of memory.
And this can cause issues because simply laying the arguments out one
after the other can result in misaligned arguments. In #5007 for example
there is an
intargument followed by adouble4argument, so thedouble4argument ends up with the wrong alignment, only being alignedon a 4 bytes boundary following from the
int.It is possible to adjust this and fixup the alignment when laying out
the local kernel arguments in the CUDA plugin, however before this patch
the only information in the plugin would be the total size of local
memory required for the given arguments, which doesn't tell us anything
about the required alignment.
So this patch propagates the size of the elements inside of the
local accessor all the way down to the PI plugin through
piKernelSetArg, and tweaks the local argument layout in the CUDAplugin to use the type size as alignment for local kernel arguments.
I'm not entirely sure if this is the best approach so feedback on this would be appreciated, this patch may also need to be refined for naming and/or position of the extra argument, however it does fix the issue in #5007