From dc6b52625d84dc2b9ba1246a5ea6e62b5c2c5e40 Mon Sep 17 00:00:00 2001 From: Uwe Dolinsky Date: Wed, 2 Jul 2025 12:45:40 +0100 Subject: [PATCH] [NATIVECPU] add noalias to state struct pointer --- llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp | 2 ++ .../test/check_device_code/native_cpu/native_cpu_builtins.cpp | 4 ++-- 2 files changed, 4 insertions(+), 2 deletions(-) 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) });