[SYCL] Emit fpbuiltin version of function only for function with FP arguments. #17253
Conversation
floating point type arguments.
zahiraam
changed the title
zahiraam
had a problem deploying
to
WindowsCILock
GitHub Actions
Error
zahiraam
changed the title
zahiraam
had a problem deploying
to
WindowsCILock
GitHub Actions
Failure
zahiraam
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
| [=](id<1> wiID) { | ||
| // NFPA: call spir_func float @fdiv(ptr noundef byval({{.*}}) align 1 {{.*}}, ptr noundef byval({{.*}}) align 1 {{.*}}) | ||
| // NFPA-FAST: call reassoc nnan ninf nsz arcp afn spir_func nofpclass(nan inf) float @fdiv(ptr noundef byval({{.*}}) align 1 {{.*}}, ptr noundef byval({{.*}}) align 1 {{.*}}) | ||
| a[0] = fdiv(HalfValue1, HalfValue2); |
There was a problem hiding this comment.
I'm not sure, if it is a real-world example, do e have fdiv builtin in SYCL/OpenCL/C++?. What about a[0] = Value1 / Value2;, can the issue be observed in this case and does your patch fix it?
There was a problem hiding this comment.
@MrSidims Thanks for the review. Fixed it.
clang/lib/CodeGen/CGCall.cpp
Outdated
| // as fpbuiltins. | ||
| if (!getLangOpts().OffloadFP32PrecSqrt || | ||
| !getLangOpts().OffloadFP32PrecDiv) { | ||
| for (auto Arg : IRCallArgs) { |
There was a problem hiding this comment.
nit (if applicable): to avoid some static analyzers complains about 'auto causes copy'.
| for (auto Arg : IRCallArgs) { | |
| for (auto &Arg : IRCallArgs) { |
This comment was marked as outdated.
This comment was marked as outdated.
Sorry, something went wrong.
|
|
||
| #include "sycl.hpp" | ||
|
|
||
| namespace sycl { | ||
| class half {}; |
There was a problem hiding this comment.
IMHO it's important, that we properly define sycl::half class aka add to it half storage. So / operator mustn't be extern - instead it should have a definition. Otherwise the code here doesn't test the issue - I'm pretty sure, that _ZdvN4sycl4halfES0_ would never become llvm.fpbuiltin intrinsic anyway.
There was a problem hiding this comment.
I have added a minimalist class half.
clang/lib/CodeGen/CGCall.cpp
Outdated
| if (hasFPAccuracyFuncMap || hasFPAccuracyVal || isFp32SqrtFunction) { | ||
| bool ArgsTypeIsFloat = true; | ||
| // In sycl mode, functions' arguments of type half are expanded | ||
| // to pointer types. Exclude these functions from bein emitted |
There was a problem hiding this comment.
| // to pointer types. Exclude these functions from bein emitted | |
| // to pointer types. Exclude these functions from being emitted |
|
@intel/llvm-gatekeepers Can this be merged in please? |
|
@sarnex Thanks. |
|
sure |
When the arguments of a function are of half types, such as
sqrt(half_type v)they are expanded to:load ptr addrspace(4), ptr addrspace(4) %a.addr.ascast, align 8, !tbaa !5This will generate a crash in the BE due to type mismatch. The expected type of the function is
float.This patch fixes the issue by restricting the emission of
fpbuiltinonly to functions that take floating point values arguments.