From f935ea0e6485f0dd28e26741984973ec51e35316 Mon Sep 17 00:00:00 2001
From: "Cai, Justin" <justin.cai@intel.com>
Date: Fri, 2 Dec 2022 09:30:52 -0800
Subject: [PATCH] [SYCLomatic #378] Added default stream test

Signed-off-by: Cai, Justin <justin.cai@intel.com>
---
 features/config/TEMPLATE_defaultStream.xml    |  9 ++++
 .../driverStreamAndEvent/defaultStream.cu     | 54 +++++++++++++++++++
 features/features.xml                         |  1 +
 features/test_feature.py                      |  2 +-
 4 files changed, 65 insertions(+), 1 deletion(-)
 create mode 100644 features/config/TEMPLATE_defaultStream.xml
 create mode 100644 features/feature_case/driverStreamAndEvent/defaultStream.cu

diff --git a/features/config/TEMPLATE_defaultStream.xml b/features/config/TEMPLATE_defaultStream.xml
new file mode 100644
index 000000000..dc009f2f9
--- /dev/null
+++ b/features/config/TEMPLATE_defaultStream.xml
@@ -0,0 +1,9 @@
+<?xml version="1.0" encoding="UTF-8"?>
+
+<test driverID="test_feature" name="TEMPLATE">
+    <description>test</description>
+    <files>
+      <file path="feature_case/driverStreamAndEvent/defaultStream.cu" />
+    </files>
+    <rules />
+</test>
diff --git a/features/feature_case/driverStreamAndEvent/defaultStream.cu b/features/feature_case/driverStreamAndEvent/defaultStream.cu
new file mode 100644
index 000000000..c2adba2c0
--- /dev/null
+++ b/features/feature_case/driverStreamAndEvent/defaultStream.cu
@@ -0,0 +1,54 @@
+#include <iostream>
+#include <cstdint>
+#include <cuda_runtime.h>
+
+struct stream_wrapper {
+  cudaStream_t s;
+  stream_wrapper() : s(0) {}
+  void set(cudaStream_t t) { s = t; }
+  void set_int(uintptr_t p) { s = (cudaStream_t) p; }
+};
+
+int main() {
+  stream_wrapper wr {};
+  int x, y;
+  int res = 0;
+  int i = 0;
+
+  auto run = [&]() {
+    x = -1;
+    y = 42;
+    cudaMemcpyAsync(&x, &y, sizeof(int), cudaMemcpyDefault, wr.s);
+    cudaStreamSynchronize(wr.s);
+    if (x != y) {
+      std::cout << "default stream fail " << i << "\n";
+      res = 1;
+    }
+    ++i;
+  };
+
+  run();
+
+  wr.set(cudaStreamDefault);
+  run();
+  
+  wr.set(cudaStreamLegacy);
+  run();
+
+  wr.set(cudaStreamPerThread);
+  run();
+
+  wr.set_int(0);
+  run();
+
+  cudaStream_t s;
+  cudaStreamCreate(&s);
+  wr.set_int((uintptr_t) s);
+  run();
+
+  if (!res) {
+    std::cout << "default stream success\n";
+  }
+
+  return res;  
+}
diff --git a/features/features.xml b/features/features.xml
index 07bfee6c7..fd5fbb018 100644
--- a/features/features.xml
+++ b/features/features.xml
@@ -391,6 +391,7 @@
     <test testName="libcu_atomic" configFile="config/TEMPLATE_libcu.xml" />
     <test testName="pointer_attributes" configFile="config/TEMPLATE_pointer_attributes.xml" />
     <test testName="image" configFile="config/TEMPLATE_image.xml" />
+    <test testName="defaultStream" configFile="config/TEMPLATE_defaultStream.xml" />
     <!-- <test testName="math_intel_specific" configFile="config/TEMPLATE_math_intel_specific.xml" splitGroup="double" /> -->
   </tests>
 </suite>
diff --git a/features/test_feature.py b/features/test_feature.py
index 3b76c7363..efde2a234 100644
--- a/features/test_feature.py
+++ b/features/test_feature.py
@@ -37,7 +37,7 @@
               'cusolver_test1', 'cusolver_test2', 'thrust_op', 'cublas-extension', 'cublas_v1_runable', 'thrust_minmax_element',
               'thrust_is_sorted', 'thrust_partition', 'thrust_remove_copy', 'thrust_unique_copy', 'thrust_transform_exclusive_scan',
               'thrust_set_difference', 'thrust_set_difference_by_key', 'thrust_set_intersection_by_key', 'thrust_stable_sort',
-              'thrust_tabulate', 'thrust_for_each_n', 'device_info']
+              'thrust_tabulate', 'thrust_for_each_n', 'device_info', 'defaultStream']
 
 def setup_test():
     return True