[NATIVECPU][PERF] add noalias to state struct pointer (#19288)

Adding the `noalias` attribute to the state struct pointer as it cannot
alias with any other pointer parameter.

Author: uwedolinsky

llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp

@@ -161,6 +161,8 @@ Function *cloneFunctionAndAddParam(Function *OldF, Type *T,
   if (!OldF->isDeclaration())
     CloneFunctionInto(NewF, OldF, VMap,
                       CloneFunctionChangeType::LocalChangesOnly, ReturnInst);
+  if (StateArgTLS == nullptr)
+    NewF->addParamAttr(Args.size() - 1, llvm::Attribute::NoAlias);
   return NewF;
 }
 

sycl/test/check_device_code/native_cpu/native_cpu_builtins.cpp

@@ -43,15 +43,15 @@ int main() {
   deviceQueue.submit([&](sycl::handler &h) {
     h.parallel_for<Test2>(
         r2, [=](sycl::nd_item<2> ndi) { acc[ndi.get_global_id(1)] = 42; });
-    // CHECK: @_ZTS5Test2.NativeCPUKernel(ptr {{.*}}%0, ptr {{.*}}%1, ptr addrspace(1) %2)
+    // CHECK: @_ZTS5Test2.NativeCPUKernel(ptr {{.*}}%0, ptr {{.*}}%1, ptr addrspace(1) noalias %2)
     // CHECK: call{{.*}}__dpcpp_nativecpu_get_global_id(i32 0, ptr addrspace(1) %2)
   });
   sycl::nd_range<3> r3({1, 1, 1}, {1, 1, 1});
   deviceQueue.submit([&](sycl::handler &h) {
     h.parallel_for<Test3>(r3, [=](sycl::nd_item<3> ndi) {
       acc[ndi.get_global_id(2)] = ndi.get_global_range(0);
     });
-    // CHECK: @_ZTS5Test3.NativeCPUKernel(ptr {{.*}}%0, ptr {{.*}}%1, ptr addrspace(1) %2)
+    // CHECK: @_ZTS5Test3.NativeCPUKernel(ptr {{.*}}%0, ptr {{.*}}%1, ptr addrspace(1) noalias %2)
     // CHECK-DAG: call{{.*}}__dpcpp_nativecpu_get_global_range(i32 2, ptr addrspace(1) %2)
     // CHECK-DAG: call{{.*}}__dpcpp_nativecpu_get_global_id(i32 0, ptr addrspace(1) %2)
   });