Merge pull request #445 from ntrost57/fix/bsrilu0

Auto-submit by Jenkins

Author: rocm-ci

CHANGELOG.md

@@ -30,6 +30,7 @@ Documentation for rocSPARSE is available at
 ### Resolved issues
 * Fixed an issue in `rocsparse_spgemm`, `rocsparse_[s|d|c|z]csrgemm`, and `rocsparse_[s|d|c|z]bsrgemm` where incorrect results could be produced when rocSPARSE was built with optimization level `O0`. This was caused by a bug in the hash tables that could allow keys to be inserted twice.
 * Fixed an issue in the routine `rocsparse_spgemm` when using `rocsparse_spgemm_stage_symbolic` and `rocsparse_spgemm_stage_numeric`, where the routine would crash when `alpha` and `beta` were passed as host pointers and where `beta != 0`.
+* Fixed an issue in `rocsparse_bsrilu0` where the algorithm was running out of bounds of the `bsr_val` array.
 
 ### Upcoming changes
 

library/src/precond/bsrilu0_device.h

@@ -172,7 +172,9 @@ namespace rocsparse
 
                         // Load BSR block from row k into shared memory
                         sdata1[threadIdx.y][threadIdx.x]
-                            = bsr_val[BSR_IND(k, threadIdx.x, threadIdx.y, dir)];
+                            = (threadIdx.x < block_dim && threadIdx.y < block_dim)
+                                  ? bsr_val[BSR_IND(k, threadIdx.x, threadIdx.y, dir)]
+                                  : static_cast<T>(0);
 
                         // Make sure all writes to shared memory are visible
                         __threadfence_block();