diff --git a/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp b/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp index fa68506218c80..c7955b41f2d6d 100644 --- a/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp +++ b/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; } diff --git a/sycl/test/check_device_code/native_cpu/native_cpu_builtins.cpp b/sycl/test/check_device_code/native_cpu/native_cpu_builtins.cpp index 310c4cc0e2f6e..e46bd95351920 100644 --- a/sycl/test/check_device_code/native_cpu/native_cpu_builtins.cpp +++ b/sycl/test/check_device_code/native_cpu/native_cpu_builtins.cpp @@ -43,7 +43,7 @@ int main() { deviceQueue.submit([&](sycl::handler &h) { h.parallel_for( 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}); @@ -51,7 +51,7 @@ int main() { h.parallel_for(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) });