From 96d733a0aa84a0615450eff86c4070689b5bcdab Mon Sep 17 00:00:00 2001
From: LwhJesse <256257451+LwhJesse@users.noreply.github.com>
Date: Thu, 14 May 2026 06:14:01 +0800
Subject: [PATCH] Fix CUDA block-sparse matvec contribution routing

---
 Common/src/linear_algebra/CSysMatrixGPU.cu | 7 ++-----
 1 file changed, 2 insertions(+), 5 deletions(-)

diff --git a/Common/src/linear_algebra/CSysMatrixGPU.cu b/Common/src/linear_algebra/CSysMatrixGPU.cu
index 7e0c81ca54f..e346586ae72 100644
--- a/Common/src/linear_algebra/CSysMatrixGPU.cu
+++ b/Common/src/linear_algebra/CSysMatrixGPU.cu
@@ -43,16 +43,13 @@ __global__ void GPUMatrixVectorProductAdd(matrixType* matrix, vectorType* vec, v
 
    if(row<nPointDomain && threadNo<activeThreads)
    {
-      vectorType res = 0.0;
-
       for(int index = d_row_ptr[row] * nVar * nEqn + threadNo; index < d_row_ptr[row+1] * nVar * nEqn; index+=activeThreads)
       {
          int blockCol = index%nEqn;
          int blockNo = index/(nVar * nEqn);
-         res += matrix[index] * vec[(d_col_ind[blockNo])*nVar + blockCol];
+         int trueBlockRow = (index % (nVar * nEqn)) / nEqn;
+         atomicAdd(&prod[row * nVar + trueBlockRow], matrix[index] * vec[(d_col_ind[blockNo])*nVar + blockCol]);
       }
-
-      atomicAdd(&prod[row * nVar + blockRow], res);
    }
 }