[InferAddressSpaces] Add InferAddressSpaces pass to pipeline for SPIR

[ghost]

Clang generates a ton of 'addrspacecast' instructions to cast temporary variables allocated to store kernel arguments in the addrspace(4) address space as well as to load elements of the arguments:

define weak_odr dso_local spir_kernel void @_ZTSZZ17compute_(float noundef %_arg_, %"class.cl::sycl::range.0"* noundef byval(%"class.cl::sycl::range.0") align 8 %_arg_1) {
entry:
; storing:
  %0 = alloca %class.anon, align 8
  %1 = addrspacecast %class.anon* %0 to %class.anon addrspace(4)*
  %2 = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %1, i64 0, i32 0
  store float %_arg_, float addrspace(4)* %2, align 8, !tbaa !9

; loading
  %3 = getelementptr inbounds %"class.cl::sycl::range.0", %"class.cl::sycl::range.0"* %_arg_1, i64 0, i32 0, i32 0, i64 0
  %4 = addrspacecast i64* %3 to i64 addrspace(4)*
  %5 = load i64, i64 addrspace(4)* %4, align 8
}

For the SPIR/SPIR-V target, addrspace(4) is the generic address space and these addrspacecast instructions can be safely removed from the code:

define weak_odr dso_local spir_kernel void @_ZTSZZ17compute_(float noundef %_arg_, %"class.cl::sycl::range.0"* noundef byval(%"class.cl::sycl::range.0") align 8 %_arg_1) {
entry:
; storing
  %0 = alloca %class.anon, align 8
  %1 = getelementptr inbounds %class.anon, %class.anon* %0, i64 0, i32 0
  store float %_arg_, float* %1, align 8, !tbaa !9

; loading
  %2 = getelementptr inbounds %"class.cl::sycl::range.0", %"class.cl::sycl::range.0"* %_arg_1, i64 0, i32 0, i32 0, i64 0
  %3 = load i64, i64* %2, align 8
}

To perform this removing, the InferAddressSpaces pass has been added to the clang optimization pipeline for SPIR and SPIR-V targets. This pass should be run after the other optimization passes (both function and module) and, it is very important, after inlining to let the pass "understand" from which address space as many as possible variables came and eliminate as many as possible addrspacecast instructions.

The elimination of redundant addrspacecast instruction decreases the size of the generated SPIR-V module and therefore makes less pressure on the backend JIT compilers.

[ghost]

/summary:run