[SYCL][FPGA] Expose value_type and min_capacity from sycl::ext::intel::experimental::pipe

[pcolberg]

This applies commit e1619fa64616a05ae962952b7bc6517af9f74dda to the experimental implementation to match the supported implementation.

https://github.com/intel/llvm/pull/5471 https://github.com/intel/llvm/pull/5895

Cc: @tyoungsc @shuoniu-intel @rho180 @zibaiwan

[pcolberg]

The second commit in https://github.com/intel/llvm/pull/5896 stacks on top of https://github.com/intel/llvm/pull/5895. I will drop the first commit in https://github.com/intel/llvm/pull/5896 once https://github.com/intel/llvm/pull/5895 is merged.

[pcolberg]

@tyoungsc I copied the changes to the specification and the implementation verbatim from your commit e1619fa64616a05ae962952b7bc6517af9f74dda, but noticed that the type of min_capacity is size_t in the specification while it is int32_t in the implementation. Do you know where this difference stems from? Are the types identical in practice? Should either be changed to match the either?

[tyoungsc]

@tyoungsc I copied the changes to the specification and the implementation verbatim from your commit e1619fa, but noticed that the type of min_capacity is size_t in the specification while it is int32_t in the implementation. Do you know where this difference stems from? Are the types identical in practice? Should either be changed to match the either?

I don't know where this stems from :(. No, they will not be identical, size_t will likely be a uint64_t. I also have no clue why my changes caused the build to break and am not treating this with high priority.

[gmlueck]

Needless to say, the spec and the implementation should agree. I see that you also use size_t for the map and unmap member functions of the host pipe class. Should those types match the type of min_capacity?

[tyoungsc]

Needless to say, the spec and the implementation should agree. I see that you also use size_t for the map and unmap member functions of the host pipe class. Should those types match the type of min_capacity?

IMO they should all match and be unsigned, since a negative value for any of them doesn't make sense. Whether we choose uint32_t or uint64_t (size_t), I don't have a preference. Since the spec currently uses size_t for all, I would say we should switch all to size_t.

[pcolberg]

Needless to say, the spec and the implementation should agree. I see that you also use size_t for the map and unmap member functions of the host pipe class. Should those types match the type of min_capacity?

IMO they should all match and be unsigned, since a negative value for any of them doesn't make sense. Whether we choose uint32_t or uint64_t (size_t), I don't have a preference. Since the spec currently uses size_t for all, I would say we should switch all to size_t.

Thanks @gmlueck and @tyoungsc!

@GarveyJoe do you agree with using size_t for min_capacity in the implementation, or do you see a reason for uint32_t, e.g., hardware limitations, efficiency, etc., in which case the specification would need to be adjusted?

[pcolberg]

The SPIRV API uses int32_t for data type size, alignment, and min_capacity, see ConstantPipeStorage and __spirv_ReadPipe/__spirv_WritePipe. So really the question is whether int32_t should be exposed to the user or not. Since the choice has already been made, we might want to keep it as is and update the specification to match the implementation, as to not break the supported ABI.

kernel_readable_io_pipe and kernel_writeable_io_pipe do use size_t for the min_capacity template parameter, but expose it as an int32_t static member.

[GarveyJoe]

Mike and I must have missed this when initially working on the pipes spec. The SPIR-V spec is clear that the capacity of a pipe must be a 32-bit unsigned integer but at the same time we don't want to break ABI compatibility.

As a compromise we could keep the template parameter as a size_t and make the member a size_t but static_assert if the user tries to use a min_capacity that doesn't fit in an int32_t stating that our compiler can't handle pipes that large. Any existing code that uses a capacity that doesn't fit in an int32_t is presumably broken right now anyways since it can't be translated into SPIR-V so this shouldn't cause any regressions in compiling code.

If we do this we should probably add to the spec a sentence saying that the maximum minimum_capacity supported by a device is implementation defined and that the implementation must emit a diagnostic if that maximum is exceeded. We could also require some reasonable minimum maximum minimum_capacity. There is lots of precedent for this in both C++ and SPIR-V, such as maximum array size (array declarations take a size_t but an implementation doesn't need to be able to create an array of size equal to maximum value a size_t can hold).