Merge from 'sycl' to 'sycl-web'

Author: iclsrc

clang/lib/CodeGen/CGCall.cpp

@@ -5688,10 +5688,11 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
     if (!getLangOpts().FPAccuracyFuncMap.empty() ||
         !getLangOpts().FPAccuracyVal.empty()) {
       const auto *FD = dyn_cast_if_present<FunctionDecl>(TargetDecl);
-      assert(FD && "expecting a function");
-      CI = EmitFPBuiltinIndirectCall(IRFuncTy, IRCallArgs, CalleePtr, FD);
-      if (CI)
-        return RValue::get(CI);
+      if (FD) {
+        CI = EmitFPBuiltinIndirectCall(IRFuncTy, IRCallArgs, CalleePtr, FD);
+        if (CI)
+          return RValue::get(CI);
+      }
     }
     CI = Builder.CreateCall(IRFuncTy, CalleePtr, IRCallArgs, BundleList);
   } else {

sycl/doc/design/KernelFusionJIT.md

@@ -2,8 +2,7 @@
 
 ## Context
 
-To support the [user driven kernel fusion extension](https://github.com/intel/llvm/pull/7098) (a presentation can be found [here](https://github.com/oneapi-src/oneAPI-tab/blob/main/tab-dpcpp-onedpl/presentations/oneAPI-TAB-20220727-Kernel-Fusion.pdf)).
-Currently, only targets able to consume SPIR-V are supported.
+To support the [user driven kernel fusion extension](https://github.com/intel/llvm/pull/7098) (a presentation can be found [here](https://github.com/oneapi-src/oneAPI-tab/blob/main/language/presentations/oneAPI-TAB-20220727-Kernel-Fusion.pdf)).
 
 The basic workflow is shown in the diagram below
 
@@ -284,3 +283,13 @@ flag during static compilation.
 During the fusion process at runtime, the JIT will load the LLVM IR and
 finalize the fused kernel to the final target. More information is available
 [here](./CompilerAndRuntimeDesign.md#kernel-fusion-support).
+
+### Unsupported SYCL constructs
+
+The following SYCL API constructs are currently not officially supported for
+kernel fusion and should be considered untested/unsupported: 
+
+- Reductions
+- `sycl::stream`
+- Specialization constants and `sycl::kernel_handler`
+- Images (`sycl::unsampled_image` and `sycl::sampled_image`)

sycl/doc/extensions/experimental/sycl_ext_codeplay_kernel_fusion.asciidoc

@@ -73,12 +73,6 @@ on an API for kernel fusion in SYCL to inform a future extension proposal
 addressing the mentioned problems. 
 ====
 
-[NOTE]
-====
-This extension is currently being implemented in {dpcpp} only for kernels in
-SPIRV format.
-====
-
 == Overview
 
 Every kernel launch in SYCL carries an overhead due to memory traffic and device

sycl/source/detail/scheduler/commands.cpp

@@ -2736,7 +2736,7 @@ pi_int32 ExecCGCommand::enqueueImpCommandBuffer() {
         Req->MDims, Req->MAccessRange,
         /*DstOffset=*/{0, 0, 0}, Req->MElemSize, std::move(MSyncPointDeps),
         &OutSyncPoint);
-
+    MEvent->setSyncPoint(OutSyncPoint);
     return PI_SUCCESS;
   }
   case CG::CGTYPE::CopyPtrToAcc: {
@@ -2750,7 +2750,7 @@ pi_int32 ExecCGCommand::enqueueImpCommandBuffer() {
         /*SrcOffset*/ {0, 0, 0}, Req->MElemSize, AllocaCmd->getMemAllocation(),
         Req->MDims, Req->MMemoryRange, Req->MAccessRange, Req->MOffset,
         Req->MElemSize, std::move(MSyncPointDeps), &OutSyncPoint);
-
+    MEvent->setSyncPoint(OutSyncPoint);
     return PI_SUCCESS;
   }
   default:

sycl/test-e2e/AOT/early_aot.cpp

@@ -0,0 +1,74 @@
+// Test early-AOT behaviors with -fsycl -fno-sycl-rdc.  This targets spir64_gen
+
+// REQUIRES: ocloc, gpu
+// UNSUPPORTED: cuda, hip
+
+// Build the early AOT device binaries
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -fno-sycl-rdc -c -DADD_CPP %s -o %t_add.o
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -fno-sycl-rdc -c -DSUB_CPP %s -o %t_sub.o
+// RUN: %clangxx -fsycl -DMAIN_CPP %s %t_add.o %t_sub.o -o %t.out
+
+// RUN: %{run} %t.out
+
+#ifdef MAIN_CPP
+// main.cpp
+
+#include "sycl/sycl.hpp"
+#include <iostream>
+
+void add(sycl::queue q, int *result, int a, int b);
+void sub(sycl::queue q, int *result, int a, int b);
+
+int main() {
+  sycl::queue q;
+  int *result = sycl::malloc_host<int>(2, q);
+  if (!result)
+    std::cout << "Error: failed to allocate USM host memory\n";
+
+  try {
+    add(q, &(result[0]), 2, 1);
+  } catch (sycl::exception const &e) {
+    std::cout
+        << "Caught synchronous SYCL exception while launching add kernel:\n"
+        << e.what() << "\n";
+    std::terminate();
+  }
+  try {
+    sub(q, &(result[1]), 2, 1);
+  } catch (sycl::exception const &e) {
+    std::cout
+        << "Caught synchronous SYCL exception while launching sub kernel:\n"
+        << e.what() << "\n";
+    std::terminate();
+  }
+  q.wait();
+
+  // Check the results
+  if (!(result[0] == 3 && result[1] == 1)) {
+    std::cout << "FAILED\n";
+    return 1;
+  }
+  std::cout << "PASSED\n";
+  return 0;
+}
+
+#endif // MAIN_CPP
+
+#ifdef ADD_CPP
+// add.cpp
+#include "sycl/sycl.hpp"
+
+void add(sycl::queue q, int *result, int a, int b) {
+  q.single_task<class add_dummy>([=] { *result = a + b; });
+}
+
+#endif // ADD_CPP
+
+#ifdef SUB_CPP
+// sub.cpp
+#include "sycl/sycl.hpp"
+
+void sub(sycl::queue q, int *result, int a, int b) {
+  q.single_task<class sub_dummy>([=] { *result = a - b; });
+}
+#endif // SUB_CPP

sycl/test-e2e/Matrix/Legacy/XMX8/element_wise_all_ops_half.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: aspect-fp16
 // REQUIRES: matrix-xmx8
+// REQUIRES: matrix-fp16
 
 // RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
 // RUN: %{run} %t.out

sycl/test-e2e/Matrix/Legacy/XMX8/joint_matrix_half.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: aspect-fp16
 // REQUIRES: matrix-xmx8
+// REQUIRES: matrix-fp16
 
 // RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
 // RUN: %{run} %t.out

sycl/test-e2e/Matrix/Legacy/element_wise_all_ops_half.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: aspect-fp16
 // REQUIRES: matrix
+// REQUIRES: matrix-fp16
 
 // RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
 // RUN: %{run} %t.out

sycl/test-e2e/Matrix/Legacy/joint_matrix_half.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: aspect-fp16
 // REQUIRES: matrix
+// REQUIRES: matrix-fp16
 
 // RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
 // RUN: %{run} %t.out

sycl/test-e2e/Matrix/SG32/element_wise_all_ops_half.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: aspect-fp16
 // REQUIRES: matrix,gpu
+// REQUIRES: matrix-fp16
 
 // RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
 // RUN: %{run} %t.out

sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: aspect-fp16
 // REQUIRES: matrix
+// REQUIRES: matrix-fp16
 
 // RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
 // RUN: %{run} %t.out

sycl/test-e2e/Matrix/XMX8/element_wise_all_ops_half.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: aspect-fp16
 // REQUIRES: matrix-xmx8
+// REQUIRES: matrix-fp16
 
 // RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
 // RUN: %{run} %t.out

sycl/test-e2e/Matrix/XMX8/joint_matrix_half.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: aspect-fp16
 // REQUIRES: matrix-xmx8
+// REQUIRES: matrix-fp16
 
 // RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
 // RUN: %{run} %t.out

sycl/test-e2e/Matrix/element_wise_all_ops_half.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: aspect-fp16
 // REQUIRES: matrix
+// REQUIRES: matrix-fp16
 
 // RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
 // RUN: %{run} %t.out

sycl/test-e2e/Matrix/joint_matrix_half.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: aspect-fp16
 // REQUIRES: matrix
+// REQUIRES: matrix-fp16
 
 // RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
 // RUN: %{run} %t.out