python使用cuda后端时kernel如果第一个参数是ti.types.ndarray()会报错

代码如下:

@ti.kernel
def ttt(b:ti.types.ndarray(),a:int):
    for i in range(a):
        pass

ti.init(ti.cuda)
ttt(np.zeros([5,5]),5)

报错信息如下:

      4         pass
      6 ti.init(ti.cuda)
----> 7 ttt(np.zeros([5,5]),5)

File d:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py:1023, in _kernel_impl..wrapped(*args, **kwargs)
   1020 @functools.wraps(_func)
   1021 def wrapped(*args, **kwargs):
   1022     try:
-> 1023         return primal(*args, **kwargs)
   1024     except (TaichiCompilationError, TaichiRuntimeError) as e:
   1025         raise type(e)('\n' + str(e)) from None

File d:\Programs\Python\Python38\lib\site-packages\taichi\lang\shell.py:27, in _shell_pop_print..new_call(*args, **kwargs)
     25 @functools.wraps(old_call)
     26 def new_call(*args, **kwargs):
---> 27     ret = old_call(*args, **kwargs)
     28     # print's in kernel won't take effect until ti.sync(), discussion:
     29     # https://github.com/taichi-dev/taichi/pull/1303#discussion_r444897102
     30     print(_ti_core.pop_python_print_buffer(), end='')

File d:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py:950, in Kernel.__call__(self, *args, **kwargs)
    948     impl.current_cfg().opt_level = 1
    949 key = self.ensure_compiled(*args)
--> 950 return self.runtime.compiled_functions[key](*args)

File d:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py:853, in Kernel.get_function_body..func__(*args)
    851 except Exception as e:
    852     e = handle_exception_from_cpp(e)
--> 853     raise e from None
    855 ret = None
    856 ret_dt = self.return_type

File d:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py:850, in Kernel.get_function_body..func__(*args)
    845     raise TaichiRuntimeError(
    846         f"The number of elements in kernel arguments is too big! Do not exceed 64 on {_ti_core.arch_name(impl.current_cfg().arch)} backend."
    847     )
    849 try:
--> 850     t_kernel(launch_ctx)
    851 except Exception as e:
    852     e = handle_exception_from_cpp(e)

RuntimeError: [jit_cuda.cpp:taichi::lang::convert@67] Assertion failure: std::isalpha(new_name[i]) || std::isdigit(new_name[i]) || new_name[i] == '_' || new_name[i] == '.'

[Taichi] version 1.5.0, llvm 15.0.1, commit 7b885c28, win, python 3.8.8
cuda11.6,驱动版本512.15

可以先试试把驱动升级到最新么? 我这里找不到相同版本的驱动,或者你可以试试 516 的驱动,这个驱动我们会日常测试,刚才试了下也没有重现你的问题

driver 链接私信你了

用了你发的526的驱动,代码如下

import taichi as ti
import numpy as np
@ti.kernel
def ttt(b:ti.types.ndarray(),a:int):
    for i in range(a):
        pass

ti.init(ti.cuda)
ttt(np.zeros([5,5]),5)

在命令行运行后输出为

[Taichi] version 1.5.0, llvm 15.0.1, commit 7b885c28, win, python 3.8.8
[Taichi] Starting on arch=cuda
[E 05/09/23 15:22:19.549 23020] [jit_cuda.cpp:taichi::lang::convert@67] Assertion failure: std::isalpha(new_name[i]) || std::isdigit(new_name[i]) || new_name[i] == '_' || new_name[i] == '.'


Traceback (most recent call last):
  File "e:/Programs/ttt.py", line 9, in <module>
    ttt(np.zeros([5,5]),5)
  File "D:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py", line 1023, in wrapped
    return primal(*args, **kwargs)
  File "D:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py", line 950, in __call__
    return self.runtime.compiled_functions[key](*args)
  File "D:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py", line 853, in func__
    raise e from None
  File "D:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py", line 850, in func__
    t_kernel(launch_ctx)
RuntimeError: [jit_cuda.cpp:taichi::lang::convert@67] Assertion failure: std::isalpha(new_name[i]) || std::isdigit(new_name[i]) || new_name[i] == '_' || new_name[i] == '.'

目前我在尝试使用cuda后端时还遇到了更多奇怪的bug,报错都是这个,可能多定义一个field都会报错,这导致目前cuda后端对我而言几乎不可用,似乎是同一个问题导致的?

很奇怪的是我把中间的pass换成对一个数累加就不报错了,代码如下:

import taichi as ti
import numpy as np
@ti.kernel
def ttt(b:ti.types.ndarray(),a:int):
    for i in range(a):
        c[None]+=1

ti.init(ti.cuda)
c=ti.field(dtype=int, shape=())
ttt(np.zeros([5,5]),5)
print(c[None])

输出如下:

[Taichi] version 1.5.0, llvm 15.0.1, commit 7b885c28, win, python 3.8.8
[Taichi] Starting on arch=cuda
5

正在做一个 fix,估计2天后可以出现在 nightly 里面,麻烦先等等

好的,我又发现一个类似的bug,代码如下:

import taichi as ti
@ti.kernel
def ttt(m:int,n:int):
    for i in range(n):
        for j in range(m):
            t[j,i]=ti.cast(0,ti.i8)
ti.init(ti.cuda)
a=2000
b=41
c=5
t=ti.field(dtype=ti.i8,shape=(a,b))
ttt(a,c)
print(t.to_numpy())

输出的报错和上面几乎一样:

