[SYCL][Fusion] Test kernel fusion and optimization (intel#1535)

Test different scenarios for kernel fusion, including creation of the fused kernel by the JIT compiler and performance optimizations such as dataflow internalization.

Automatically detect availability of the kernel fusion extension in the DPC++ build in `lit.cfg.py` and make it available for `REQUIRES` clauses.

Spec: intel#7098
Implementation: intel#7831

Signed-off-by: Lukas Sommer <[email protected]>

Author: sommerlukas

SYCL/KernelFusion/complete_fusion.cpp

@@ -0,0 +1,67 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+// REQUIRES: fusion
+
+// Test complete fusion without any internalization
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+int main() {
+  constexpr size_t dataSize = 512;
+  int in1[dataSize], in2[dataSize], in3[dataSize], tmp[dataSize], out[dataSize];
+
+  for (size_t i = 0; i < dataSize; ++i) {
+    in1[i] = i * 2;
+    in2[i] = i * 3;
+    in3[i] = i * 4;
+    tmp[i] = -1;
+    out[i] = -1;
+  }
+
+  queue q{ext::codeplay::experimental::property::queue::enable_fusion{}};
+
+  {
+    buffer<int> bIn1{in1, range{dataSize}};
+    buffer<int> bIn2{in2, range{dataSize}};
+    buffer<int> bIn3{in3, range{dataSize}};
+    buffer<int> bTmp{tmp, range{dataSize}};
+    buffer<int> bOut{out, range{dataSize}};
+
+    ext::codeplay::experimental::fusion_wrapper fw{q};
+    fw.start_fusion();
+
+    assert(fw.is_in_fusion_mode() && "Queue should be in fusion mode");
+
+    q.submit([&](handler &cgh) {
+      auto accIn1 = bIn1.get_access(cgh);
+      auto accIn2 = bIn2.get_access(cgh);
+      auto accTmp = bTmp.get_access(cgh);
+      cgh.parallel_for<class KernelOne>(
+          dataSize, [=](id<1> i) { accTmp[i] = accIn1[i] + accIn2[i]; });
+    });
+
+    q.submit([&](handler &cgh) {
+      auto accTmp = bTmp.get_access(cgh);
+      auto accIn3 = bIn3.get_access(cgh);
+      auto accOut = bOut.get_access(cgh);
+      cgh.parallel_for<class KernelTwo>(
+          dataSize, [=](id<1> i) { accOut[i] = accTmp[i] * accIn3[i]; });
+    });
+
+    fw.complete_fusion({ext::codeplay::experimental::property::no_barriers{}});
+
+    assert(!fw.is_in_fusion_mode() &&
+           "Queue should not be in fusion mode anymore");
+  }
+
+  // Check the results
+  for (size_t i = 0; i < dataSize; ++i) {
+    assert(out[i] == (20 * i * i) && "Computation error");
+  }
+
+  return 0;
+}

SYCL/lit.cfg.py

@@ -427,6 +427,23 @@
     else:
         lit_config.warning("Couldn't find pre-installed AOT device compiler " + aot_tool)
 
+# Check if kernel fusion is available by compiling a small program that will
+# be ill-formed (compilation stops with non-zero exit code) if the feature
+# test macro for kernel fusion is not defined.
+check_fusion_file = 'check_fusion.cpp'
+with open(check_fusion_file, 'w') as ff:
+    ff.write('#include <sycl/sycl.hpp>\n')
+    ff.write('#ifndef SYCL_EXT_CODEPLAY_KERNEL_FUSION\n')
+    ff.write('#error \"Feature test for fusion failed\"\n')
+    ff.write('#endif // SYCL_EXT_CODEPLAY_KERNEL_FUSION\n')
+    ff.write('int main() { return 0; }\n')
+
+status = subprocess.getstatusoutput(config.dpcpp_compiler + ' -fsycl  ' +
+                                    check_fusion_file)
+if status[0] == 0:
+    lit_config.note('Kernel fusion extension enabled')
+    config.available_features.add('fusion')
+
 # Set timeout for a single test
 try:
     import psutil