Repository: taichi-dev/taichi (language: python)

Issue:
不一定能复现的奇怪错误
<!-- We've collected some common issue solutions in https://docs.taichi-lang.org/docs/install. Make sure you've check them out first. Hopefully they could address your problem. -->

**Describe the bug**
用taichi实现了一下opencv的pyrDown函数，计算上就是一个带有reflectpad2d的conv2d，kernel是一个5*5的高斯核。

taichi能正常实现，并且与opencv的结果对比误差也是在范围内的。

但我想一次性处理多张图，所以我加了一个batch size维度，来一次性处理多张图，这个增加了batch的代码在“一些电脑”上报错了。

**To Reproduce**

不带batch size的只处理一张图的能用的代码：

```py
import os
os.environ["CUDA_VISIBLE_DEVICES"] = '0'

import torch
import taichi as ti
ti.init(arch=ti.cpu,debug=True)

@ti.kernel
def pyrDown_ti(input:ti.types.ndarray(),
               weight:ti.types.ndarray(),
               output:ti.types.ndarray(),
               src_size:ti.i32, dst_size:ti.i32):
    for oh, ow, wh, ww in ti.ndrange(dst_size, dst_size, 5, 5):
        ih:ti.i32 = 2*oh+wh-2
        iw:ti.i32 = 2*ow+ww-2
        if ih < 0: ih = ti.abs(ih)
        if iw < 0: iw = ti.abs(iw)
        if ih >= src_size: ih = src_size-ih-2
        if iw >= src_size: iw = src_size-iw-2
        output[oh,ow] += input[ih,iw] * weight[wh,ww]

input = torch.rand(256,256).cuda()
weight = torch.rand(5,5).cuda()
output = torch.rand(128,128).cuda()

pyrDown_ti(input, weight, output, 256, 128)
```

带batch一次处理多张图但会报错的代码：

```py
import os
os.environ["CUDA_VISIBLE_DEVICES"] = '0'

import torch
import taichi as ti
ti.init(arch=ti.cpu,debug=True)

@ti.kernel
def pyrDown_ti(input:ti.types.ndarray(),
               weight:ti.types.ndarray(),
               output:ti.types.ndarray(),
               src_size:ti.i32, dst_size:ti.i32):
    for oh, ow, b, wh, ww in ti.ndrange(dst_size, dst_size, 256, 5, 5):
        ih:ti.i32 = 2*oh+wh-2
        iw:ti.i32 = 2*ow+ww-2
        if ih < 0: ih = ti.abs(ih)
        if iw < 0: iw = ti.abs(iw)
        if ih >= src_size: ih = src_size-ih-2
        if iw >= src_size: iw = src_size-iw-2
        output[oh,ow,b] += input[ih,iw,b] * weight[wh,ww]

input = torch.rand(256,256,256).cuda()
weight = torch.rand(5,5).cuda()
output = torch.rand(128,128,256).cuda()

pyrDown_ti(input, weight, output, 256, 128)
```

**Log/Screenshots**
在四台电脑上进行了测试，其中三台电脑失败了。

成功的设备：
```
win11
1*RTX3060
NVIDIA-SMI 516.59       Driver Version: 516.59       CUDA Version: 11.7
torch1.13.1+cu117
[Taichi] version 1.4.1, llvm 15.0.1, commit e67c674e, win, python 3.8.10
[Taichi] Starting on arch=cuda
```

失败的设备：
```
设备一：
Linux server 5.15.0-60-generic #66~20.04.1-Ubuntu SMP Wed Jan 25 09:41:30 UTC 2023 x86_64 x86_64 x86_64 GNU/Linux
4*RTX3090
NVIDIA-SMI 470.161.03   Driver Version: 470.161.03   CUDA Version: 11.4
torch 1.13.1+cu117
[Taichi] version 1.4.1, llvm 15.0.4, commit e67c674e, linux, python 3.9.16
[Taichi] Starting on arch=x64

设备二：
Linux server 5.15.0-58-generic #64~20.04.1-Ubuntu SMP Fri Jan 6 16:42:31 UTC 2023 x86_64 x86_64 x86_64 GNU/Linux
4*RTX3090
NVIDIA-SMI 470.161.03   Driver Version: 470.161.03   CUDA Version: 11.4
torch 1.12.1+cu113
[Taichi] version 1.4.1, llvm 15.0.4, commit e67c674e, linux, python 3.8.13
[Taichi] Starting on arch=x64

设备三：
Linux xxx 5.19.0-32-generic #33~22.04.1-Ubuntu SMP PREEMPT_DYNAMIC Mon Jan 30 17:03:34 UTC 2 x86_64 x86_64 x86_64 GNU/Linux
1*RTX4090
NVIDIA-SMI 525.78.01    Driver Version: 525.78.01    CUDA Version: 12.0
torch 1.13.1+cu117
[Taichi] version 1.4.1, llvm 15.0.4, commit e67c674e, linux, python 3.9.16
[Taichi] Starting on arch=x64
```

