Merge branch 'sycl' into hip-support-graph

.github/ISSUE_TEMPLATE/1-bug-report.yml

@@ -0,0 +1,42 @@
+name: Bug report
+description: File a bug report to help us improve
+labels: ["bug"]
+body:
+  - type: textarea
+    id: description
+    attributes:
+      label: Describe the bug
+      description: A clear and concise description of what the bug is.
+    validations:
+      required: true
+  - type: textarea
+    id: how-to-reproduce
+    attributes:
+      label: To reproduce
+      description: Please describe the steps to reproduce the behavior
+      placeholder: |
+        1. Include a code snippet that is as short as possible
+        2. Specify the command which should be used to compile the program
+        3. Specify the command which should be used to launch the program
+        4. Indicate what is wrong and what was expected
+    validations:
+      required: false
+  - type: textarea
+    id: environment
+    attributes:
+      label: Environment
+      description: Please complete the following information
+      placeholder: |
+        - OS: [e.g Windows/Linux]
+        - Target device and vendor: [e.g. Intel GPU]
+        - DPC++ version: [e.g. commit hash or output of `clang++ --version`]
+        - Dependencies version: [e.g. the output of `sycl-ls --verbose`]
+    validations:
+      required: false
+  - type: textarea
+    id: additional
+    attributes:
+      label: Additional context
+      description: Add any other context about the problem here
+    validations:
+      required: false

.github/ISSUE_TEMPLATE/2-bug-report-cuda.yml

@@ -0,0 +1,42 @@
+name: CUDA-specific bug report
+description: Please use this template in case of CUDA-specific issue
+labels: ["bug", "cuda"]
+body:
+  - type: textarea
+    id: description
+    attributes:
+      label: Describe the bug
+      description: A clear and concise description of what the bug is.
+    validations:
+      required: true
+  - type: textarea
+    id: how-to-reproduce
+    attributes:
+      label: To reproduce
+      description: Please describe the steps to reproduce the behavior
+      placeholder: |
+        1. Include code snippet as short as possible
+        2. Specify the command which should be used to compile the program
+        3. Specify the command which should be used to launch the program
+        4. Indicate what is wrong and what was expected
+    validations:
+      required: false
+  - type: textarea
+    id: environment
+    attributes:
+      label: Environment
+      description: Please complete the following information
+      placeholder: |
+        - OS: [e.g Windows/Linux]
+        - Target device and vendor: [e.g. Nvidia GPU]
+        - DPC++ version: [e.g. commit hash or output of `clang++ --version`]
+        - Dependencies version: [e.g. the output of `nvidia-smi` and `sycl-ls --verbose`]
+    validations:
+      required: false
+  - type: textarea
+    id: additional
+    attributes:
+      label: Additional context
+      description: Add any other context about the problem here
+    validations:
+      required: false

.github/ISSUE_TEMPLATE/3-bug-report-hip.yml

@@ -0,0 +1,42 @@
+name: HIP-specific bug report
+description: Please use this template in case of HIP-specific issue
+labels: ["bug", "hip"]
+body:
+  - type: textarea
+    id: description
+    attributes:
+      label: Describe the bug
+      description: A clear and concise description of what the bug is.
+    validations:
+      required: true
+  - type: textarea
+    id: how-to-reproduce
+    attributes:
+      label: To reproduce
+      description: Please describe the steps to reproduce the behavior
+      placeholder: |
+        1. Include code snippet as short as possible
+        2. Specify the command which should be used to compile the program
+        3. Specify the command which should be used to launch the program
+        4. Indicate what is wrong and what was expected
+    validations:
+      required: false
+  - type: textarea
+    id: environment
+    attributes:
+      label: Environment
+      description: Please complete the following information
+      placeholder: |
+        - OS: [e.g Windows/Linux]
+        - Target device and vendor: [e.g. AMD GPU]
+        - DPC++ version: [e.g. commit hash or output of `clang++ --version`]
+        - Dependencies version: [e.g. the output of `rocm-smi` and `sycl-ls --verbose`]
+    validations:
+      required: false
+  - type: textarea
+    id: additional
+    attributes:
+      label: Additional context
+      description: Add any other context about the problem here
+    validations:
+      required: false

.github/ISSUE_TEMPLATE/4-feature-request.yml

@@ -0,0 +1,39 @@
+name: Feature request
+description: Suggest an idea/improvement for this project
+labels: ["enhancement"]
+body:
+  - type: textarea
+    id: description
+    attributes:
+      label: Is your feature request related to a problem? Please describe
+      description: |
+        A clear and concise description of what the problem is.
+        Include reproducer or code/pseudo-code example.
+        Include specific environment details where problem occurs.
+    validations:
+      required: true
+  - type: textarea
+    id: desired-solution
+    attributes:
+      label: Describe the solution you would like
+      description: |
+        A clear and concise description of what you want to happen.
+    validations:
+      required: false
+  - type: textarea
+    id: alternative-solution
+    attributes:
+      label: Describe alternatives you have considered
+      description: |
+        A clear and concise description of any alternative solutions or features
+        you have considered.
+    validations:
+      required: false
+  - type: textarea
+    id: additional
+    attributes:
+      label: Additional context
+      description: |
+        Add any other context about the feature request here.
+    validations:
+      required: false

