Model AddrSpaceCasts in SCEV


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];

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.