Merge branch 'sycl' into e2e-split

Author: ayylol

clang/lib/Driver/ToolChains/Clang.cpp

@@ -11471,6 +11471,12 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
     if (Args.hasArg(options::OPT_fsycl_embed_ir))
       CmdArgs.push_back(Args.MakeArgString("-sycl-embed-ir"));
 
+    if (Args.hasFlag(options::OPT_fsycl_allow_device_image_dependencies,
+                     options::OPT_fno_sycl_allow_device_image_dependencies,
+                     false))
+      CmdArgs.push_back(
+          Args.MakeArgString("-sycl-allow-device-image-dependencies"));
+
     // Formulate and add any offload-wrapper and AOT specific options. These
     // are additional options passed in via -Xsycl-target-linker and
     // -Xsycl-target-backend.

clang/lib/Driver/ToolChains/SYCL.h

@@ -1,4 +1,4 @@
-//===--- SYCL.h - SYCL ToolChain Implementations -----------------*- C++ -*-===//
+//===--- SYCL.h - SYCL ToolChain Implementations ----------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -158,8 +158,8 @@ SmallVector<std::string, 8> getDeviceLibraries(const Compilation &C,
                                                bool IsSpirvAOT);
 
 // Populates the SYCL device traits macros.
-void populateSYCLDeviceTraitsMacrosArgs(Compilation &C,
-    const llvm::opt::ArgList &Args,
+void populateSYCLDeviceTraitsMacrosArgs(
+    Compilation &C, const llvm::opt::ArgList &Args,
     const SmallVectorImpl<std::pair<const ToolChain *, StringRef>> &Targets);
 
 bool shouldDoPerObjectFileLinking(const Compilation &C);
@@ -180,11 +180,11 @@ class LLVM_LIBRARY_VISIBILITY Linker : public Tool {
 private:
   /// \return llvm-link output file name.
   const char *constructLLVMLinkCommand(Compilation &C, const JobAction &JA,
-                             const InputInfo &Output,
-                             const llvm::opt::ArgList &Args,
-                             llvm::StringRef SubArchName,
-                             llvm::StringRef OutputFilePrefix,
-                             const InputInfoList &InputFiles) const;
+                                       const InputInfo &Output,
+                                       const llvm::opt::ArgList &Args,
+                                       llvm::StringRef SubArchName,
+                                       llvm::StringRef OutputFilePrefix,
+                                       const InputInfoList &InputFiles) const;
 };
 
 /// Directly call FPGA Compiler and Linker
@@ -241,7 +241,7 @@ template <auto GPUArh> std::optional<StringRef> isGPUTarget(StringRef Target) {
   if (Target.starts_with(GPUArh)) {
     return resolveGenDevice(Target);
   }
-  return  std::nullopt;
+  return std::nullopt;
 }
 
 } // end namespace gen
