From 7f62842382f65c37465bf35bb7b109e0d013d0e2 Mon Sep 17 00:00:00 2001 From: James Sandham <33790278+jsandham@users.noreply.github.com> Date: Fri, 22 Dec 2023 00:27:54 -0700 Subject: [PATCH] Cherrypick adaptive fix from #688 PR (#363) --- library/src/level2/csrmv_device.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/library/src/level2/csrmv_device.h b/library/src/level2/csrmv_device.h index 9ebfe6a78..1c14c26c3 100644 --- a/library/src/level2/csrmv_device.h +++ b/library/src/level2/csrmv_device.h @@ -441,11 +441,10 @@ ROCSPARSE_DEVICE_ILF void csrmvn_adaptive_device(bool conj, // For every other workgroup, wg_flags[first_wg_in_row] holds the value they wait on. // If your flag == first_wg's flag, you spin loop. // The first workgroup will eventually flip this flag, and you can move forward. - __syncthreads(); + __threadfence(); while(gid != first_wg_in_row && lid == 0 && ((rocsparse_atomic_max(&wg_flags[first_wg_in_row], 0U)) == compare_value)) ; - __syncthreads(); // After you've passed the barrier, update your local flag to make sure that // the next time through, you know what to wait on.