!!! 在失败的设备上，不管是x64还是cuda都是失败的!!!

下面是ti.init(arch=ti.gpu,debug=True)的错误log：
```
[Taichi] version 1.4.1, llvm 15.0.4, commit e67c674e, linux, python 3.9.16
[Taichi] Starting on arch=cuda
[E 02/23/23 08:02:34.294 3137491] [cuda_driver.h:operator()@88] CUDA Error CUDA_ERROR_ILLEGAL_ADDRESS: an illegal memory access was encountered while calling stream_synchronize (cuStreamSynchronize)


***********************************
* Taichi Compiler Stack Traceback *
***********************************
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x32b1ad7) [0x7f1dca754ad7]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x5e10ed) [0x7f1dc7a840ed]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x9ed6af) [0x7f1dc7e906af]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x9559fe) [0x7f1dc7df89fe]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x941782) [0x7f1dc7de4782]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so: taichi::lang::Kernel::operator()(taichi::lang::Kernel::LaunchContextBuilder&)
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x899273) [0x7f1dc7d3c273]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x765b80) [0x7f1dc7c08b80]
python(+0x149dc7) [0x5625fdc86dc7]
python(_PyObject_MakeTpCall+0x2fc) [0x5625fdc6f2bc]
python(+0x147ba5) [0x5625fdc84ba5]
python(+0xa0a14) [0x5625fdbdda14]
python(_PyObject_MakeTpCall+0x2fc) [0x5625fdc6f2bc]
python(_PyEval_EvalFrameDefault+0x4d62) [0x5625fdc6aef2]
python(+0x12827a) [0x5625fdc6527a]
python(_PyFunction_Vectorcall+0xd5) [0x5625fdc76c05]
python(_PyEval_EvalFrameDefault+0x402b) [0x5625fdc6a1bb]
python(+0x12827a) [0x5625fdc6527a]
python(_PyObject_FastCallDictTstate+0x13e) [0x5625fdc6eb1e]
python(_PyObject_Call_Prepend+0xe3) [0x5625fdc82283]
python(+0x2110d9) [0x5625fdd4e0d9]
python(PyObject_Call+0x158) [0x5625fdc851a8]
python(_PyEval_EvalFrameDefault+0x402b) [0x5625fdc6a1bb]
python(+0x12827a) [0x5625fdc6527a]
python(_PyFunction_Vectorcall+0xd5) [0x5625fdc76c05]
python(_PyEval_EvalFrameDefault+0x3d8) [0x5625fdc66568]
python(+0x12827a) [0x5625fdc6527a]
python(_PyEval_EvalCodeWithName+0x47) [0x5625fdc64f07]
python(PyEval_EvalCodeEx+0x39) [0x5625fdc64eb9]
python(PyEval_EvalCode+0x1b) [0x5625fdd158bb]
python(+0x206adc) [0x5625fdd43adc]
python(+0x202c24) [0x5625fdd3fc24]
python(+0x94360) [0x5625fdbd1360]
python(PyRun_SimpleFileExFlags+0x1b6) [0x5625fdd39656]
python(Py_RunMain+0x38a) [0x5625fdd369ca]
python(Py_BytesMain+0x39) [0x5625fdd09359]
/lib/x86_64-linux-gnu/libc.so.6: __libc_start_main
python(+0x1cc251) [0x5625fdd09251]

Internal error occurred. Check out this page for possible solutions:
https://docs.taichi-lang.org/docs/install
Traceback (most recent call last):
  File "/home/wjx/cross_line/test/bug.py", line 26, in <module>
    pyrDown_ti(input, weight, output, 256, 128)
  File "/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/lang/kernel_impl.py", line 974, in wrapped
    return primal(*args, **kwargs)
  File "/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/lang/kernel_impl.py", line 901, in __call__
    return self.runtime.compiled_functions[key](*args)
  File "/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/lang/kernel_impl.py", line 826, in func__
    raise e from None
  File "/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/lang/kernel_impl.py", line 823, in func__
    t_kernel(launch_ctx)
RuntimeError: [cuda_driver.h:operator()@88] CUDA Error CUDA_ERROR_ILLEGAL_ADDRESS: an illegal memory access was encountered while calling stream_synchronize (cuStreamSynchronize)
[E 02/23/23 08:02:34.404 3137491] [cuda_driver.h:operator()@88] CUDA Error CUDA_ERROR_ILLEGAL_ADDRESS: an illegal memory access was encountered while calling stream_synchronize (cuStreamSynchronize)


***********************************
* Taichi Compiler Stack Traceback *
***********************************
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x32b1ad7) [0x7f1dca754ad7]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x5e10ed) [0x7f1dc7a840ed]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x9a56fc) [0x7f1dc7e486fc]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so: taichi::lang::Program::finalize()
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so: taichi::lang::Program::~Program()
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x877860) [0x7f1dc7d1a860]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x76c68b) [0x7f1dc7c0f68b]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x76c39c) [0x7f1dc7c0f39c]
python(+0x1259e6) [0x5625fdc629e6]
python(+0x208263) [0x5625fdd45263]
python(+0x11eb44) [0x5625fdc5bb44]
python(_PyGC_CollectNoFail+0x37) [0x5625fdd447b7]
python(Py_FinalizeEx+0x190) [0x5625fdd430e0]
python(Py_RunMain+0x112) [0x5625fdd36752]
python(Py_BytesMain+0x39) [0x5625fdd09359]
/lib/x86_64-linux-gnu/libc.so.6: __libc_start_main
python(+0x1cc251) [0x5625fdd09251]

Internal error occurred. Check out this page for possible solutions:
https://docs.taichi-lang.org/docs/install
terminate called after throwing an instance of 'std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >'
[E 02/23/23 08:02:34.435 3137491] Received signal 6 (Aborted)


***********************************
* Taichi Compiler Stack Traceback *
***********************************
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x32b1ad7) [0x7f1dca754ad7]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x6922e4) [0x7f1dc7b352e4]
/lib/x86_64-linux-gnu/libc.so.6(+0x43090) [0x7f1e9c60a090]
/lib/x86_64-linux-gnu/libc.so.6: gsignal
/lib/x86_64-linux-gnu/libc.so.6: abort
/lib/x86_64-linux-gnu/libstdc++.so.6(+0x9e911) [0x7f1e7cffd911]
/lib/x86_64-linux-gnu/libstdc++.so.6(+0xaa38c) [0x7f1e7d00938c]
/lib/x86_64-linux-gnu/libstdc++.so.6(+0xaa3f7) [0x7f1e7d0093f7]
/home/wjx/miniconda3/envs/torch/lib/python3.9/site-packages/taichi/_lib/core/taichi_python.cpython-39-x86_64-linux-gnu.so(+0x44ed2b) [0x7f1dc78f1d2b]

Internal error occurred. Check out this page for possible solutions:
https://docs.taichi-lang.org/docs/install
```


