[SYCL][NATIVECPU] Remove unneeded accesses to state thread_local (#17215)

No longer generate unneeded reads from state struct thread_local.

Also removed unused local.

Author: uwedolinsky

llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp

@@ -210,10 +210,24 @@ static Function *getReplaceFunc(Module &M, StringRef Name, const Use &U,
 }
 
 static Value *getStateArg(Function *F, llvm::Constant *StateTLS) {
-  // Todo: we should probably cache the state thread local load here
-  // to avoid re-emitting it for each builtin
   if (StateTLS) {
-    IRBuilder<> BB(&*F->getEntryBlock().getFirstInsertionPt());
+    // Find previous read from thread_local, if any
+    const auto IP = F->getEntryBlock().getFirstInsertionPt();
+    if (IP.isValid()) {
+      if (const CallInst *I = dyn_cast<CallInst>(&*IP)) {
+        if (I->getIntrinsicID() == Intrinsic::threadlocal_address &&
+            I->getOperand(0) == StateTLS) {
+          const auto Next = std::next(IP);
+          if (Next.isValid()) {
+            if (LoadInst *LI = dyn_cast<LoadInst>(&*Next)) {
+              if (LI->getPointerOperand() == I)
+                return LI;
+            }
+          }
+        }
+      }
+    }
+    IRBuilder<> BB(&*IP);
     llvm::Value *V = BB.CreateThreadLocalAddress(StateTLS);
     return BB.CreateLoad(StateTLS->getType(), V);
   }
@@ -336,7 +350,6 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M,
 #ifdef NATIVECPU_USE_OCK
   {
     SmallSet<Function *, 5> RemovableFuncs;
-    SmallVector<Function *, 5> WrapperFuncs;
 
     for (auto &OldF : OldKernels) {
       // If vectorization occurred, at this point we have a wrapper function

sycl/test/check_device_code/native_cpu/native_cpu_builtins.cpp

@@ -28,9 +28,11 @@ int main() {
     // CHECK: call{{.*}}__dpcpp_nativecpu_get_global_id(i32 0, ptr addrspace(1) %2)
     // CHECK-NOT: @llvm.threadlocal
 
-    // CHECK-TL:      %[[VAL1:.*]] = call ptr addrspace(1) @llvm.threadlocal.address.p1(ptr addrspace(1) @_ZL28nativecpu_thread_local_state)
-    // CHECK-TL-NEXT  %[[VAL2:.*]] = load ptr addrspace(1), ptr addrspace(1) %VAL1, align 8
-    // CHECK-TL-NEXT  %{{.*}} = call i64 @__dpcpp_nativecpu_get_wg_size(i32 0, ptr addrspace(1) %VAL2)
+    // CHECK-TL: define void @_ZTSN4sycl3_V16detail19__pf_kernel_wrapperI5Test1EE.NativeCPUKernel({{.*}}
+    // CHECK-TL-NEXT:entry:
+    // CHECK-TL-NEXT:  %[[VAL1:.*]] = call ptr addrspace(1) @llvm.threadlocal.address.p1(ptr addrspace(1) @_ZL28nativecpu_thread_local_state)
+    // CHECK-TL-NEXT:  %[[VAL2:.*]] = load ptr addrspace(1), ptr addrspace(1) %[[VAL1]], align 8
+    // CHECK-TL-NEXT:  %{{.*}} = call i64 @__dpcpp_nativecpu_get_wg_size(i32 0, ptr addrspace(1) %[[VAL2]])
 
     // CHECK-TL:      %{{.*}} = call ptr addrspace(1) @llvm.threadlocal.address.p1(ptr addrspace(1) @_ZL28nativecpu_thread_local_state)
     // CHECK-TL-DAG: store ptr addrspace(1) %{{.*}}, ptr addrspace(1) %{{.*}}, align 8