Review comments

Author: PietroGhg

clang/lib/CodeGen/BackendUtil.cpp

@@ -47,7 +47,6 @@
 #include "llvm/Passes/PassPlugin.h"
 #include "llvm/Passes/StandardInstrumentations.h"
 #include "llvm/ProfileData/InstrProfCorrelator.h"
-#include "llvm/SYCLLowerIR/CheckNDRangeSYCLNativeCPU.h"
 #include "llvm/SYCLLowerIR/CleanupSYCLMetadata.h"
 #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
@@ -1166,11 +1165,8 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
       if (LangOpts.EnableDAEInSpirKernels)
         MPM.addPass(DeadArgumentEliminationSYCLPass());
 
-      // We have to schedule the pass here because the native cpu pipeline
-      // is ran as part of a separate clang invocation, but we want the information
-      // in sycl-post-link.
       if (LangOpts.SYCLIsNativeCPU)
-        MPM.addPass(CheckNDRangeSYCLNativeCPUPass());
+        llvm::sycl::utils::addSYCLNativeCPUEarlyPasses(MPM);
       // Rerun aspect propagation without warning diagnostics.
       MPM.addPass(
           SYCLPropagateAspectsUsagePass(/*FP64ConvEmu=*/CodeGenOpts.FP64ConvEmu,

clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp

@@ -713,6 +713,8 @@ class BinaryWrapper {
                                     "__sycl_native_cpu_decls");
     auto *EntriesBegin = ConstantExpr::getGetElementPtr(GVar->getValueType(), GVar,
                                                  getSizetConstPair(0u, 0u));
+
+    // Add Native CPU specific properties to the nativecpu_program struct
     Constant *PropValue = NullPtr;
     if (NativeCPUProps.has_value()) {
       auto PropsOrErr = addSYCLPropertySetToModule(*NativeCPUProps);
@@ -724,6 +726,11 @@ class BinaryWrapper {
       auto T = addStructArrayToModule({S}, getSyclPropSetTy());
       PropValue = T.first;
     }
+
+    // Create the nativecpu_program struct.
+    // We add it to a ConstantArray of length 1 because the SYCL runtime expects a
+    // non-zero sized binary image, and this allows it to point the end of the binary
+    // image to the end of the array.
     auto *Program = ConstantStruct::get(NCPUProgramT, {EntriesBegin, PropValue});
     ArrayType *ProgramATy = ArrayType::get(NCPUProgramT, 1);
     Constant *CPA = ConstantArray::get(ProgramATy, {Program});

llvm/include/llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h

@@ -11,13 +11,19 @@
 //===----------------------------------------------------------------------===//
 #pragma once
 #include "llvm/ADT/Twine.h"
+#include "llvm/IR/Module.h"
 #include "llvm/IR/PassManager.h"
 #include "llvm/Passes/OptimizationLevel.h"
 
 namespace llvm {
 namespace sycl {
 namespace utils {
 
+
+// Used to schedule passes in the device compiler cc1 invocation for
+// Native CPU.
+void addSYCLNativeCPUEarlyPasses(ModulePassManager &MPM);
+
 void addSYCLNativeCPUBackendPasses(ModulePassManager &MPM,
                                    ModuleAnalysisManager &MAM,
                                    OptimizationLevel OptLevel);

llvm/lib/SYCLNativeCPUUtils/CheckNDRangeSYCLNativeCPU.cpp

@@ -1,4 +1,4 @@
-//===------ PrepareSYCLNativeCPU.cpp - SYCL Native CPU Preparation Pass ---===//
+//- CheckNDRangeSYCLNativeCPU.cpp - Check if a kernel uses nd_range features -//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -23,10 +23,19 @@
 
 using namespace llvm;
 
-// TODO: add other bts
-static std::array<const char *, 5> ndFunctions{
-    "_Z23__spirv_WorkgroupSize_xv", "_Z23__spirv_NumWorkgroups_xv",
-    "_Z21__spirv_WorkgroupId_xv", "_Z27__spirv_LocalInvocationId_xv",
+static std::array<const char *, 13> NdFunctions{
+    "_Z23__spirv_WorkgroupSize_xv",
+    "_Z23__spirv_WorkgroupSize_yv",
+    "_Z23__spirv_WorkgroupSize_zv",
+    "_Z23__spirv_NumWorkgroups_xv",
+    "_Z23__spirv_NumWorkgroups_yv",
+    "_Z23__spirv_NumWorkgroups_zv",
+    "_Z21__spirv_WorkgroupId_xv",
+    "_Z21__spirv_WorkgroupId_yv",
+    "_Z21__spirv_WorkgroupId_zv",
+    "_Z27__spirv_LocalInvocationId_xv",
+    "_Z27__spirv_LocalInvocationId_yv",
+    "_Z27__spirv_LocalInvocationId_zv",
     "_Z22__spirv_ControlBarrierjjj"};
 
 static void addNDRangeMetadata(Function &F, bool Value) {
@@ -56,8 +65,8 @@ CheckNDRangeSYCLNativeCPUPass::run(Module &M, ModuleAnalysisManager &MAM) {
         for (auto &I : BB) {
           if (auto CI = dyn_cast<CallInst>(&I)) {
             auto CalleeName = CI->getCalledFunction()->getName();
-            if (std::find(ndFunctions.begin(), ndFunctions.end(), CalleeName) !=
-                ndFunctions.end()) {
+            if (std::find(NdFunctions.begin(), NdFunctions.end(), CalleeName) !=
+                NdFunctions.end()) {
               IsNDRange = true;
               break;
             }

llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp

@@ -131,3 +131,7 @@ void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(
     MPM.addPass(DumpIR());
   }
 }
+
+void llvm::sycl::utils::addSYCLNativeCPUEarlyPasses(ModulePassManager &MPM) {
+        MPM.addPass(CheckNDRangeSYCLNativeCPUPass());
+}