@@ -280,9 +280,10 @@ class LLVM_LIBRARY_VISIBILITY SYCLToolChain : public ToolChain {
   llvm::opt::DerivedArgList *
   TranslateArgs(const llvm::opt::DerivedArgList &Args, StringRef BoundArch,
                 Action::OffloadKind DeviceOffloadKind) const override;
-  void addClangTargetOptions(const llvm::opt::ArgList &DriverArgs,
-                         llvm::opt::ArgStringList &CC1Args,
-                         Action::OffloadKind DeviceOffloadKind) const override;
+  void
+  addClangTargetOptions(const llvm::opt::ArgList &DriverArgs,
+                        llvm::opt::ArgStringList &CC1Args,
+                        Action::OffloadKind DeviceOffloadKind) const override;
   void AddImpliedTargetArgs(const llvm::Triple &Triple,
                             const llvm::opt::ArgList &Args,
                             llvm::opt::ArgStringList &CmdArgs,
@@ -324,7 +325,8 @@ class LLVM_LIBRARY_VISIBILITY SYCLToolChain : public ToolChain {
   CXXStdlibType GetCXXStdlibType(const llvm::opt::ArgList &Args) const override;
   void AddSYCLIncludeArgs(const llvm::opt::ArgList &DriverArgs,
                           llvm::opt::ArgStringList &CC1Args) const override;
-  void AddClangSystemIncludeArgs(const llvm::opt::ArgList &DriverArgs,
+  void
+  AddClangSystemIncludeArgs(const llvm::opt::ArgList &DriverArgs,
                             llvm::opt::ArgStringList &CC1Args) const override;
   void AddClangCXXStdlibIncludeArgs(
       const llvm::opt::ArgList &Args,
@@ -357,7 +359,8 @@ inline bool isSYCLNativeCPU(const llvm::opt::ArgList &Args) {
   return false;
 }
 
-inline bool isSYCLNativeCPU(const llvm::Triple &HostT, const llvm::Triple &DevT) {
+inline bool isSYCLNativeCPU(const llvm::Triple &HostT,
+                            const llvm::Triple &DevT) {
   return HostT == DevT;
 }
 

clang/lib/Sema/SemaSYCL.cpp

@@ -3950,13 +3950,26 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   }
 
   // Default inits the type, then calls the init-method in the body.
+  // A type may not have a public default constructor as per its spec so
+  // typically if this is the case the default constructor will be private and
+  // in such cases we must manually override the access specifier from private
+  // to public just for the duration of this default initialization.
+  // TODO: Revisit this approach once https://github.com/intel/llvm/issues/16061
+  // is closed.
   bool handleSpecialType(FieldDecl *FD, QualType Ty) {
+    const auto *RecordDecl = Ty->getAsCXXRecordDecl();
+    AccessSpecifier DefaultConstructorAccess;
+    auto DefaultConstructor =
+        std::find_if(RecordDecl->ctor_begin(), RecordDecl->ctor_end(),
+                     [](auto it) { return it->isDefaultConstructor(); });
+    DefaultConstructorAccess = DefaultConstructor->getAccess();
+    DefaultConstructor->setAccess(AS_public);
+
     addFieldInit(FD, Ty, std::nullopt,
                  InitializationKind::CreateDefault(KernelCallerSrcLoc));
-
+    DefaultConstructor->setAccess(DefaultConstructorAccess);
     addFieldMemberExpr(FD, Ty);
 
-    const auto *RecordDecl = Ty->getAsCXXRecordDecl();
     createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts);
     CXXMethodDecl *FinalizeMethod =
         getMethodByName(RecordDecl, FinalizeMethodName);
@@ -3970,9 +3983,17 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   }
 
   bool handleSpecialType(const CXXBaseSpecifier &BS, QualType Ty) {
-    const auto *RecordDecl = Ty->getAsCXXRecordDecl();
+    const auto *BaseRecordDecl = BS.getType()->getAsCXXRecordDecl();
+    AccessSpecifier DefaultConstructorAccess;
+    auto DefaultConstructor =
+        std::find_if(BaseRecordDecl->ctor_begin(), BaseRecordDecl->ctor_end(),
+                     [](auto it) { return it->isDefaultConstructor(); });
+    DefaultConstructorAccess = DefaultConstructor->getAccess();
+    DefaultConstructor->setAccess(AS_public);
+
     addBaseInit(BS, Ty, InitializationKind::CreateDefault(KernelCallerSrcLoc));
-    createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts);
+    DefaultConstructor->setAccess(DefaultConstructorAccess);
+    createSpecialMethodCall(BaseRecordDecl, getInitMethodName(), BodyStmts);
     return true;
   }
 
@@ -4669,16 +4690,21 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
   bool handleSyclSpecialType(const CXXRecordDecl *RD,
                              const CXXBaseSpecifier &BC,
                              QualType FieldTy) final {
-    const auto *AccTy =
-        cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
-    assert(AccTy->getTemplateArgs().size() >= 2 &&
-           "Incorrect template args for Accessor Type");
-    int Dims = static_cast<int>(
-        AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
-    int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
-    Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
-                        CurOffset +
-                            offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
+    if (isSyclAccessorType(FieldTy)) {
+      const auto *AccTy =
+          cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
+      assert(AccTy->getTemplateArgs().size() >= 2 &&
+             "Incorrect template args for Accessor Type");
+      int Dims = static_cast<int>(
+          AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
+      int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
+      Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
+                          CurOffset +
+                              offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
+    } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
+      addParam(FieldTy, SYCLIntegrationHeader::kind_work_group_memory,
+               offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
+    }
     return true;
   }
 

clang/test/Driver/sycl-offload-new-driver.c

@@ -195,3 +195,19 @@
 // RUN:          --offload-new-driver 2>&1 \
 // RUN:   | FileCheck -check-prefix NVPTX_CUDA_PATH %s
 // NVPTX_CUDA_PATH: clang-linker-wrapper{{.*}} "--cuda-path={{.*}}Inputs/CUDA_80/usr/local/cuda"
+
+/// Check for -sycl-allow-device-image-dependencies transmission to clang-linker-wrapper tool
+// RUN: %clangxx -fsycl -###  --offload-new-driver \
+// RUN:          -fsycl-allow-device-image-dependencies %s 2>&1 \
+// RUN:  | FileCheck -check-prefix CHECK_DYNAMIC_LINKING %s
+// CHECK_DYNAMIC_LINKING: clang-linker-wrapper{{.*}} "-sycl-allow-device-image-dependencies"
+
+/// Check that -sycl-allow-device-image-dependencies is not passed to clang-linker-wrapper tool
+// RUN: %clangxx -fsycl -### --offload-new-driver \
+// RUN:          -fno-sycl-allow-device-image-dependencies %s 2>&1 \
+// RUN:  | FileCheck -check-prefix CHECK_NO_DYNAMIC_LINKING %s
+
+/// Check that -sycl-allow-device-image-dependencies is not passed to clang-linker-wrapper tool
+// RUN: %clangxx -fsycl -### --offload-new-driver %s 2>&1 \
+// RUN:  | FileCheck -check-prefix CHECK_NO_DYNAMIC_LINKING %s
+// CHECK_NO_DYNAMIC_LINKING-NOT: clang-linker-wrapper{{.*}} "-sycl-allow-device-image-dependencies"

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -448,6 +448,23 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<T, AS, access::decorated::legacy> {
   pointer_t m_Pointer;
 };
 
+// Dummy implementation of work_group_memory for use in SemaSYCL tests.
+template <typename DataT>
+class __attribute__((sycl_special_class))
+__SYCL_TYPE(work_group_memory) work_group_memory {
+
+// Default constructor for objects later initialized with __init member.
+  work_group_memory() = default;
+
+public:
+  work_group_memory(handler &CGH) {}
+
+  void __init(__attribute((opencl_local)) DataT *Ptr) { this->Ptr = Ptr; }
+  void use() const {}
+private:
+  __attribute((opencl_local)) DataT *Ptr;
+};
+
 namespace ext {
 namespace oneapi {
 namespace experimental {

clang/test/SemaSYCL/work_group_memory_inheritance.cpp

@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -sycl-std=2020 %s | FileCheck %s
+
+// Check that AST is correctly generated for kernel arguments that inherit from work group memory.
+
+#include "sycl.hpp"
+
+sycl::queue myQueue;
+
+struct WorkGroupMemoryDerived :
+                         sycl::work_group_memory<int> {
+};
+
+int main() {
+  myQueue.submit([&](sycl::handler &h) {
+    WorkGroupMemoryDerived DerivedObject{ h };
+    h.parallel_for<class kernel>([=] {
+          DerivedObject.use();
+        });
+  });
+  return 0;
+}
+
+// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (__local int *)'
+// CHECK-NEXT: ParmVarDecl {{.*}}used _arg__base '__local int *'
+// CHECK-NEXT: CompoundStmt {{.*}}
+// CHECK-NEXT: DeclStmt {{.*}}
+// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel {{.*}} cinit
+// CHECK-NEXT: InitListExpr {{.*}}
+// CHECK-NEXT: InitListExpr {{.*}} 'WorkGroupMemoryDerived'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::work_group_memory<int>' 'void () noexcept'
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__local int *)' lvalue .__init {{.*}}
+// CHECK-NEXT: MemberExpr {{.*}} 'WorkGroupMemoryDerived' lvalue .DerivedObject
+// CHECK-NEXT: DeclRefExpr {{.*}} lvalue Var {{.*}} '__SYCLKernel'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__local int *' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__local int *' lvalue ParmVar {{.*}} '_arg__base' '__local int *'
+// CHECK-NEXT: CompoundStmt {{.*}}
+// CHECK-NEXT: CXXOperatorCallExpr {{.*}} 'void' '()'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'auto (*)() const -> void' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr {{.*}}'auto () const -> void' lvalue CXXMethod {{.*}} 'operator()' 'auto () const -> void'
+// CHECK-NEXT: ImplicitCastExpr {{.*}}
+// CHECK-NEXT: DeclRefExpr {{.*}}lvalue Var {{.*}} '__SYCLKernel'
+

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

@@ -676,6 +676,7 @@ getTripleBasedSYCLPostLinkOpts(const ArgList &Args,
   if ((!Args.hasFlag(OPT_no_sycl_remove_unused_external_funcs,
                      OPT_sycl_remove_unused_external_funcs, false) &&
        !SYCLNativeCPU) &&
+      !Args.hasArg(OPT_sycl_allow_device_image_dependencies) &&
       !Triple.isNVPTX() && !Triple.isAMDGPU())
     PostLinkArgs.push_back("-emit-only-kernels-as-entry-points");
 
@@ -1111,8 +1112,10 @@ wrapSYCLBinariesFromFile(std::vector<module_split::SplitModule> &SplitModules,
     if (!MBOrDesc)
       return createFileError(SI.ModuleFilePath, MBOrDesc.getError());
 
-    StringRef ImageTarget = IsEmbeddedIR ? StringRef(EmbeddedIRTarget) : StringRef(RegularTarget);
-    Images.emplace_back(std::move(*MBOrDesc), SI.Properties, SI.Symbols, ImageTarget);
+    StringRef ImageTarget =
+        IsEmbeddedIR ? StringRef(EmbeddedIRTarget) : StringRef(RegularTarget);
+    Images.emplace_back(std::move(*MBOrDesc), SI.Properties, SI.Symbols,
+                        ImageTarget);
   }
 
   LLVMContext C;
@@ -1197,7 +1200,8 @@ static Expected<StringRef> runCompile(StringRef &InputFile,
 static Expected<StringRef>
 runWrapperAndCompile(std::vector<module_split::SplitModule> &SplitModules,
                      const ArgList &Args, bool IsEmbeddedIR = false) {
-  auto OutputFile = sycl::wrapSYCLBinariesFromFile(SplitModules, Args, IsEmbeddedIR);
+  auto OutputFile =
+      sycl::wrapSYCLBinariesFromFile(SplitModules, Args, IsEmbeddedIR);
   if (!OutputFile)
     return OutputFile.takeError();
   // call to clang
@@ -2416,8 +2420,8 @@ Expected<SmallVector<StringRef>> linkAndWrapDeviceFiles(
         // of sycl-post-link (filetable referencing LLVM Bitcode + symbols)
         // through the offload wrapper and link the resulting object to the
         // application.
-        auto OutputFile =
-            sycl::runWrapperAndCompile(SplitModules, LinkerArgs, /* IsEmbeddedIR */ true);
+        auto OutputFile = sycl::runWrapperAndCompile(SplitModules, LinkerArgs,
+                                                     /* IsEmbeddedIR */ true);
         if (!OutputFile)
           return OutputFile.takeError();
         WrappedOutput.push_back(*OutputFile);

clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td

@@ -243,3 +243,8 @@ Flags<[WrapperOnlyOption]>,  HelpText<"Embed LLVM IR for runtime kernel fusion">
 def sycl_dump_device_code_EQ : Joined<["--", "-"], "sycl-dump-device-code=">,
   Flags<[WrapperOnlyOption]>,
   HelpText<"Path to the folder where the tool dumps SPIR-V device code. Other formats aren't dumped.">;
+
+// Options to enable/disable device dynamic linking.
+def sycl_allow_device_image_dependencies : Flag<["--", "-"], "sycl-allow-device-image-dependencies">,
+  Flags<[WrapperOnlyOption, HelpHidden]>,
+  HelpText<"Allow dependencies between device code images">;

libdevice/nativecpu_utils.cpp

@@ -17,7 +17,7 @@
 #include "device.h"
 #include <cstdint>
 #include <sycl/__spirv/spirv_ops.hpp>
-#include <sycl/types.hpp>
+#include <sycl/vector.hpp>
 
 // including state definition from Native CPU UR adapter
 #include "nativecpu_state.hpp"

sycl-jit/jit-compiler/CMakeLists.txt

@@ -19,10 +19,12 @@ add_llvm_library(sycl-jit
    BitReader
    Core
    Support
+   Option
    Analysis
    IPO
    TransformUtils
    Passes
+   IRReader
    Linker
    ScalarOpts
    InstCombine