Use n_vecs = 2 for minimum/maximum to work around a bug in CPU driver for AMD hardware

oleksandr-pavlyk · oleksandr-pavlyk · commit bd52b277a64f · 2024-10-27T21:58:20.000-05:00
diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/maximum.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/maximum.hpp
@@ -206,7 +206,7 @@ sycl::event maximum_contig_impl(sycl::queue &exec_q,
 {
     using resTy = typename MaximumOutputType<argTy1, argTy2>::value_type;
     constexpr auto vec_sz = VecSize_v<argTy1, argTy2, resTy>;
-    constexpr unsigned int n_vecs = 1u;
+    constexpr unsigned int n_vecs = 2u;
 
     return elementwise_common::binary_contig_impl<
         argTy1, argTy2, MaximumOutputType, MaximumContigFunctor,
diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/minimum.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/minimum.hpp
@@ -206,7 +206,7 @@ sycl::event minimum_contig_impl(sycl::queue &exec_q,
 {
     using resTy = typename MinimumOutputType<argTy1, argTy2>::value_type;
     constexpr auto vec_sz = VecSize_v<argTy1, argTy2, resTy>;
-    constexpr unsigned int n_vecs = 1u;
+    constexpr unsigned int n_vecs = 2u;
 
     return elementwise_common::binary_contig_impl<
         argTy1, argTy2, MinimumOutputType, MinimumContigFunctor,

Original file line number	Diff line number	Diff line change
`@@ -206,7 +206,7 @@ sycl::event maximum_contig_impl(sycl::queue &exec_q,`
`206`	`206`	`{`
`207`	`207`	`using resTy = typename MaximumOutputType<argTy1, argTy2>::value_type;`
`208`	`208`	`constexpr auto vec_sz = VecSize_v<argTy1, argTy2, resTy>;`
`209`		`- constexpr unsigned int n_vecs = 1u;`
	`209`	`+ constexpr unsigned int n_vecs = 2u;`
`210`	`210`
`211`	`211`	`return elementwise_common::binary_contig_impl<`
`212`	`212`	`argTy1, argTy2, MaximumOutputType, MaximumContigFunctor,`
Original file line number	Diff line number	Diff line change
`@@ -206,7 +206,7 @@ sycl::event minimum_contig_impl(sycl::queue &exec_q,`
`206`	`206`	`{`
`207`	`207`	`using resTy = typename MinimumOutputType<argTy1, argTy2>::value_type;`
`208`	`208`	`constexpr auto vec_sz = VecSize_v<argTy1, argTy2, resTy>;`
`209`		`- constexpr unsigned int n_vecs = 1u;`
	`209`	`+ constexpr unsigned int n_vecs = 2u;`
`210`	`210`
`211`	`211`	`return elementwise_common::binary_contig_impl<`
`212`	`212`	`argTy1, argTy2, MinimumOutputType, MinimumContigFunctor,`