From 5ade20654070e71db363c3b5a7a1d85dccff9548 Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Wed, 14 Jul 2021 12:55:55 +0100 Subject: [PATCH 1/2] [SYCL][CUDA] Unoptimized stream regression test Previously the CUDA backend would fail due to invalid atomic memory orders not being optimized out. The use of these invalid memory orders have been removed in recent changes, so this commit adds a regression test to make sure they do not resurface. Signed-off-by: Steffen Larsen --- SYCL/Regression/unoptimized_stream.cpp | 15 +++++++++++++++ 1 file changed, 15 insertions(+) create mode 100644 SYCL/Regression/unoptimized_stream.cpp diff --git a/SYCL/Regression/unoptimized_stream.cpp b/SYCL/Regression/unoptimized_stream.cpp new file mode 100644 index 0000000000..cdcf1f50ad --- /dev/null +++ b/SYCL/Regression/unoptimized_stream.cpp @@ -0,0 +1,15 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -O0 -o %t.out +// RUN: %HOST_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +#include + +int main() { + sycl::queue q; + q.submit([&](sycl::handler &cgh) { + sycl::stream os(1024, 256, cgh); + cgh.single_task([=]() { os << "test"; }); + }); + return 0; +} From 4437c559ef4c6fd7cfa0122b9ffba242164b1c82 Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Tue, 20 Jul 2021 10:00:15 +0100 Subject: [PATCH 2/2] Add test motivation note Signed-off-by: Steffen Larsen --- SYCL/Regression/unoptimized_stream.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/SYCL/Regression/unoptimized_stream.cpp b/SYCL/Regression/unoptimized_stream.cpp index cdcf1f50ad..e9d800bdc2 100644 --- a/SYCL/Regression/unoptimized_stream.cpp +++ b/SYCL/Regression/unoptimized_stream.cpp @@ -3,6 +3,11 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +// NOTE: The libclc target used by the CUDA backend used to generate atomic load +// variants that were unsupported by NVPTX. Even if they were not used +// directly, sycl::stream and other operations would keep the invalid +// operations in when optimizations were disabled. + #include int main() {