spr workarround for AdaptiveCpp to 'avoid' group broadcast instruction

codeplaysoftware · Jan 29, 2024 · affe6e0 · affe6e0
1 parent 199717f
commit affe6e0
Show file tree

Hide file tree

Showing 3 changed files with 14 additions and 12 deletions.
diff --git a/benchmark/portblas/CMakeLists.txt b/benchmark/portblas/CMakeLists.txt
@@ -77,15 +77,13 @@ if(${BLAS_ENABLE_EXTENSIONS})
 endif()
 
 # Skip these benchmarks for AdaptiveCpp for SPIRV/OpenCL targets
-# that use SYCL 2020 features like reduction or hang during execution
-# (https://github.com/AdaptiveCpp/AdaptiveCpp/issues/1309)
+# that use SYCL 2020 features like group reduction or hang 
+# during execution (https://github.com/AdaptiveCpp/AdaptiveCpp/issues/1309)
 set(ADAPTIVE_CPP_SKIP
   blas1/asum.cpp
   blas1/dot.cpp
   blas1/sdsdot.cpp
   blas1/nrm2.cpp
-  blas2/spr.cpp
-  blas2/spr2.cpp
   blas2/trsv.cpp
   blas2/tbsv.cpp
   blas2/tpsv.cpp

diff --git a/src/operations/blas2/spr.hpp b/src/operations/blas2/spr.hpp
@@ -101,15 +101,20 @@ typename rhs_1_t::value_t Spr<Single, isUpper, lhs_t, rhs_1_t, rhs_2_t>::eval(
 
   index_t row = 0, col = 0;
 
+#ifndef __HIPSYCL__
   if (!id) {
+#endif
     Spr<Single, isUpper, lhs_t, rhs_1_t, rhs_2_t>::compute_row_col(
         global_idx, N_, row, col);
+#ifndef __HIPSYCL__
   }
 
   row = cl::sycl::group_broadcast(ndItem.get_group(), row);
   col = cl::sycl::group_broadcast(ndItem.get_group(), col);
+#endif
 
   if (global_idx < lhs_size) {
+#ifndef __HIPSYCL__
     if constexpr (isUpper) {
       if (id) {
         row += id;
@@ -127,6 +132,7 @@ typename rhs_1_t::value_t Spr<Single, isUpper, lhs_t, rhs_1_t, rhs_2_t>::eval(
         }
       }
     }
+#endif
 
     value_t lhs_val = lhs_.eval(global_idx);
     value_t rhs_1_val = rhs_1_.eval(row);
@@ -135,7 +141,8 @@ typename rhs_1_t::value_t Spr<Single, isUpper, lhs_t, rhs_1_t, rhs_2_t>::eval(
       value_t rhs_1_val_second = rhs_1_.eval(col);
       value_t rhs_2_val_second = rhs_2_.eval(row);
       lhs_.eval(global_idx) = rhs_1_val * rhs_2_val * alpha_ +
-            rhs_1_val_second * rhs_2_val_second * alpha_ + lhs_val;
+                              rhs_1_val_second * rhs_2_val_second * alpha_ +
+                              lhs_val;
     } else
       lhs_.eval(global_idx) = rhs_1_val * rhs_2_val * alpha_ + lhs_val;
   }
@@ -161,9 +168,8 @@ Spr<Single, isUpper, lhs_t, rhs_1_t, rhs_2_t>::adjust_access_displacement() {
 
 template <bool Single, bool isUpper, typename lhs_t, typename rhs_1_t,
           typename rhs_2_t>
-PORTBLAS_INLINE
-    typename Spr<Single, isUpper, lhs_t, rhs_1_t, rhs_2_t>::index_t
-    Spr<Single, isUpper, lhs_t, rhs_1_t, rhs_2_t>::get_size() const {
+PORTBLAS_INLINE typename Spr<Single, isUpper, lhs_t, rhs_1_t, rhs_2_t>::index_t
+Spr<Single, isUpper, lhs_t, rhs_1_t, rhs_2_t>::get_size() const {
   return rhs_1_.get_size();
 }
 template <bool Single, bool isUpper, typename lhs_t, typename rhs_1_t,

diff --git a/test/unittest/CMakeLists.txt b/test/unittest/CMakeLists.txt
@@ -74,16 +74,14 @@ set(SYCL_UNITTEST_SRCS
 )
 
 # Skip these tests for AdaptiveCpp for SPIRV/OpenCL targets 
-# that use SYCL 2020 features like group reduction/broadcast 
-# or hang during execution (https://github.com/AdaptiveCpp/AdaptiveCpp/issues/1309)
+# that use SYCL 2020 features like group reduction or hang 
+# during execution (https://github.com/AdaptiveCpp/AdaptiveCpp/issues/1309)
 set(ADAPTIVE_CPP_SKIP 
   ${PORTBLAS_UNITTEST}/blas1/blas1_asum_test.cpp
   ${PORTBLAS_UNITTEST}/blas1/blas1_sdsdot_test.cpp
   ${PORTBLAS_UNITTEST}/blas1/blas1_nrm2_test.cpp
   ${PORTBLAS_UNITTEST}/blas1/blas1_dot_test.cpp
   ${PORTBLAS_UNITTEST}/blas1/blas1_rot_test.cpp
-  ${PORTBLAS_UNITTEST}/blas2/blas2_spr_test.cpp
-  ${PORTBLAS_UNITTEST}/blas2/blas2_spr2_test.cpp
   # Hang during execution (without failing)
   ${PORTBLAS_UNITTEST}/blas3/blas3_trsm_test.cpp
 )