[Taichi] version 1.5.0, llvm 15.0.1, commit 7b885c28, win, python 3.8.8
[Taichi] Starting on arch=cuda
[E 05/09/23 18:02:43.373 4960] [jit_cuda.cpp:taichi::lang::convert@67] Assertion failure: std::isalpha(new_name[i]) || std::isdigit(new_name[i]) || new_name[i] == '_' || new_name[i] == '.'


Traceback (most recent call last):
  File "e:/Programs/more_field_en.py", line 12, in <module>
    ttt(a,c)
  File "D:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py", line 1023, in wrapped
    return primal(*args, **kwargs)
  File "D:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py", line 950, in __call__
    return self.runtime.compiled_functions[key](*args)
  File "D:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py", line 853, in func__
    raise e from None
  File "D:\Programs\Python\Python38\lib\site-packages\taichi\lang\kernel_impl.py", line 850, in func__
    t_kernel(launch_ctx)
RuntimeError: [jit_cuda.cpp:taichi::lang::convert@67] Assertion failure: std::isalpha(new_name[i]) || std::isdigit(new_name[i]) || new_name[i] == '_' || new_name[i] == '.'

奇怪的是我把b=41换成b=42,或者把t[j,i]=ti.cast(0,ti.i8)改成t[j,i]=ti.cast(1,ti.i8)就不报错了
例如以下代码:

import taichi as ti
@ti.kernel
def ttt(m:int,n:int):
    for i in range(n):
        for j in range(m):
            t[j,i]=ti.cast(1,ti.i8)
ti.init(ti.cuda)
a=2000
b=41
c=5
t=ti.field(dtype=ti.i8,shape=(a,b))
ttt(a,c)
print(t.to_numpy())

输出为

[Taichi] version 1.5.0, llvm 15.0.1, commit 7b885c28, win, python 3.8.8
[Taichi] Starting on arch=cuda
[[1 1 1 ... 0 0 0]
 [1 1 1 ... 0 0 0]
 [1 1 1 ... 0 0 0]
 ...
 [1 1 1 ... 0 0 0]
 [1 1 1 ... 0 0 0]
 [1 1 1 ... 0 0 0]]

不清楚造成这些bug的原因是否一致

能不能试试今天的 nightly ?

pip install -i https://pypi.taichi.graphics/simple/ taichi-nightly

上面的问题解决了。

不过测试时发现了另一个问题,taichi的cpu后端和opengl后端都是支持非英文kernel名的,而cuda后端不支持,以前版本在切换到cuda后端运行报错后可以直接改名解决,但nightly版本似乎会缓存编译结果,导致如果使用非英文的kernel名,就算改成英文kernel名后也依然会报错,我没找到清空编译缓存的方法只能通过修改kernel解决。
测试代码如下:

import taichi as ti
@ti.kernel
def 测试():
    print(1)

ti.init(ti.cuda)
try:
    测试()
except RuntimeError as e:
    print(e)

@ti.kernel
def ttt():
    print(1)

ttt()

报错输出为:

[Taichi] version 1.7.0, llvm 15.0.1, commit 5ac2756a, win, python 3.10.0
[Taichi] Starting on arch=cuda
[W 05/10/23 13:08:13.087 18256] [taichi/rhi/cuda/cuda_driver.h:taichi::lang::CUDADriverFunction<void * *,void *,char const *>::call_with_warning@85] CUDA Error CUDA_ERROR_NOT_FOUND: named symbol not found while calling module_get_function (cuModuleGetFunction)
[E 05/10/23 13:08:13.087 18256] [taichi/runtime/cuda/jit_cuda.h:taichi::lang::JITModuleCUDA::lookup_function@54] Cannot look up function 测试_c74_0_kernel_0_serial


[taichi/runtime/cuda/jit_cuda.h:taichi::lang::JITModuleCUDA::lookup_function@54] Cannot look up function 测试_c74_0_kernel_0_serial
[W 05/10/23 13:08:13.089 18256] [taichi/rhi/cuda/cuda_driver.h:taichi::lang::CUDADriverFunction<void * *,void *,char const *>::call_with_warning@85] CUDA Error CUDA_ERROR_NOT_FOUND: named symbol not found while calling module_get_function (cuModuleGetFunction)
[E 05/10/23 13:08:13.090 18256] [taichi/runtime/cuda/jit_cuda.h:taichi::lang::JITModuleCUDA::lookup_function@54] Cannot look up function 测试_c74_0_kernel_0_serial


Traceback (most recent call last):
  File "e:\Programs\m.py", line 16, in <module>
    ttt()
  File "D:\ProgramData\Miniconda3\envs\taichi\lib\site-packages\taichi\lang\kernel_impl.py", line 962, in wrapped
    return primal(*args, **kwargs)
  File "D:\ProgramData\Miniconda3\envs\taichi\lib\site-packages\taichi\lang\kernel_impl.py", line 894, in __call__
    return self.launch_kernel(kernel_cpp, *args)
  File "D:\ProgramData\Miniconda3\envs\taichi\lib\site-packages\taichi\lang\kernel_impl.py", line 806, in launch_kernel
    raise e from None
  File "D:\ProgramData\Miniconda3\envs\taichi\lib\site-packages\taichi\lang\kernel_impl.py", line 803, in launch_kernel
    prog.launch_kernel(compiled_kernel_data, launch_ctx)
RuntimeError: [taichi/runtime/cuda/jit_cuda.h:taichi::lang::JITModuleCUDA::lookup_function@54] Cannot look up function 测试_c74_0_kernel_0_serial

性能调优 | Taichi Docs 参考下这个把

关掉离线缓存后问题解决了,不过cuda后端下有没有可能支持非英文kernel名?

在意这个的话麻烦在 github 提个 issue 吧,不过这个优先级应该不会很高

明白