**Additional comments**
If possible, please also consider attaching the output of command `ti diagnose`. This produces the detailed environment information and hopefully helps us diagnose faster.

If you have local commits (e.g. compile fixes before you reproduce the bug), please make sure you first make a PR to fix the build errors and then report the bug.



Discussion / hints:
@EdVince There's an real out of bound access when your `iw==src_size (iw=256)`, `if iw >= src_size: iw = src_size-iw-2` will make it `-2` and that's an invalid memory access. 

But this is a valid bug report as `debug=True` should have caught this so we'll look into a fix, thanks for reporting the issue! 

Fix patch:
diff --git a/taichi/transforms/check_out_of_bound.cpp b/taichi/transforms/check_out_of_bound.cpp
index 3607d08cf1dda..009b104cea506 100644
--- a/taichi/transforms/check_out_of_bound.cpp
+++ b/taichi/transforms/check_out_of_bound.cpp
@@ -40,6 +40,59 @@ class CheckOutOfBound : public BasicStmtVisitor {
     // TODO: implement bound check here for other situations.
   }
 
+  void visit(ExternalPtrStmt *stmt) override {
+    if (is_done(stmt))
+      return;
+    auto new_stmts = VecStatement();
+    auto zero = new_stmts.push_back<ConstStmt>(TypedConstant(0));
+    Stmt *result = new_stmts.push_back<ConstStmt>(TypedConstant(true));
+    std::string msg = fmt::format(
+        "[kernel={}] Out of bound access to ndarray at arg {} with indices [",
+        kernel_name, stmt->base_ptr->as<ArgLoadStmt>()->arg_id);
+    std::vector<Stmt *> args;
+    int flattened_element = 1;
+    for (int i = 0; i < stmt->element_shape.size(); i++) {
+      flattened_element *= stmt->element_shape[i];
+    }
+    for (int i = 0; i < stmt->indices.size(); i++) {
+      auto lower_bound = zero;
+      auto check_lower_bound = new_stmts.push_back<BinaryOpStmt>(
+          BinaryOpType::cmp_ge, stmt->indices[i], lower_bound);
+      Stmt *upper_bound{nullptr};
+      // TODO: Simplify logic here since SOA layout for ndarray is deprecated
+      if ((stmt->element_dim < 0 && i == (stmt->indices.size() - 1)) ||
+          (stmt->element_dim > 0 && i == 0)) {
+        upper_bound =
+            new_stmts.push_back<ConstStmt>(TypedConstant(flattened_element));
+      } else {
+        auto axis = stmt->element_dim <= 0 ? i : (i - stmt->element_dim);
+        upper_bound = new_stmts.push_back<ExternalTensorShapeAlongAxisStmt>(
+            /*axis=*/axis,
+            /*arg_id=*/stmt->base_ptr->as<ArgLoadStmt>()->arg_id);
+      }
+      auto check_upper_bound = new_stmts.push_back<BinaryOpStmt>(
+          BinaryOpType::cmp_lt, stmt->indices[i], upper_bound);
+      auto check_i = new_stmts.push_back<BinaryOpStmt>(
+          BinaryOpType::bit_and, check_lower_bound, check_upper_bound);
+      result = new_stmts.push_back<BinaryOpStmt>(BinaryOpType::bit_and, result,
+                                                 check_i);
+
+      auto input_index = stmt->indices[i];
+      args.emplace_back(input_index);
+    }
+
+    for (int i = 0; i < stmt->indices.size(); i++) {
+      if (i > 0)
+        msg += ", ";
+      msg += "%d";
+    }
+    msg += "]\n" + stmt->tb;
+
+    new_stmts.push_back<AssertStmt>(result, msg, args);
+    modifier.insert_before(stmt, std::move(new_stmts));
+    set_done(stmt);
+  }
+
   void visit(GlobalPtrStmt *stmt) override {
     if (is_done(stmt))
       return;


Test patch:
diff --git a/tests/python/test_ndarray.py b/tests/python/test_ndarray.py
index 43794649d4a73..856cf97eb2560 100644
--- a/tests/python/test_ndarray.py
+++ b/tests/python/test_ndarray.py
@@ -716,3 +716,57 @@ def test():
 
     for i in range(300):
         test()
+
+
+@test_utils.test(arch=supported_archs_taichi_ndarray,
+                 require=ti.extension.assertion,
+                 debug=True,
+                 check_out_of_bound=True,
+                 gdb_trigger=False)
+def test_scalar_ndarray_oob():
+    @ti.kernel
+    def access_arr(input: ti.types.ndarray(), x: ti.i32) -> ti.f32:
+        return input[x]
+
+    input = np.random.randn(4)
+
+    # Works
+    access_arr(input, 1)
+
+    with pytest.raises(AssertionError, match=r'Out of bound access'):
+        access_arr(input, 4)
+
+    with pytest.raises(AssertionError, match=r'Out of bound access'):
+        access_arr(input, -1)
+
+
+# SOA layout for ndarray is deprecated so no need to test
+@test_utils.test(arch=supported_archs_taichi_ndarray,
+                 require=ti.extension.assertion,
+                 debug=True,
+                 check_out_of_bound=True,
+                 gdb_trigger=False)
+def test_matrix_ndarray_oob():
+    @ti.kernel
+    def access_arr(input: ti.types.ndarray(), p: ti.i32, q: ti.i32, x: ti.i32,
+                   y: ti.i32) -> ti.f32:
+        return input[p, q][x, y]
+
+    input = ti.ndarray(dtype=ti.math.mat2, shape=(4, 5))
+
+    # Works
+    access_arr(input, 2, 3, 0, 1)
+
+    # element_shape
+    with pytest.raises(AssertionError, match=r'Out of bound access'):
+        access_arr(input, 2, 3, 2, 1)
+    # field_shape[0]
+    with pytest.raises(AssertionError, match=r'Out of bound access'):
+        access_arr(input, 4, 4, 0, 1)
+    with pytest.raises(AssertionError, match=r'Out of bound access'):
+        access_arr(input, -3, 4, 1, 1)
+    # field_shape[1]
+    with pytest.raises(AssertionError, match=r'Out of bound access'):
+        access_arr(input, 3, 5, 0, 1)
+    with pytest.raises(AssertionError, match=r'Out of bound access'):
+        access_arr(input, 2, -10, 1, 1)