[SYCL][NATIVECPU] Emit LocalInvocationId free function on Native CPU (#…

…12810)

Similarly to the CUDA and Hip backends, the Native CPU pass pipeline
expects the work item builtins to be of the "free function" kind, this
PR ensures that for code that goes through `parallel_for_work_group`.

This PR fixes the `e2e/Basic/accessor/accessor.cpp` test on Native CPU.

llvm/include/llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h

@@ -47,6 +47,10 @@ inline llvm::Twine addSYCLNativeCPUSuffix(StringRef S) {
  return llvm::Twine(S, SYCLNATIVECPUSUFFIX);
 }
 
+inline bool isSYCLNativeCPU(const Module &M) {
+ return M.getModuleFlag("is-native-cpu") != nullptr;
+}
+
 } // namespace utils
 } // namespace sycl
 } // namespace llvm

llvm/lib/SYCLLowerIR/LowerWGScope.cpp

@@ -97,6 +97,7 @@
 #include "llvm/IR/Module.h"
 #include "llvm/InitializePasses.h"
 #include "llvm/Pass.h"
+#include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h"
 #include "llvm/Support/CommandLine.h"
 
 #ifndef NDEBUG
@@ -898,7 +899,7 @@ GlobalVariable *spirv::createWGLocalVariable(Module &M, Type *T,
 // Return a value equals to 0 if and only if the local linear id is 0.
 Value *spirv::genPseudoLocalID(Instruction &Before, const Triple &TT) {
  Module &M = *Before.getModule();
- if (TT.isNVPTX() || TT.isAMDGCN()) {
+ if (TT.isNVPTX() || TT.isAMDGCN() || sycl::utils::isSYCLNativeCPU(M)) {
  LLVMContext &Ctx = Before.getContext();
  Type *RetTy = getSizeTTy(M);