message: Ensure there is no segfaults in the kernels.
diff: @@ -161,14 +161,15 @@ __global__ __launch_bounds__(default_block_size) void step_2_kernel(
  {
  const size_type tidx = blockDim.x * blockIdx.x + threadIdx.x;
  const size_type col = tidx % padding;
  - if (tidx < num_rows * padding) {
  + if (col >= num_cols || tidx >= num_rows * padding) {
  + return;
  + }
  alpha[col] = (beta[col] != zero<ValueType>()) ? rho[col] / beta[col]
  : zero<ValueType>();
  s[tidx] = (alpha[col] == zero<ValueType>())
  ? r[tidx]
  : r[tidx] - alpha[col] * v[tidx];
  }
  -}
  template <typename ValueType>
  @@ -202,7 +203,10 @@ __global__ __launch_bounds__(default_block_size) void step_3_kernel(
  {
  const size_type tidx = blockDim.x * blockIdx.x + threadIdx.x;
  const size_type col = tidx % padding;
  - if (tidx < num_rows * padding && tidx < num_rows * padding_b) {
  + if (col >= num_cols || tidx >= num_rows * padding ||
  + tidx >= num_rows * padding_b) {
  + return;
  + }
  omega[col] = (beta[col] != zero<ValueType>()) ? omega[col] / beta[col]
  : zero<ValueType>();
  // x[tidx] = (omega[col] == zero<ValueType>()) ? x[tidx] : x[tidx] +
  @@ -212,7 +216,6 @@ __global__ __launch_bounds__(default_block_size) void step_3_kernel(
  // * q[tidx];
  r[tidx] = s[tidx] - omega[col] * t[tidx];
  }
  -}
  template <typename ValueType>
  void step_3(matrix::Dense<ValueType> *x, matrix::Dense<ValueType> *r,
  