Merge branch 'sycl' into run-mode

Author: ayylol

.github/workflows/sycl-containers.yaml

@@ -51,6 +51,13 @@ jobs:
             file: ubuntu2204_intel_drivers
             tag: unstable
             build_args: "use_latest=true"
+          - name: Build + Intel Drivers Ubuntu 22.04 Docker image
+            file: ubuntu2204_intel_drivers
+            tag: alldeps
+            build_args: |
+              base_image=ghcr.io/intel/llvm/ubuntu2204_build
+              base_tag=latest
+              use_latest=false
     steps:
       - name: Checkout
         uses: actions/checkout@v4

.github/workflows/sycl-linux-run-tests.yml

@@ -185,8 +185,8 @@ jobs:
       with:
         path: khronos_sycl_cts
         repository: 'KhronosGroup/SYCL-CTS'
-        ref: 'SYCL-2020'
-        default_branch: 'SYCL-2020'
+        ref: 'main'
+        default_branch: 'main'
         cache_path: "/__w/repo_cache/"
     - name: SYCL CTS GIT submodules init
       if: inputs.tests_selector == 'cts'

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5194,7 +5194,10 @@ void Clang::ConstructHostCompilerJob(Compilation &C, const JobAction &JA,
   if (HostCompilerDefArg) {
     ExecPath = HostCompilerDefArg->getValue();
     if (!ExecPath.empty() && ExecPath == llvm::sys::path::stem(ExecPath))
-      ExecPath = TC.GetProgramPath(ExecPath.c_str());
+      // Use PATH to find executable passed in from -fsycl-host-compiler.
+      if (llvm::ErrorOr<std::string> Prog =
+              llvm::sys::findProgramByName(ExecPath))
+        ExecPath = *Prog;
   }
 
   // Add any user-specified arguments.

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -1504,18 +1504,13 @@ static void parseTargetOpts(StringRef ArgString, const llvm::opt::ArgList &Args,
 void SYCLToolChain::TranslateGPUTargetOpt(const llvm::opt::ArgList &Args,
                                           llvm::opt::ArgStringList &CmdArgs,
                                           OptSpecifier Opt_EQ) const {
-  for (auto *A : Args) {
-    if (A->getOption().matches(Opt_EQ)) {
-      if (auto GpuDevice =
-              tools::SYCL::gen::isGPUTarget<tools::SYCL::gen::AmdGPU>(
-                  A->getValue())) {
-        StringRef ArgString;
-        SmallString<64> OffloadArch("--offload-arch=");
-        OffloadArch += GpuDevice->data();
-        ArgString = OffloadArch;
-        parseTargetOpts(ArgString, Args, CmdArgs);
-        A->claim();
-      }
+  if (const Arg *TargetArg = Args.getLastArg(Opt_EQ)) {
+    StringRef Val = TargetArg->getValue();
+    if (auto GpuDevice =
+            tools::SYCL::gen::isGPUTarget<tools::SYCL::gen::AmdGPU>(Val)) {
+      SmallString<64> OffloadArch("--offload-arch=");
+      OffloadArch += GpuDevice->data();
+      parseTargetOpts(OffloadArch, Args, CmdArgs);
     }
   }
 }

clang/test/Driver/sycl-host-compiler-old-model.cpp

@@ -90,3 +90,13 @@
 // CHECK_SAVE_TEMPS-NEXT: g++{{.*}} "[[APPEND_CPP]]" "-c"
 // CHECK_SAVE_TEMPS-SAME: "-o" "[[HOST_OBJ:.+\.o]]"
 // CHECK_SAVE_TEMPS-NEXT: clang-offload-bundler{{.*}} "-input=[[DEVICE_BC]]" "-input=[[HOST_OBJ]]"
+
+/// Test to verify binary from PATH is used
+// RUN: rm -rf %t && mkdir -p %t/test_path
+// RUN: touch %t/test_path/clang++ && chmod +x %t/test_path/clang++
+// RUN: env "PATH=%t/test_path%{pathsep}%PATH%" \
+// RUN: %clangxx -### -fsycl -fsycl-host-compiler=clang++ \
+// RUN:   -fsycl-host-compiler-options=-DDUMMY_OPT --no-offload-new-driver \
+// RUN:   %s 2>&1 \
+// RUN: | FileCheck -check-prefix=PATH_CHECK %s
+// PATH_CHECK: {{(/|\\\\)}}test_path{{(/|\\\\)}}clang++{{.*}} "-DDUMMY_OPT"

clang/test/Driver/sycl-offload.cpp

@@ -104,3 +104,9 @@
 // CHECK_FSYCL_FP64_CONV_EMU_WIN-NOT: clang{{.*}} "-cc1" "-triple x86_64-unknown-linux-gnu" {{.*}} "-fsycl-fp64-conv-emu"
 // CHECK_FSYCL_FP64_CONV_EMU_WIN-DAG: clang{{.*}} "-cc1" "-triple" "spir64_gen{{.*}}" "-fsycl-fp64-conv-emu"
 // CHECK_FSYCL_FP64_CONV_EMU_WIN-DAG: clang-offload-packager{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch=,kind=sycl,compile-opts={{.*}}-options -ze-fp64-gen-conv-emu{{.*}}"
+
+/// Compilation checks to make sure an early empty -fsycl-targets does not
+/// crash.
+// RUN: %clangxx -### -fsycl -fsycl-targets= -fsycl-targets=spir64 %s 2>&1 \
+// RUN:  | FileCheck %s -check-prefix=CHECK_SPIR64
+// CHECK_SPIR64: clang{{.*}} "-cc1" "-triple" "spir64-unknown-unknown"{{.*}} "-fsycl-is-device"

llvm/lib/SYCLLowerIR/LowerWGScope.cpp

@@ -1001,6 +1001,8 @@ Instruction *spirv::genWGBarrier(Instruction &Before, const Triple &TT) {
   FunctionCallee FC =
       M.getOrInsertFunction(Name, Attr, RetTy, ScopeTy, ScopeTy, SemanticsTy);
   assert(FC.getCallee() && "spirv intrinsic creation failed");
+  if (TT.isSPIROrSPIRV())
+    cast<Function>(FC.getCallee())->setCallingConv(CallingConv::SPIR_FUNC);
 
   IRBuilder<> Bld(Ctx);
   Bld.SetInsertPoint(&Before);
@@ -1011,5 +1013,7 @@ Instruction *spirv::genWGBarrier(Instruction &Before, const Triple &TT) {
                    asUInt(spirv::MemorySemantics::WorkgroupMemory));
   auto BarrierCall = Bld.CreateCall(FC, {ArgExec, ArgMem, ArgSema});
   BarrierCall->addFnAttr(llvm::Attribute::Convergent);
+  if (TT.isSPIROrSPIRV())
+    BarrierCall->setCallingConv(CallingConv::SPIR_FUNC);
   return BarrierCall;
 }

llvm/lib/SYCLLowerIR/SYCLJointMatrixTransform.cpp

@@ -22,16 +22,43 @@ namespace {
 static constexpr char ACCESS_CHAIN[] = "_Z19__spirv_AccessChain";
 static constexpr char MATRIX_TYPE[] = "spirv.CooperativeMatrixKHR";
 
-// This routine extracts spirv.CooperativeMatrixKHR target extension type
-// from sycl::joint_matrix class object if it's used in __spirv_AccessChain
-// function call. It's necessary because otherwise OpAccessChain indices would
-// be wrong.
+// This function finds all calls to __spirv_AccessChain function and transforms
+// its users and operands to make LLVM IR more SPIR-V friendly.
 bool transformAccessChain(Function *F) {
   bool ModuleChanged = false;
   for (auto I : F->users()) {
     auto *CI = dyn_cast<CallInst>(I);
     if (!CI)
       continue;
+
+    // This is a W/A for bfloat16 and tf32 types - they are represented in SYCL
+    // as structures with int16/float storages. It means, that in LLVM IR
+    // user of CallInst to __spirv_AccessChain function would be not load/store
+    // instruction, but a zero GEP. This zero GEP is no-op, but can confuse a
+    // SPIR-V consumer, so lets remove it here.
+    auto *Unique = CI->getUniqueUndroppableUser();
+    if (auto *GEP = dyn_cast_or_null<GetElementPtrInst>(Unique)) {
+      if (GEP->hasAllZeroIndices()) {
+        GEP->replaceAllUsesWith(CI);
+        GEP->dropAllReferences();
+        GEP->eraseFromParent();
+      }
+    }
+
+    // It can happen that the optimizer can remove duplicated or dead uses
+    // of CallInst to __spirv_AccessChain function. But it can't remove
+    // __spirv_AccessChain call itself as it's a call to external function.
+    // Lets clean such calls.
+    if (CI->getNumUses() == 0) {
+      CI->dropAllReferences();
+      CI->eraseFromParent();
+      continue;
+    }
+
+    // This routine extracts spirv.CooperativeMatrixKHR target extension type
+    // from sycl::joint_matrix class object if it's used in __spirv_AccessChain
+    // function call. It's necessary because otherwise OpAccessChain indices
+    // would be wrong.
     Instruction *Ptr =
         dyn_cast<Instruction>(CI->getArgOperand(0)->stripPointerCasts());
     if (!Ptr || !isa<AllocaInst>(Ptr))

llvm/lib/SYCLLowerIR/SYCLVirtualFunctionsAnalysis.cpp

@@ -71,37 +71,23 @@ void checkKernel(const Function *F, const CallGraphTy &CG) {
 void computeFunctionToKernelsMappingImpl(Function *Kernel, const Function *F,
                                          const CallGraphTy &CG,
                                          FuncToFuncMapTy &Mapping) {
+  Mapping[F].insert(Kernel);
   CallGraphTy::const_iterator It = CG.find(F);
   // It could be that the function itself is a leaf and doesn't call anything
   if (It == CG.end())
     return;
 
-  Mapping[F].insert(Kernel);
-
   const SmallPtrSet<Value *, 8> &Callees = It->getSecond();
   for (const Value *V : Callees)
     if (auto *Callee = dyn_cast<Function>(V))
       computeFunctionToKernelsMappingImpl(Kernel, Callee, CG, Mapping);
 }
 
+// Compute a map from functions used by a kernel to that kernel.
+// For simplicity we also consider a kernel to be using itself.
 void computeFunctionToKernelsMapping(Function *Kernel, const CallGraphTy &CG,
                                      FuncToFuncMapTy &Mapping) {
-  // For simplicity we also consider a kernel to be using itself
-  Mapping[Kernel].insert(Kernel);
-
-  CallGraphTy::const_iterator It = CG.find(Kernel);
-  // It could be that the kernel doesn't call anything
-  if (It == CG.end())
-    return;
-
-  const SmallPtrSet<Value *, 8> &Callees = It->getSecond();
-  for (const Value *V : Callees) {
-    auto *Callee = dyn_cast<Function>(V);
-    if (!Callee)
-      continue;
-    Mapping[Callee].insert(Kernel);
-    computeFunctionToKernelsMappingImpl(Kernel, Callee, CG, Mapping);
-  }
+  computeFunctionToKernelsMappingImpl(Kernel, Kernel, CG, Mapping);
 }
 
 void collectVTablesThatUseFunction(

llvm/lib/SYCLLowerIR/SpecConstants.cpp

@@ -475,6 +475,7 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName,
   auto *FT = FunctionType::get(RetTy, ArgTys, false /*isVarArg*/);
   std::string FunctionName = mangleFuncItanium(BaseFunctionName, FT);
   Module *M = InsertBefore->getFunction()->getParent();
+  bool IsSPIROrSPIRV = llvm::Triple(M->getTargetTriple()).isSPIROrSPIRV();
 
   if (RetTy->isIntegerTy(1)) {
     assert(ArgTys.size() == 2 && "Expected a scalar spec constant");
@@ -500,6 +501,11 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName,
 
       auto *Call =
           CallInst::Create(NewFT, NewFC.getCallee(), Args, "", InsertBefore);
+      if (IsSPIROrSPIRV) {
+        cast<Function>(NewFC.getCallee())
+            ->setCallingConv(CallingConv::SPIR_FUNC);
+        Call->setCallingConv(CallingConv::SPIR_FUNC);
+      }
       return CastInst::CreateTruncOrBitCast(Call, RetTy, "tobool",
                                             InsertBefore);
     }
@@ -520,7 +526,12 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName,
   // types? Is it necessary?
 
   FunctionCallee FC = M->getOrInsertFunction(FunctionName, FT);
-  return CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore);
+  auto *Call = CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore);
+  if (IsSPIROrSPIRV) {
+    cast<Function>(FC.getCallee())->setCallingConv(CallingConv::SPIR_FUNC);
+    Call->setCallingConv(CallingConv::SPIR_FUNC);
+  }
+  return Call;
 }
 
 Instruction *emitSpecConstant(unsigned NumericID, Type *Ty,