fix overflow in thread index calculation

upsj · yhmtsai · upsj · commit 863383bb7977 · 2019-12-17T13:46:00.000+01:00
Co-Authored-By: Yuhsiang M. Tsai &lt;yhmtsai@gmail.com&gt;
diff --git a/common/matrix/csr_kernels.hpp.inc b/common/matrix/csr_kernels.hpp.inc
@@ -501,7 +501,7 @@ __global__ __launch_bounds__(default_block_size) void spgeam_nnz(
     IndexType num_rows, IndexType *nnz)
 {
     constexpr auto sentinel = device_numeric_limits<IndexType>::max;
-    auto row = threadIdx.x + blockDim.x * blockIdx.x;
+    auto row = threadIdx.x + blockDim.x * static_cast<size_type>(blockIdx.x);
     if (row >= num_rows) {
         return;
     }
@@ -544,7 +544,7 @@ __global__ __launch_bounds__(default_block_size) void spgeam(
     const IndexType *c_row_ptrs, IndexType *c_col_idxs, ValueType *c_vals)
 {
     constexpr auto sentinel = device_numeric_limits<IndexType>::max;
-    auto row = threadIdx.x + blockDim.x * blockIdx.x;
+    auto row = threadIdx.x + blockDim.x * static_cast<size_type>(blockIdx.x);
     if (row >= num_rows) {
         return;
     }

Original file line number	Diff line number	Diff line change
`@@ -501,7 +501,7 @@ __global__ __launch_bounds__(default_block_size) void spgeam_nnz(`
`501`	`501`	`IndexType num_rows, IndexType *nnz)`
`502`	`502`	`{`
`503`	`503`	`constexpr auto sentinel = device_numeric_limits<IndexType>::max;`
`504`		`- auto row = threadIdx.x + blockDim.x * blockIdx.x;`
	`504`	`+ auto row = threadIdx.x + blockDim.x * static_cast<size_type>(blockIdx.x);`
`505`	`505`	`if (row >= num_rows) {`
`506`	`506`	`return;`
`507`	`507`	`}`
`@@ -544,7 +544,7 @@ __global__ __launch_bounds__(default_block_size) void spgeam(`
`544`	`544`	`const IndexType c_row_ptrs, IndexType c_col_idxs, ValueType *c_vals)`
`545`	`545`	`{`
`546`	`546`	`constexpr auto sentinel = device_numeric_limits<IndexType>::max;`
`547`		`- auto row = threadIdx.x + blockDim.x * blockIdx.x;`
	`547`	`+ auto row = threadIdx.x + blockDim.x * static_cast<size_type>(blockIdx.x);`
`548`	`548`	`if (row >= num_rows) {`
`549`	`549`	`return;`
`550`	`550`	`}`