csrsm bugfix (#475) (#325)

* csrsm bugfix * same fix for bsrsm
ROCm · Jun 20, 2023 · 690915f · 690915f
1 parent 8252a37
commit 690915f
Show file tree

Hide file tree

Showing 3 changed files with 16 additions and 15 deletions.
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -5,6 +5,7 @@ Full documentation for rocSPARSE is available at [rocsparse.readthedocs.io](http
 ## rocSPARSE 2.5.2 for ROCm 5.6.0
 ### Improved
 - Fixed a memory leak in csritsv
+- Fixed a bug in csrsm and bsrsm
 
 ## rocSPARSE 2.5.1 for ROCm 5.5.0
 ### Added

diff --git a/library/src/level3/bsrsm_device_large.h b/library/src/level3/bsrsm_device_large.h
@@ -105,12 +105,12 @@ __launch_bounds__(BLOCKSIZE) ROCSPARSE_KERNEL
             }
         }
 
-        // Wait for spin looping thread to finish as the whole block depends on this row
-        __syncthreads();
-
         // Make sure updated X is visible globally
         __threadfence();
 
+        // Wait for spin looping thread to finish as the whole block depends on this row
+        __syncthreads();
+
         // Local sum computation
 
         // Do not run out of bounds
@@ -170,12 +170,12 @@ __launch_bounds__(BLOCKSIZE) ROCSPARSE_KERNEL
         }
     }
 
-    // Wait for all threads to finish writing into global memory before we mark the row "done"
-    __syncthreads();
-
     // Make sure X is written to global memory before setting row is done flag
     __threadfence();
 
+    // Wait for all threads to finish the threadfence before we mark the row "done"
+    __syncthreads();
+
     if(row < mb && threadIdx.x == 0)
     {
         // Write "row is done" flag
@@ -267,12 +267,12 @@ __launch_bounds__(BLOCKSIZE) ROCSPARSE_KERNEL
             }
         }
 
-        // Wait for spin looping thread to finish as the whole block depends on this row
-        __syncthreads();
-
         // Make sure updated X is visible globally
         __threadfence();
 
+        // Wait for spin looping thread to finish as the whole block depends on this row
+        __syncthreads();
+
         // Local sum computation
 
         // Do not run out of bounds
@@ -332,12 +332,12 @@ __launch_bounds__(BLOCKSIZE) ROCSPARSE_KERNEL
         }
     }
 
-    // Wait for all threads to finish writing into global memory before we mark the row "done"
-    __syncthreads();
-
     // Make sure X is written to global memory before setting row is done flag
     __threadfence();
 
+    // Wait for all threads to finish the threadfence before we mark the row "done"
+    __syncthreads();
+
     if(row < mb && threadIdx.x == 0)
     {
         // Write "row is done" flag

diff --git a/library/src/level3/csrsm_device.h b/library/src/level3/csrsm_device.h
@@ -231,12 +231,12 @@ __device__ void csrsm_device(rocsparse_operation transB,
         B[idx_B] = local_sum;
     }
 
-    // Wait for all threads to finish writing into global memory before we mark the row "done"
-    __syncthreads();
-
     // Make sure B is written to global memory before setting row is done flag
     __threadfence();
 
+    // Wait for all threads to finish the threadfence before we mark the row "done"
+    __syncthreads();
+
     if(hipThreadIdx_x == 0)
     {
         // Write the "row is done" flag