From 82b40f72f4967cdb20eced30141bbf864d06c1d5 Mon Sep 17 00:00:00 2001
From: "Victor A. P. Magri" <50467563+victorapm@users.noreply.github.com>
Date: Tue, 2 May 2023 11:21:51 -0400
Subject: [PATCH] Fix device OpenMP build (#894)

* Remove SyncCompute call to fix compilation with device omp
* Fix hypre_SeqVectorAxpyzDevice implementation for device omp
* Add warning for function not implemented for device omp
---
 src/seq_mv/vector_device.c | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/src/seq_mv/vector_device.c b/src/seq_mv/vector_device.c
index d56c7eab1..7053e5de2 100644
--- a/src/seq_mv/vector_device.c
+++ b/src/seq_mv/vector_device.c
@@ -92,7 +92,6 @@ hypre_SeqVectorScaleDevice( HYPRE_Complex alpha,
    }
 #endif
 
-   hypre_SyncComputeStream(hypre_handle());
    hypre_GpuProfilingPopRange();
 
    return hypre_error_flag;
@@ -172,7 +171,7 @@ hypre_SeqVectorAxpyzDevice( HYPRE_Complex  alpha,
    #pragma omp target teams distribute parallel for private(i) is_device_ptr(z_data, y_data, x_data)
    for (i = 0; i < total_size; i++)
    {
-      z_data[i] = alpha * x_data[i] + beta * z_data[i];
+      z_data[i] = alpha * x_data[i] + beta * y_data[i];
    }
 #endif
 
@@ -190,6 +189,7 @@ hypre_SeqVectorElmdivpyDevice( hypre_Vector *x,
                                HYPRE_Int    *marker,
                                HYPRE_Int     marker_val )
 {
+#if defined(HYPRE_USING_GPU)
    HYPRE_Complex  *x_data        = hypre_VectorData(x);
    HYPRE_Complex  *b_data        = hypre_VectorData(b);
    HYPRE_Complex  *y_data        = hypre_VectorData(y);
@@ -199,8 +199,6 @@ hypre_SeqVectorElmdivpyDevice( hypre_Vector *x,
    HYPRE_Int       size          = hypre_VectorSize(b);
 
    hypre_GpuProfilingPushRange("SeqVectorElmdivpyDevice");
-
-#if defined(HYPRE_USING_GPU)
    if (num_vectors_b == 1)
    {
       if (num_vectors_x == 1)
@@ -246,6 +244,8 @@ hypre_SeqVectorElmdivpyDevice( hypre_Vector *x,
    hypre_SyncComputeStream(hypre_handle());
    hypre_GpuProfilingPopRange();
 
+#elif defined(HYPRE_USING_OPENMP)
+   hypre_error_w_msg(HYPRE_ERROR_GENERIC, "Not implemented for device OpenMP!\n");
 #endif
 
    return hypre_error_flag;