.github/ISSUE_TEMPLATE/config.yml

@@ -0,0 +1,4 @@
+contact_links:
+  - name: Ask community a question
+    url: https://github.com/intel/llvm/discussions/categories/q-a
+    about: Please use Q&A Discussions category instead of Issues to ask questions

.github/workflows/sycl-detect-changes.yml

@@ -6,9 +6,6 @@ on:
       filters:
         description: Matched filters
         value: ${{ jobs.need_check.outputs.filters }}
-      arc_tests:
-        description: Tests to run on Arc
-        value: ${{ jobs.need_check.outputs.arc_tests }}
 
 jobs:
   need_check:
@@ -19,7 +16,6 @@ jobs:
     timeout-minutes: 3
     outputs:
       filters: ${{ steps.result.outputs.result }}
-      arc_tests: ${{ steps.arc_tests.outputs.arc_tests }}
     steps:
       - name: Check file changes
         uses: dorny/paths-filter@0bc4621a3135347011ad047f9ecf449bf72ce2bd
@@ -89,12 +85,4 @@ jobs:
             return ["llvm", "llvm_spirv", "clang", "sycl_fusion", "xptifw", "libclc", "sycl", "ci", "esimd"];
 
       - run: echo '${{ steps.result.outputs.result }}'
-
-      - name: Set Arc tests
-        id: arc_tests
-        run: |
-          if [ "${{ contains(steps.result.outputs.result, 'esimd') }}" == "true" ]; then
-            echo 'arc_tests="(ESIMD|InvokeSimd|Matrix)/"' >> "$GITHUB_OUTPUT"
-          else
-            echo 'arc_tests="Matrix/"' >> "$GITHUB_OUTPUT"
-          fi
+

.github/workflows/sycl-linux-precommit.yml

@@ -47,9 +47,28 @@ jobs:
       build_image: "ghcr.io/intel/llvm/ubuntu2204_build:7ed894ab0acc8ff09262113fdb08940d22654a30"
       changes: ${{ needs.detect_changes.outputs.filters }}
 
-  test:
+  determine_arc_tests:
+    name: Decide which Arc tests to run
     needs: [build, detect_changes]
     if: ${{ always() && !cancelled() && needs.build.outputs.build_conclusion == 'success' }}
+    # Github's ubuntu-* runners are slow to allocate. Use our CUDA runner since
+    # we don't use it for anything right now.
+    runs-on: [Linux, build]
+    timeout-minutes: 3
+    outputs:
+      arc_tests: ${{ steps.arc_tests.outputs.arc_tests }}
+    steps:
+      - name: Determine Arc tests
+        id: arc_tests
+        run: |
+          if [ "${{ contains(needs.detect_changes.outputs.filters, 'esimd') }}" == "true" ]; then
+            echo 'arc_tests="(ESIMD|InvokeSimd|Matrix)/"' >> "$GITHUB_OUTPUT"
+          else
+            echo 'arc_tests="Matrix/"' >> "$GITHUB_OUTPUT"
+          fi
+  test:
+    needs: [build, detect_changes, determine_arc_tests]
+    if: ${{ always() && !cancelled() && needs.build.outputs.build_conclusion == 'success' }}
     strategy:
       fail-fast: false
       matrix:
@@ -76,7 +95,7 @@ jobs:
             reset_gpu: true
             install_drivers: ${{ contains(needs.detect_changes.outputs.filters, 'drivers') }}
             extra_lit_opts: --param matrix-xmx8=True --param gpu-intel-dg2=True
-            env: '{"LIT_FILTER":${{ needs.detect_changes.outputs.arc_tests }} }'
+            env: '{"LIT_FILTER":${{ needs.determine_arc_tests.outputs.arc_tests }} }'
     uses: ./.github/workflows/sycl-linux-run-tests.yml
     with:
       name: ${{ matrix.name }}

clang/include/clang/Basic/LangOptions.def

@@ -303,7 +303,7 @@ LANGOPT(SYCLEnableIntHeaderDiags, 1, 0, "Enable diagnostics that require the "
         "SYCL integration header")
 LANGOPT(SYCLAllowVirtualFunctions, 1, 0,
         "Allow virtual functions calls in code for SYCL device")
