SPIR-V backend: Type u8 not supported (but GPU/driver supports u8) #8758
Description
Hi,
thanks a lot for this great program! Here is a bug I found, would be amazing if anyone can help me out.
Describe the bug
My GPU/driver supports int8/int16 just fine.
The crash is from Taichi’s Vulkan SPIR-V backend: Type u8 not supported.
(So it’s Taichi’s codegen limitation on Vulkan, not my system.)
To Reproduce
import taichi as ti
ti.init(arch=ti.vulkan, debug=True) # prints device + enabled features
n = 16
x8 = ti.field(dtype=ti.u8, shape=n)
x16 = ti.field(dtype=ti.u16, shape=n)
@ti.kernel
def compute():
for i in range(n):
# exercise arithmetic (int8 native; int16 storage+arith)
x8[i] = (i * 7 + 3) # wraps modulo 256 automatically
x16[i] = i * 300 + 123
compute()
print("u8 :", x8.to_numpy())
print("u16:", x16.to_numpy())Log/Screenshots
[Taichi] version 1.7.4, llvm 15.0.1, commit b4b956fd, win, python 3.12.4
[Taichi] Starting on arch=vulkan
[W 08/25/25 15:39:52.544 14960] [program.cpp:taichi::lang::Program::Program@152] Out-of-bound access checking is not supported on arch=vulkan
[W 08/25/25 15:39:52.606 14960] [vulkan_program.cpp:taichi::lang::VulkanProgramImpl::materialize_runtime@113] Enabling vulkan validation layer in debug mode
RHI Error: Vulkan validation layer: 2, vkCreateDevice(): Internal Error, DebugPrintf is being disabled. Details:
GPU Shader Instrumentation requires timelineSemaphore to manage when command buffers are submitted at queue submit time.
[W 08/25/25 15:39:52.729 14960] TaichiWarning
While compiling `compute_c80_0`, File "taichi_test.py", line 13, in compute:
x8[i] = (i * 7 + 3) # wraps modulo 256 automatically
^^^^^^^^^^^^^^^^^^^^
Assign may lose precision: u8 <- i32
[W 08/25/25 15:39:52.729 14960] TaichiWarning
While compiling `compute_c80_0`, File "taichi_test.py", line 14, in compute:
x16[i] = i * 300 + 123
^^^^^^^^^^^^^^^^^^^^^^
Assign may lose precision: u16 <- i32
[E 08/25/25 15:39:52.736 14960] [spirv_ir_builder.cpp:taichi::lang::spirv::IRBuilder::get_primitive_type@317] Type u8 not supported.
***********************************
* Taichi Compiler Stack Traceback *
***********************************
0x7ff9da1978ad: taichi::Time::wait_until in taichi_python.cp312-win_amd64.pyd
0x7ff9da171fd0: taichi::lang::Texture::get_device_allocation_ptr_as_int in taichi_python.cp312-win_amd64.pyd
0x7ff9da171d83: taichi::lang::Texture::get_device_allocation_ptr_as_int in taichi_python.cp312-win_amd64.pyd
0x7ff9da1725da: taichi::lang::Texture::get_device_allocation_ptr_as_int in taichi_python.cp312-win_amd64.pyd
0x7ff9d9fbf9de: taichi::lang::Callable::pop_argpack_stack in taichi_python.cp312-win_amd64.pyd
0x7ff9db37f6a4: taichi::Logger::error in taichi_python.cp312-win_amd64.pyd
0x7ff9dade0984: taichi::lang::vulkan::VulkanDevice::wait_idle in taichi_python.cp312-win_amd64.pyd
0x7ff9dada45ca: taichi::lang::vulkan::VulkanDevice::wait_idle in taichi_python.cp312-win_amd64.pyd
0x7ff9d9f56ab0: taichi::lang::gfx::GfxRuntime::used_in_kernel in taichi_python.cp312-win_amd64.pyd
0x7ff9dad957d5: taichi::lang::vulkan::VulkanDevice::wait_idle in taichi_python.cp312-win_amd64.pyd
0x7ff9d9dc3bb0: taichi::lang::Texture::operator= in taichi_python.cp312-win_amd64.pyd
0x7ff9dad83e2c: taichi::lang::vulkan::VulkanDevice::wait_idle in taichi_python.cp312-win_amd64.pyd
0x7ff9dad8c0ef: taichi::lang::vulkan::VulkanDevice::wait_idle in taichi_python.cp312-win_amd64.pyd
0x7ff9dad8b86c: taichi::lang::vulkan::VulkanDevice::wait_idle in taichi_python.cp312-win_amd64.pyd
0x7ff9dad3dabe: taichi::lang::vulkan::VulkanDevice::wait_idle in taichi_python.cp312-win_amd64.pyd
0x7ff9dad18234: taichi::lang::vulkan::VulkanDevice::wait_idle in taichi_python.cp312-win_amd64.pyd
0x7ff9dad18077: taichi::lang::vulkan::VulkanDevice::wait_idle in taichi_python.cp312-win_amd64.pyd
0x7ff9dad1a2a8: taichi::lang::vulkan::VulkanDevice::wait_idle in taichi_python.cp312-win_amd64.pyd
0x7ff9da009a47: taichi::lang::Program::compile_kernel in taichi_python.cp312-win_amd64.pyd
0x7ff9da5fdc9a: taichi::ui::vulkan::Gui::render_pass in taichi_python.cp312-win_amd64.pyd
0x7ff9da3ffc22: taichi::ui::vulkan::Gui::render_pass in taichi_python.cp312-win_amd64.pyd
0x7ff9da3e534b: taichi::ui::vulkan::Gui::render_pass in taichi_python.cp312-win_amd64.pyd
0x7ff9da5e0239: taichi::ui::vulkan::Gui::render_pass in taichi_python.cp312-win_amd64.pyd
0x7ff9da385c8f: taichi::ui::vulkan::Gui::render_pass in taichi_python.cp312-win_amd64.pyd
0x7ff9da2dd2e5: taichi::Time::wait_until in taichi_python.cp312-win_amd64.pyd
0x7ffa15cb0913: PyCFunction_GetFlags in python312.dll
0x7ffa15c63f3d: PyObject_MakeTpCall in python312.dll
0x7ffa15c64265: PyObject_Vectorcall in python312.dll
0x7ffa15d9fcb3: PyEval_EvalFrameDefault in python312.dll
0x7ffa15c64534: PyFunction_Vectorcall in python312.dll
0x7ffa15c66b91: PyCell_Set in python312.dll
0x7ffa15c6722f: PyMethod_Self in python312.dll
0x7ffa15c641b3: PyVectorcall_Function in python312.dll
0x7ffa15c642f6: PyObject_Call in python312.dll
0x7ffa15d9ffac: PyEval_EvalFrameDefault in python312.dll
0x7ffa15c64534: PyFunction_Vectorcall in python312.dll
0x7ffa15c63bf4: PyObject_FastCallDictTstate in python312.dll
0x7ffa15c647f2: PyObject_Call_Prepend in python312.dll
0x7ffa15cdc9d1: PyType_Ready in python312.dll
0x7ffa15c64386: PyObject_Call in python312.dll
0x7ffa15d9ffac: PyEval_EvalFrameDefault in python312.dll
0x7ffa15d99606: PyEval_EvalCode in python312.dll
0x7ffa15e37026: PyRun_FileExFlags in python312.dll
0x7ffa15e37133: PyRun_FileExFlags in python312.dll
0x7ffa15e36d58: PyRun_StringFlags in python312.dll
0x7ffa15e33a41: PyRun_SimpleFileObject in python312.dll
0x7ffa15bdf1d1: Py_gitidentifier in python312.dll
0x7ffa15bdfe24: Py_gitidentifier in python312.dll
0x7ffa15be0348: Py_RunMain in python312.dll
0x7ffa15be03d2: Py_Main in python312.dll
0x7ff61a591490: OPENSSL_Applink in python.exe
0x7ffa7d337374: BaseThreadInitThunk in KERNEL32.DLL
0x7ffa7d63cc91: RtlUserThreadStart in ntdll.dll
Internal error occurred. Check out this page for possible solutions:
https://docs.taichi-lang.org/docs/install
Traceback (most recent call last):
File "taichi_test.py", line 16, in <module>
compute()
File "\.venv\Lib\site-packages\taichi\lang\kernel_impl.py", line 1113, in wrapped
return primal(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^^^
File "\.venv\Lib\site-packages\taichi\lang\kernel_impl.py", line 1045, in __call__
return self.launch_kernel(kernel_cpp, *args)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "\.venv\Lib\site-packages\taichi\lang\kernel_impl.py", line 976, in launch_kernel
raise e from None
File "\.venv\Lib\site-packages\taichi\lang\kernel_impl.py", line 969, in launch_kernel
compiled_kernel_data = prog.compile_kernel(prog.config(), prog.get_device_caps(), t_kernel)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
RuntimeError: [spirv_ir_builder.cpp:taichi::lang::spirv::IRBuilder::get_primitive_type@317] Type u8 not supported.
Additional comments
[Taichi] version 1.7.4, llvm 15.0.1, commit b4b956f, win, python 3.12.4
** Taichi Programming Language **
Docs: https://docs.taichi-lang.org/
GitHub: https://github.com/taichi-dev/taichi/
Forum: https://forum.taichi.graphics/
Taichi system diagnose:
python: 3.12.4 | packaged by Anaconda, Inc. | (main, Jun 18 2024, 15:03:56) [MSC v.1929 64 bit (AMD64)]
system: win32
executable: .venv\Scripts\python.exe
platform: Windows-10-10.0.19045-SP0
architecture: 64bit WindowsPE
uname: uname_result(system='Windows', node='UOW-C3XNZL2', release='10', version='10.0.19045', machine='AMD64')
.venv\Lib\site-packages\taichi\tools\diagnose.py:20: DeprecationWarning: 'locale.getdefaultlocale' is deprecated and slated for removal in Python 3.15. Use setlocale(), getencoding() and getlocale() instead.
print(f'locale: {".".join(locale.getdefaultlocale())}')
locale: en_AU.cp1252
PATH: ***
PYTHONPATH: ***
lsb_release not available: [WinError 2] The system cannot find the file specified
import: <module 'taichi' from '\.venv\Lib\site-packages\taichi\init.py'>
cpu: True
metal: False
opengl: True
[W 08/25/25 15:43:02.539 18180] [cuda_driver.cpp:taichi::lang::CUDADriverBase::load_lib@36] nvcuda.dll lib not found.
cuda: False
vulkan: True
glewinfo not available: [WinError 2] The system cannot find the file specified
nvidia-smi not available: [WinError 2] The system cannot find the file specified
[Taichi] version 1.7.4, llvm 15.0.1, commit b4b956f, win, python 3.12.4
[Taichi] version 1.7.4, llvm 15.0.1, commit b4b956f, win, python 3.12.4
[Taichi] Starting on arch=x64
[Taichi] version 1.7.4, llvm 15.0.1, commit b4b956f, win, python 3.12.4
[Taichi] Starting on arch=opengl
[W 08/25/25 15:43:07.709 17380] [cuda_driver.cpp:taichi::lang::CUDADriverBase::load_lib@36] nvcuda.dll lib not found.
[W 08/25/25 15:43:07.710 17380] [misc.py:adaptive_arch_select@758] Arch=[<Arch.cuda: 3>] is not supported, falling back to CPU
[Taichi] version 1.7.4, llvm 15.0.1, commit b4b956f, win, python 3.12.4
[Taichi] Starting on arch=x64
[Taichi] version 1.7.4, llvm 15.0.1, commit b4b956f, win, python 3.12.4
** Taichi Programming Language **
Docs: https://docs.taichi-lang.org/
GitHub: https://github.com/taichi-dev/taichi/
Forum: https://forum.taichi.graphics/
TAICHI EXAMPLES
+------------------------------------------------------------------------------------+
| 0: ad_gravity | 25: karman_vortex_street | 50: patterns |
| 1: circle_packing_image | 26: keyboard | 51: pbf2d |
| 2: comet | 27: laplace | 52: physarum |
| 3: cornell_box | 28: laplace_equation | 53: poisson_disk_sampling |
| 4: diff_sph | 29: mandelbrot_zoom | 54: print_offset |
| 5: differential_evolution | 30: marching_squares | 55: rasterizer |
| 6: euler | 31: mass_spring_3d_ggui | 56: regression |
| 7: eulerfluid2d | 32: mass_spring_game | 57: sdf_renderer |
| 8: explicit_activation | 33: mass_spring_game_ggui | 58: simple_derivative |
| 9: export_mesh | 34: mciso_advanced | 59: simple_texture |
| 10: export_ply | 35: mgpcg | 60: simple_uv |
| 11: export_videos | 36: mgpcg_advanced | 61: snow_phaseField |
| 12: fem128 | 37: minimal | 62: stable_fluid |
| 13: fem128_ggui | 38: minimization | 63: stable_fluid_ggui |
| 14: fem99 | 39: mpm128 | 64: stable_fluid_graph |
| 15: fractal | 40: mpm128_ggui | 65: taichi_bitmasked |
| 16: fractal3d_ggui | 41: mpm3d | 66: taichi_dynamic |
| 17: fullscreen | 42: mpm3d_ggui | 67: taichi_logo |
| 18: game_of_life | 43: mpm88 | 68: taichi_ngp |
| 19: gui_image_io | 44: mpm88_graph | 69: taichi_sparse |
| 20: gui_widgets | 45: mpm99 | 70: texture_graph |
| 21: implicit_fem | 46: mpm_lagrangian_forces | 71: tutorial |
| 22: implicit_mass_spring | 47: nbody | 72: two_stream_instability |
| 23: initial_value_problem | 48: odop_solar | 73: vortex_rings |
| 24: jacobian | 49: oit_renderer | 74: waterwave |
+------------------------------------------------------------------------------------+
42
Running example minimal ...
[Taichi] Starting on arch=x64
42.0
Running time: 0.36s
Consider attaching this log when maintainers ask about system information.
Running time: 12.91s