intel
[SYCL][Matrix] Replace size_t with uint32_t to specify joint matrix rows and columns
Currently we are using size_t for the mentioned purpose. And it feels a bit odd to use a type, which size might be different depending on your environment to specify parameters of the hardware instruction(s) for close-to-hardware feature in heterogeneous programming model. Also 64-bits seem to be an overkill to specify number of Rows of a matrix, we are unlikely to ever hit 32-bit limitations.
Note, that in SPIR-V https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/KHR/SPV_KHR_cooperative_matrix.asciidoc Rows/Cols are explicitly set to be constant instruction with scalar 32-bit integer type, so right now if we want to ensure conformance of our implementation we need to check for 32-bit overflow and then cast size_t to uint32_t for these parameters.
The reason to use size_t for "joint_matrix" shape rows and cols and things like: stride and coordinates in the joint_matrix_apply, is to be consistent with what is currently using in SYCL spec (sycl::marray, sub_group_sizes) and C++ (a span extent type is "size_t").
How these are handled in SPIRV? Adding @gmlueck to get his opinion on this.
Rows/Cols are explicitly set to be constant instruction with scalar 32-bit integer type, so right now if we want to ensure conformance of our implementation we need to check for 32-bit overflow and then cast size_t to uint32_t for these parameters.
Are we just talking about some extra validation in the LLVM-to-SPIRV translator? If so, this doesn't sound too onerous.
If using size_t was causing a performance slowdown, I would be more swayed. In that case, I'd suggest using int or unsigned as the type, though, rather than uint32_t. Using uint32_t just because that is the representation in SPIR-V seems like we would be leaking implementation details into the API.
Hi! There have been no updates for at least the last 60 days, though the ticket has assignee(s).
@dkhaldi, could I ask you to take one of the following actions? :)
- Please provide an update if you have any (or just a small comment if you don't have any yet).
- OR mark this issue with the 'confirmed' label if you have confirmed the problem/request and our team should work on it.
- OR close the issue if it has been resolved.
- OR take any other suitable action.
Thanks!
@MrSidims , as far as I remember, you are already fixing this at the SPIRV level to make sure things are aligned with the cooperative matrix spec. But I don't think we are planning any change on the SYCL level. Should we close this issue now?