-LANGOPT(SYCLIsNativeCPU      , 1, 0, "Generate code for SYCL NativeCPU")
+LANGOPT(SYCLIsNativeCPU      , 1, 0, "Generate code for SYCL Native CPU")
 
 LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP")
 LANGOPT(OffloadUniformBlock, 1, 0, "Assume that kernels are launched with uniform block sizes (default true for CUDA/HIP and false otherwise)")

clang/lib/CodeGen/BackendUtil.cpp

@@ -962,6 +962,8 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
       PB.registerPipelineStartEPCallback([&](ModulePassManager &MPM,
                                              OptimizationLevel Level) {
         MPM.addPass(ESIMDVerifierPass(LangOpts.SYCLESIMDForceStatelessMem));
+        if (Level == OptimizationLevel::O0)
+          MPM.addPass(ESIMDRemoveOptnoneNoinlinePass());
         MPM.addPass(SYCLPropagateAspectsUsagePass(/*ExcludeAspects=*/{"fp64"}));
         MPM.addPass(SYCLPropagateJointMatrixUsagePass());
       });
@@ -1086,7 +1088,11 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
     }
 
     if (SYCLNativeCPUBackend) {
-      sycl::utils::addSYCLNativeCPUBackendPasses(MPM, MAM);
+      llvm::sycl::utils::addSYCLNativeCPUBackendPasses(MPM, MAM, Level);
+      // Run optimization passes after all the changes we made to the kernels.
+      // Todo: maybe we could find a set of relevant passes instead of re-running
+      // the full optimization pipeline.
+      MPM.addPass(PB.buildPerModuleDefaultPipeline(Level));
     }
     if (LangOpts.SYCLIsDevice) {
       MPM.addPass(SYCLMutatePrintfAddrspacePass());

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5096,7 +5096,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   bool IsFPGASYCLOffloadDevice =
       IsSYCLOffloadDevice &&
       Triple.getSubArch() == llvm::Triple::SPIRSubArch_fpga;
-  const bool IsSYCLNativeCPU = isSYCLNativeCPU(TC, C.getDefaultToolChain());
+  const bool IsSYCLNativeCPU = isSYCLNativeCPU(TC);
 
   // Perform the SYCL host compilation using an external compiler if the user
   // requested.
@@ -9851,7 +9851,7 @@ void OffloadWrapper::ConstructJob(Compilation &C, const JobAction &JA,
       TargetTripleOpt = ("llvm_" + TargetTripleOpt).str();
     }
 
-    const bool IsSYCLNativeCPU = isSYCLNativeCPU(TC, C.getDefaultToolChain());
+    const bool IsSYCLNativeCPU = isSYCLNativeCPU(TC);
     if (IsSYCLNativeCPU) {
       TargetTripleOpt = "native_cpu";
     }
@@ -10316,8 +10316,11 @@ static void addArgs(ArgStringList &DstArgs, const llvm::opt::ArgList &Alloc,
 // Partially copied from clang/lib/Frontend/CompilerInvocation.cpp
 static std::string getSYCLPostLinkOptimizationLevel(const ArgList &Args) {
   if (Arg *A = Args.getLastArg(options::OPT_O_Group)) {
+    // Pass -O2 when the user passes -O0 due to IGC
+    // debugging limitation. Note this only effects
+    // ESIMD code.
     if (A->getOption().matches(options::OPT_O0))
-      return "-O0";
+      return "-O2";
 
     if (A->getOption().matches(options::OPT_Ofast))
       return "-O3";
@@ -10381,7 +10384,7 @@ void SYCLPostLink::ConstructJob(Compilation &C, const JobAction &JA,
   if (!TCArgs.hasFlag(options::OPT_fno_sycl_remove_unused_external_funcs,
                       options::OPT_fsycl_remove_unused_external_funcs, false) &&
       !T.isNVPTX() && !T.isAMDGPU() &&
-      !isSYCLNativeCPU(getToolChain(), C.getDefaultToolChain()))
+      !isSYCLNativeCPU(getToolChain()))
     addArgs(CmdArgs, TCArgs, {"-emit-only-kernels-as-entry-points"});
 
   // OPT_fsycl_device_code_split is not checked as it is an alias to

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -401,8 +401,7 @@ const char *SYCL::Linker::constructLLVMLinkCommand(
   // instead of the original object.
   if (JA.isDeviceOffloading(Action::OFK_SYCL)) {
     bool IsRDC = !shouldDoPerObjectFileLinking(C);
-    const bool IsSYCLNativeCPU = isSYCLNativeCPU(
-        this->getToolChain(), *C.getSingleOffloadToolChain<Action::OFK_Host>());
+    const bool IsSYCLNativeCPU = isSYCLNativeCPU(this->getToolChain());
     auto isNoRDCDeviceCodeLink = [&](const InputInfo &II) {
       if (IsRDC)
         return false;