llvm
llvm copied to clipboard
intel
[ESIMD] Fix printf for simd_view to one element
The corresponding LIT test verifying the fix: https://github.com/intel/llvm-test-suite/pull/1130
Even though I see how this can be useful, it might cause confusion if a function called printf is behaving differently from the normal printf.
Even though I see how this can be useful, it might cause confusion if a function called printf is behaving differently from the normal printf.
@rolandschulz, this confusion is only for the types for which printf traits exist, and all such cases will be documented. I can't see any reason for ESIMD users to want to print simd_view object reference instead of the pointed-to element value in printf("%d", vec[i]). Does this sound reasonable to you?
Even though I see how this can be useful, it might cause confusion if a function called printf is behaving differently from the normal printf.
Yes, that is a bit of confusion, which is still much better than thinking that experimental::printf is just unusable, because it prints garbage for such a simple and natural thing as simd_val[i]
Unfortunately, returning a reference to a value holder is not an option because 'simd_val' is normally a register, reads/writes to it need to be controlled by jit-compiler via read/write intrinsics. Thus need to return a proxy object (simd_view), which would require implicit/explicit conversion to underlying element type.
The alternative is to leave experimental::printf() as is and ask users to use conversion like (int)simd_val[i], which is simple solution, but is not obvious. I've seen customers who tried printf(format, simd_val[i]), got bad output and got stuck there thinking experimental::printf is buggy/unusable.
I said that this can cause confusion because it doesn't actual solve the problem. It just works around it for the one use case. But the actual problem (that implicit conversion doesn't happen in all cases) is still there. And people might be confused why they can print vec[i] but can't pass vec[1] to a SFINAE function (e.g. template<class T, typename=std::enable_if_t< std::is_arithmetic_v<T>>> void f(T);). Users anyhow need to learn that they need to trigger conversion (e.g. +vec[i] or (int)vec[i]) when passing a simd_view to a variadic function. Therefore I'm not sure you actual help them by hiding that detail for printf. But you know the users better and I leave it up to you to decide what they prefer.
people might be confused why they can print vec[i] but can't pass vec[1] to a SFINAE function (e.g. template<class T, typename=std::enable_if_t< std::is_arithmetic_v<T>>> void f(T);).
Yes, they can be confused, but that kind of confusion is resolved easily because compiler gives comments/notes on SFINAE mismatch, something like: "note: candidate function ... ignored, could not match the parameter ...", while 'printf' silently takes simd_view and gives always wrong result.
Users anyhow need to learn that they need to trigger conversion (e.g. +vec[i] or (int)vec[i]) when passing a simd_view to a variadic function.
IMO, passing to variadic functions is not a problem. The problem may be seen only at the terminal uses of simd_view. Some variadic function template <typename... Ts> funcA(Ts... Args) may have one of 'Args' with 'simd_view' type. Passing that 'Arg' further is Ok. Using 'Arg' in integer context is Ok too. Using it in mixed type expression may be not Ok (float + simd_view) - will give compilation error with additional info. Using it in printf (as terminal use from our point of view) is definitely not Ok as produces silent error.
I see 3 possible solutions:
-
Cancel the proposed fix, and explain users the conversion is needed. It is probably acceptable solution. It though is a sort of barrier when start working with ESIMD and using printf for debugging.
-
Keep going with the current fix in sycl::ext::oneapi::experimental::printf. The difficulty here is to protect & document the idea of having 'printf_arument_type_traits', especially if it happens to be usable only for esimd/simd_view.
-
Provide printf in 'esimd' namespace, i.e. ext::intel::experimental::esimd::printf() which would have work-around for simd_view.
IMO, passing to variadic functions is not a problem.
It might be OK. It depends on what the variadic function is doing. The kind of issues we got with sycl::vec before https://github.com/intel/llvm/pull/736 was: the user tries to call template<class T>void f(T, T); with an int and a vec<int,4>[0] and doesn't understand why T can't be deduced.
I see 3 possible solutions:
I don't think the 3rd solution is a good idea. Having different printf in different namespaces with different behavior just adds to the confusion.
I think we should understand why for the GPU BE requires the intrinsics for the element access and why the solution we use for sycl::vec isn't OK. I don't see why this should be harder to optimize without the intrinsics on the GPU compared to the CPU. Even if that isn't something we can quickly implement for the GPU BE, I think we should pursue it it medium term solution. Because it would give us the best usability.
I think we should understand why for the GPU BE requires the intrinsics for the element access and why the solution we use for
sycl::vecisn't OK. I don't see why this should be harder to optimize without the intrinsics on the GPU compared to the CPU. Even if that isn't something we can quickly implement for the GPU BE, I think we should pursue it it medium term solution. Because it would give us the best usability.
The approach used for sycl::vec today has pretty disturbing comment (https://github.com/intel/llvm/blob/sycl/sycl/include/sycl/types.hpp#L901) and has the link to another doc saying that returning reference is not recommended and may be prohibited in future: http://llvm.org/docs/GetElementPtr.html#can-gep-index-into-vector-elements