Merge from 'sycl' to 'sycl-web' (12 commits)

Author: wwXing0

.github/workflows/sycl_precommit.yml

@@ -82,6 +82,7 @@ jobs:
     secrets: inherit
     with:
       build_ref: ${{ github.event.pull_request.head.sha }}
+      merge_ref: ${{ github.event.pull_request.base.sha }}
       build_cache_root: "/__w/"
       build_cache_size: "8G"
       build_artifact_suffix: "default"

clang/lib/CodeGen/CodeGenModule.cpp

@@ -2349,10 +2349,6 @@ void CodeGenModule::GenKernelArgMetadata(llvm::Function *Fn,
       Fn->setMetadata("kernel_arg_exclusive_ptr",
                       llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
     }
-    if (LangOpts.SYCLIsNativeCPU) {
-      Fn->setMetadata("kernel_arg_type",
-                      llvm::MDNode::get(VMContext, argTypeNames));
-    }
   } else {
     if (getLangOpts().OpenCL || getLangOpts().SYCLIsDevice) {
       Fn->setMetadata("kernel_arg_addr_space",

clang/lib/Driver/ToolChains/Clang.cpp

@@ -9812,7 +9812,8 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA,
         ",+SPV_INTEL_global_variable_decorations"
         ",+SPV_INTEL_fpga_buffer_location"
         ",+SPV_INTEL_fpga_argument_interfaces"
-        ",+SPV_INTEL_fpga_invocation_pipelining_attributes";
+        ",+SPV_INTEL_fpga_invocation_pipelining_attributes"
+        ",+SPV_INTEL_fpga_latency_control";
     ExtArg = ExtArg + DefaultExtArg + INTELExtArg;
     if (C.getDriver().IsFPGAHWMode())
       // Enable several extensions on FPGA H/W exclusively

clang/test/CodeGenSYCL/native_cpu_basic.cpp

@@ -1,6 +1,5 @@
 // This test checks for some basic Front End features for Native CPU:
 // * Kernel name mangling
-// * kernel_arg_type metadata node
 // * is-native-cpu module flag
 // RUN: %clang_cc1 -fsycl-is-device -S -emit-llvm -internal-isystem %S/Inputs -fsycl-is-native-cpu -o %t.ll %s 
 // RUN: FileCheck -input-file=%t.ll %s 
@@ -50,11 +49,9 @@ void gen() {
   test<float>(q);
 }
 
-// Check name mangling and kernel_arg_type metadata
-// CHECK-DAG: @_ZTS6init_aIiE_NativeCPUKernel({{.*}}){{.*}}!kernel_arg_type ![[TYPE1:[0-9]*]]
-// CHECK-DAG: @_ZTS6init_aIfE_NativeCPUKernel({{.*}}){{.*}}!kernel_arg_type ![[TYPE3:[0-9]*]]
-// CHECK-DAG: ![[TYPE1]] = !{!"int*", !"sycl::range<1>", !"sycl::range<1>", !"sycl::id<1>", !"int"}
-// CHECK-DAG: ![[TYPE3]] = !{!"float*", !"sycl::range<1>", !"sycl::range<1>", !"sycl::id<1>", !"float"}
+// Check name mangling 
+// CHECK-DAG: @_ZTS6init_aIiE_NativeCPUKernel_NativeCPUKernel({{.*}})
+// CHECK-DAG: @_ZTS6init_aIfE_NativeCPUKernel_NativeCPUKernel({{.*}})
 
 // Check Native CPU module flag
 // CHECK-DAG: !{{[0-9]*}} = !{i32 1, !"is-native-cpu", i32 1}

clang/test/Driver/sycl-spirv-ext.c

@@ -47,6 +47,7 @@
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_fpga_buffer_location
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_fpga_argument_interfaces
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_fpga_invocation_pipelining_attributes
+// CHECK-DEFAULT-SAME:,+SPV_INTEL_fpga_latency_control
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_token_type
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_bfloat16_conversion
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_joint_matrix
@@ -75,6 +76,7 @@
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_arithmetic_fence
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_fpga_buffer_location
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_fpga_argument_interfaces
+// CHECK-FPGA-HW-SAME:,+SPV_INTEL_fpga_latency_control
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_usm_storage_classes
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_runtime_aligned
 // CHECK-FPGA-HW-SAME:,+SPV_INTEL_fpga_cluster_attributes,+SPV_INTEL_loop_fuse

devops/actions/cached_checkout/action.yml

@@ -61,7 +61,7 @@ runs:
       git fetch origin $DEFAULT_BRANCH
       echo "FETCHED:"
       git log -1 origin/$DEFAULT_BRANCH
-      echo "Merging it into the current workspace"
+      echo "Merging ${{ inputs.merge_ref }} into the current workspace"
       # Set fake identity to fulfil git requirements
       git config --local user.email "[email protected]"
       git config --local user.name "GitHub Actions"

devops/containers/ubuntu2204_base.Dockerfile

@@ -13,8 +13,9 @@ RUN /install.sh
 # discover user home directory and fail a few LIT tests. Fixes UID and GID to
 # 1001, that is used as default by GitHub Actions.
 RUN groupadd -g 1001 sycl && useradd sycl -u 1001 -g 1001 -m -s /bin/bash
-# Add sycl user to video group so that it can access GPU
+# Add sycl user to video/irc groups so that it can access GPU
 RUN usermod -aG video sycl
+RUN usermod -aG irc sycl
 # Allow sycl user to run as sudo
 RUN echo "sycl  ALL=(ALL) NOPASSWD:ALL" >> /etc/sudoers
 

devops/containers/ubuntu2204_build.Dockerfile

@@ -27,8 +27,9 @@ RUN apt install -yqq libnuma-dev wget gnupg2 && \
 # discover user home directory and fail a few LIT tests. Fixes UID and GID to
 # 1001, that is used as default by GitHub Actions.
 RUN groupadd -g 1001 sycl && useradd sycl -u 1001 -g 1001 -m -s /bin/bash
-# Add sycl user to video group so that it can access GPU
+# Add sycl user to video/irc groups so that it can access GPU
 RUN usermod -aG video sycl
+RUN usermod -aG irc sycl
 
 COPY actions/cached_checkout /actions/cached_checkout
 COPY actions/cleanup /actions/cleanup

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -50,7 +50,7 @@ static const char *LegalSYCLFunctions[] = {
     "^sycl::_V1::multi_ptr<.+>::.+",
     "^sycl::_V1::nd_item<.+>::.+",
     "^sycl::_V1::group<.+>::.+",
-    "^sycl::_V1::sub_group<.+>::.+",
+    "^sycl::_V1::sub_group::.+",
     "^sycl::_V1::range<.+>::.+",
     "^sycl::_V1::kernel_handler::.+",
     "^sycl::_V1::cos",

llvm/lib/SYCLLowerIR/EmitSYCLNativeCPUHeader.cpp

@@ -33,133 +33,12 @@
 using namespace llvm;
 
 namespace {
-SmallVector<bool> getArgMask(const Function *F) {
-  SmallVector<bool> Res;
-  auto *UsedNode = F->getMetadata("sycl_kernel_omit_args");
-  if (!UsedNode) {
-    // the metadata node is not available if -fenable-sycl-dae
-    // was not set; set everything to true in the mask.
-    for (unsigned I = 0; I < F->getFunctionType()->getNumParams(); I++) {
-      Res.push_back(true);
-    }
-    return Res;
-  }
-  auto NumOperands = UsedNode->getNumOperands();
-  for (unsigned I = 0; I < NumOperands; I++) {
-    auto &Op = UsedNode->getOperand(I);
-    if (auto *CAM = dyn_cast<ConstantAsMetadata>(Op.get())) {
-      if (auto *Const = dyn_cast<ConstantInt>(CAM->getValue())) {
-        auto Val = Const->getValue();
-        Res.push_back(!Val.getBoolValue());
-      } else {
-        report_fatal_error("Unable to retrieve constant int from "
-                           "sycl_kernel_omit_args metadata node");
-      }
-    } else {
-      report_fatal_error(
-          "Error while processing sycl_kernel_omit_args metadata node");
-    }
-  }
-  return Res;
-}
 
-SmallVector<StringRef> getArgTypeNames(const Function *F) {
-  SmallVector<StringRef> Res;
-  auto *TNNode = F->getMetadata("kernel_arg_type");
-  assert(TNNode &&
-         "kernel_arg_type metadata node is required for sycl native CPU");
-  auto NumOperands = TNNode->getNumOperands();
-  for (unsigned I = 0; I < NumOperands; I++) {
-    auto &Op = TNNode->getOperand(I);
-    auto *MDS = dyn_cast<MDString>(Op.get());
-    if (!MDS)
-      report_fatal_error("error while processing kernel_arg_types metadata");
-    Res.push_back(MDS->getString());
-  }
-  return Res;
-}
-
-void emitKernelDecl(const Function *F, const SmallVector<bool> &ArgMask,
-                    const SmallVector<StringRef> &ArgTypeNames,
-                    raw_ostream &O) {
-  auto EmitArgDecl = [&](const Argument *Arg, unsigned Index) {
-    Type *ArgTy = Arg->getType();
-    if (isa<PointerType>(ArgTy))
-      return "void *";
-    return ArgTypeNames[Index].data();
-  };
-
-  auto NumParams = F->getFunctionType()->getNumParams();
-  O << "extern \"C\" void " << F->getName() << "(";
-
-  unsigned I = 0, UsedI = 0;
-  for (; I + 1 < ArgMask.size() && UsedI + 1 < NumParams; I++) {
-    if (!ArgMask[I])
-      continue;
-    O << EmitArgDecl(F->getArg(UsedI), I) << ", ";
-    UsedI++;
-  }
-
-  // parameters may have been removed.
-  bool NoUsedArgs = true;
-  for (auto &Entry : ArgMask) {
-    NoUsedArgs &= !Entry;
-  }
-  if (NoUsedArgs) {
-    O << ");\n";
-    return;
-  }
-  // find the index of the last used arg
-  while (!ArgMask[I] && I + 1 < ArgMask.size())
-    I++;
-  O << EmitArgDecl(F->getArg(UsedI), I) << ", __nativecpu_state *);\n";
-}
-
-void emitSubKernelHandler(const Function *F, const SmallVector<bool> &ArgMask,
-                          const SmallVector<StringRef> &ArgTypeNames,
-                          raw_ostream &O) {
-  SmallVector<unsigned> UsedArgIdx;
-  auto EmitParamCast = [&](Argument *Arg, unsigned Index) {
-    std::string Res;
-    llvm::raw_string_ostream OS(Res);
-    UsedArgIdx.push_back(Index);
-    if (isa<PointerType>(Arg->getType())) {
-      OS << "  void* arg" << Index << " = ";
-      OS << "MArgs[" << Index << "].getPtr();\n";
-      return OS.str();
-    }
-    auto TN = ArgTypeNames[Index].str();
-    OS << "  " << TN << " arg" << Index << " = ";
-    OS << "*(" << TN << "*)"
-       << "MArgs[" << Index << "].getPtr();\n";
-    return OS.str();
-  };
-
-  O << "\ninline static void " << F->getName() << "subhandler(";
+void emitSubKernelHandler(const Function *F, raw_ostream &O) {
+  O << "\nextern \"C\" void " << F->getName() << "subhandler(";
   O << "const sycl::detail::NativeCPUArgDesc *MArgs, "
-       "__nativecpu_state *state) {\n";
-  // Retrieve only the args that are used
-  for (unsigned I = 0, UsedI = 0;
-       I < ArgMask.size() && UsedI < F->getFunctionType()->getNumParams();
-       I++) {
-    if (ArgMask[I]) {
-      O << EmitParamCast(F->getArg(UsedI), I);
-      UsedI++;
-    }
-  }
-  // Emit the actual kernel call
-  O << "  " << F->getName() << "(";
-  if (UsedArgIdx.size() == 0) {
-    O << ");\n";
-  } else {
-    for (unsigned I = 0; I < UsedArgIdx.size() - 1; I++) {
-      O << "arg" << UsedArgIdx[I] << ", ";
-    }
-    if (UsedArgIdx.size() >= 1)
-      O << "arg" << UsedArgIdx.back();
-    O << ", state);\n";
-  }
-  O << "};\n\n";
+       "__nativecpu_state *state);\n";
+  return;
 }
 
 // Todo: maybe we could use clang-offload-wrapper for this,
@@ -254,10 +133,7 @@ PreservedAnalyses EmitSYCLNativeCPUHeaderPass::run(Module &M,
   O << "extern \"C\" void __sycl_register_lib(pi_device_binaries desc);\n";
 
   for (auto *F : Kernels) {
-    auto ArgMask = getArgMask(F);
-    auto ArgTypeNames = getArgTypeNames(F);
-    emitKernelDecl(F, ArgMask, ArgTypeNames, O);
-    emitSubKernelHandler(F, ArgMask, ArgTypeNames, O);
+    emitSubKernelHandler(F, O);
     emitSYCLRegisterLib(F, O);
   }
 

llvm/lib/SYCLLowerIR/PrepareSYCLNativeCPU.cpp

@@ -23,6 +23,7 @@
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/Instruction.h"
 #include "llvm/IR/Instructions.h"
+#include "llvm/IR/LLVMContext.h"
 #include "llvm/IR/Operator.h"
 #include "llvm/IR/Value.h"
 #include "llvm/InitializePasses.h"
@@ -63,7 +64,88 @@ void fixCallingConv(Function *F) {
     F->setLinkage(GlobalValue::LinkageTypes::WeakAnyLinkage);
 }
 
-// Clone the function and returns a new function with a new argument on type T
+// returns the indexes of the used arguments
+SmallVector<unsigned> getUsedIndexes(const Function *F) {
+  SmallVector<unsigned> res;
+  auto UsedNode = F->getMetadata("sycl_kernel_omit_args");
+  if (!UsedNode) {
+    // the metadata node is not available if -fenable-sycl-dae
+    // was not set; set everything to true
+    // Exclude one arg because we already added the state ptr
+    for (unsigned I = 0; I + 1 < F->getFunctionType()->getNumParams(); I++) {
+      res.push_back(I);
+    }
+    return res;
+  }
+  auto NumOperands = UsedNode->getNumOperands();
+  for (unsigned I = 0; I < NumOperands; I++) {
+    auto &Op = UsedNode->getOperand(I);
+    if (auto CAM = dyn_cast<ConstantAsMetadata>(Op.get())) {
+      if (auto Const = dyn_cast<ConstantInt>(CAM->getValue())) {
+        auto Val = Const->getValue();
+        if (!Val.getBoolValue()) {
+          res.push_back(I);
+        }
+      } else {
+        report_fatal_error("Unable to retrieve constant int from "
+                           "sycl_kernel_omit_args metadata node");
+      }
+    } else {
+      report_fatal_error(
+          "Error while processing sycl_kernel_omit_args metadata node");
+    }
+  }
+  return res;
+}
+
+void emitSubkernelForKernel(Function *F, Type *NativeCPUArgDescType,
+                            Type *StatePtrType) {
+  LLVMContext &Ctx = F->getContext();
+  Type *NativeCPUArgDescPtrType = PointerType::getUnqual(NativeCPUArgDescType);
+
+  // Create function signature
+  const std::string SubHandlerName = F->getName().str() + "subhandler";
+  FunctionType *FTy = FunctionType::get(
+      Type::getVoidTy(Ctx), {NativeCPUArgDescPtrType, StatePtrType}, false);
+  auto SubhFCallee = F->getParent()->getOrInsertFunction(SubHandlerName, FTy);
+  Function *SubhF = cast<Function>(SubhFCallee.getCallee());
+
+  // Emit function body, unpack kernel args
+  auto UsedIndexes = getUsedIndexes(F);
+  auto *KernelTy = F->getFunctionType();
+  // assert(UsedIndexes.size() + 1 == KernelTy->getNumParams() && "mismatch
+  // between number of params and used args");
+  IRBuilder<> Builder(Ctx);
+  BasicBlock *Block = BasicBlock::Create(Ctx, "entry", SubhF);
+  Builder.SetInsertPoint(Block);
+  unsigned NumArgs = UsedIndexes.size();
+  auto *BaseNativeCPUArg = SubhF->getArg(0);
+  SmallVector<Value *, 5> KernelArgs;
+  for (unsigned I = 0; I < NumArgs; I++) {
+    auto *Arg = F->getArg(I);
+    auto UsedI = UsedIndexes[I];
+    // Load the correct NativeCPUDesc and load the pointer from it
+    auto *Addr = Builder.CreateGEP(NativeCPUArgDescType, BaseNativeCPUArg,
+                                   {Builder.getInt64(UsedI)});
+    auto *Load = Builder.CreateLoad(PointerType::getUnqual(Ctx), Addr);
+    if (Arg->getType()->isPointerTy()) {
+      // If the arg is a pointer, just use it
+      KernelArgs.push_back(Load);
+    } else {
+      // Otherwise, load the scalar value and use that
+      auto *Scalar = Builder.CreateLoad(Arg->getType(), Load);
+      KernelArgs.push_back(Scalar);
+    }
+  }
+
+  // Call the kernel
+  // Add the nativecpu state as arg
+  KernelArgs.push_back(SubhF->getArg(1));
+  Builder.CreateCall(KernelTy, F, KernelArgs);
+  Builder.CreateRetVoid();
+}
+
+// Clones the function and returns a new function with a new argument on type T
 // added as last argument
 Function *cloneFunctionAndAddParam(Function *OldF, Type *T) {
   auto *OldT = OldF->getFunctionType();
@@ -167,6 +249,14 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M,
     ModuleChanged |= true;
   }
 
+  StructType *NativeCPUArgDescType =
+      StructType::create({PointerType::getUnqual(M.getContext())});
+  for (auto &NewK : NewKernels) {
+    emitSubkernelForKernel(NewK, NativeCPUArgDescType, StatePtrType);
+    std::string NewName = NewK->getName().str() + "_NativeCPUKernel";
+    NewK->setName(NewName);
+  }
+
   // Then we iterate over all the supported builtins, find their uses and
   // replace them with calls to our Native CPU functions.
   for (auto &Entry : BuiltinNamesMap) {