<div dir="ltr">Hi,<div><br></div><div>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.</div><div><br></div><div><font face="monospace">{<br>    cl::sycl::buffer<double, 1> A_buf(A, cl::sycl::range<1>(N));<br>    cl::sycl::buffer<double, 1> B_buf(B, cl::sycl::range<1>(N));<br>    cl::sycl::buffer<double, 1> C_buf(C, cl::sycl::range<1>(N));<br><br><br>    deviceQueue.submit([&](cl::sycl::handler& cgh) {<br>        auto A_acc = A_buf.get_access<sycl_read>(cgh);<br>        auto B_acc = B_buf.get_access<sycl_read>(cgh);<br>        auto C_acc = C_buf.get_access<sycl_write>(cgh);<br><br>        auto kern = [=]() {<br>          #pragma clang loop vectorize(enable) vectorize_width(4) interleave_count(2)<br>          for (int i = 0; i < N; i++) {<br>            C_acc[i] = A_acc[i] + B_acc[i];<br>          }<br>        };<br>        cgh.single_task<class vec_add>(kern);<br>    });<br>  }</font><br></div><div><br></div><div><br></div><div>I found that hacking the ScalarEvolution.cpp in the <font face="monospace">createSCEV(Value *V)</font> 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.<br><br>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 <font face="monospace">isNoopAddrSpaceCast(FromAS, ToAS)</font> returns true?<br><br>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.<br><br>-- <br><div dir="ltr" class="gmail_signature" data-smartmail="gmail_signature"><div dir="ltr"><div dir="ltr"><div style="color:rgb(32,33,36)"><font color="#000000">João Paulo L. de Carvalho<br>Ph.D Computer Science |  IC-UNICAMP | Campinas , SP - Brazil<br>Postdoctoral Research Fellow | University of Alberta | Edmonton, AB - Canada</font></div><div style="color:rgb(136,136,136)"><a href="mailto:joao.carvalho@ic.unicamp.br" target="_blank">joao.carvalho@ic.unicamp.br</a><br><a href="mailto:joao.carvalho@ualberta.ca" target="_blank">joao.carvalho@ualberta.ca</a></div></div></div></div></div></div>