[llvm-dev] Model AddrSpaceCasts in SCEV

[João Paulo Labegalini de Carvalho via llvm-dev]

Hi,

I noticed that simple loops (see example below) are not vectorized due to
SCEV not handling AddrSpaceCast and treating them as SCEVUnknowns. I would
like to allow the LoopAccessAnalysis pass to create dynamic array bound
checks and enable the vectorizer to generate guarded vectorized code.

{
    cl::sycl::buffer<double, 1> A_buf(A, cl::sycl::range<1>(N));
    cl::sycl::buffer<double, 1> B_buf(B, cl::sycl::range<1>(N));
    cl::sycl::buffer<double, 1> C_buf(C, cl::sycl::range<1>(N));


    deviceQueue.submit([&](cl::sycl::handler& cgh) {
        auto A_acc = A_buf.get_access<sycl_read>(cgh);
        auto B_acc = B_buf.get_access<sycl_read>(cgh);
        auto C_acc = C_buf.get_access<sycl_write>(cgh);

        auto kern = [=]() {
          #pragma clang loop vectorize(enable) vectorize_width(4)
interleave_count(2)
          for (int i = 0; i < N; i++) {
            C_acc[i] = A_acc[i] + B_acc[i];
          }
        };
        cgh.single_task<class vec_add>(kern);
    });
  }


I found that hacking the ScalarEvolution.cpp in the createSCEV(Value *V)
method and treat AddrSpaceCast as BitCasts enables the LoopAccessAnalysis
to generate array bounds checks that guard the execution of the vectorized
version of the for loop body.

Obviously this is not a correct solution, thus I ask for guidance on how to
properly handle AddrSpaceCasts. Would it be enough to rely on
TargetTransformInfo and skip AddrSpaceCasts whenever the predicate
method isNoopAddrSpaceCast(FromAS,
ToAS) returns true?

I believe that for all cases where a dynamic bounds check would be
generated it is safe to do it, right? The cases I want to properly handle
are those where the SCEV would wrongly conclude that an expression
evaluates to the same value or is equivalent to another because I wrongly
treated AddrSpaceCasts as BitCast, which does not change the pointer value.

-- 
João Paulo L. de Carvalho
Ph.D Computer Science |  IC-UNICAMP | Campinas , SP - Brazil
Postdoctoral Research Fellow | University of Alberta | Edmonton, AB - Canada
joao.carvalho at ic.unicamp.br
joao.carvalho at ualberta.ca
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20201127/2037c5f8/attachment.html>