Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

SPIR-V backend does not support float immediate number of f16 type #4757

Closed
lin-hitonami opened this issue Apr 11, 2022 · 1 comment · Fixed by #4787
Closed

SPIR-V backend does not support float immediate number of f16 type #4757

lin-hitonami opened this issue Apr 11, 2022 · 1 comment · Fixed by #4787
Labels
potential bug Something that looks like a bug but not yet confirmed vulkan Vulkan backend

Comments

@lin-hitonami
Copy link
Contributor

lin-hitonami commented Apr 11, 2022 •
edited
Loading

Describe the bug
When fixing #4741, I found that SPIR-V does not support float immediate number of f16 type, and test_f16.py::test_fractal_f16 fails.

Note: remember to add dtype=ti.f16 to the test after fixing this issue.
To Reproduce
This is an altered version of test_f16.py::test_fractal_f16 that all vectors store f16 values.

import taichi as ti
ti.init(ti.vulkan)

def test_fractal_f16():
    n = 320
    pixels = ti.field(dtype=ti.f16, shape=(n * 2, n))

    @ti.func
    def complex_sqr(z):
        return ti.Vector([z[0]**2 - z[1]**2, z[1] * z[0] * 2])

    @ti.kernel
    def paint(t: float):
        for i, j in pixels:  # Parallelized over all pixels
            c0 = ti.cast(-0.8, ti.f16)
            c1 = ti.cast(ti.cos(t) * 0.2, ti.f16)
            c = ti.Vector([-0.8, ti.cos(t) * 0.2])
            z0 = ti.cast(i / n - 1, ti.f16)
            z1 = ti.cast(j / n - 0.5, ti.f16)
            z = ti.Vector([z0, z1]) * 2
            iterations = 0
            while z.norm() < 20 and iterations < 50:
                z = complex_sqr(z) + c
                iterations += 1
            pixels[i, j] = 1 - iterations * 0.02


    paint(0.03)

test_fractal_f16()

Log/Screenshots
Please post the full log of the program (instead of just a few lines around the error message, unless the log is > 1000 lines). This will help us diagnose what's happening. For example:

[Taichi] Starting on arch=vulkan
[I 04/11/22 11:03:26.799 1318081] [vulkan_device_creator.cpp:pick_physical_device@356] Found Vulkan Device 0 (NVIDIA GeForce RTX 2060)
[I 04/11/22 11:03:26.799 1318081] [vulkan_device_creator.cpp:pick_physical_device@356] Found Vulkan Device 1 (llvmpipe (LLVM 12.0.0, 256 bits))
[I 04/11/22 11:03:26.799 1318081] [vulkan_device_creator.cpp:create_logical_device@424] Vulkan Device "NVIDIA GeForce RTX 2060" supports Vulkan 0 version 1.2.175
-------------------------------------------------------------------------------------------------------------------------------------------- Captured stdout call ---------------------------------------------------------------------------------------------------------------------------------------------
[E 04/11/22 11:03:26.854 1318081] [spirv_ir_builder.cpp:float_immediate_number@217] Type f16 not supported.


***********************************
* Taichi Compiler Stack Traceback *
***********************************
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so: taichi::Logger::error(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, bool)
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x5c9245) [0x7fcdd6257245]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x5b9d56) [0x7fcdd6247d56]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x5808f6) [0x7fcdd620e8f6]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x576c38) [0x7fcdd6204c38]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x57a003) [0x7fcdd6208003]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x576c38) [0x7fcdd6204c38]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x5ac4e4) [0x7fcdd623a4e4]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x572829) [0x7fcdd6200829]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x56f8d5) [0x7fcdd61fd8d5]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0xa32ff6) [0x7fcdd66c0ff6]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0xa8a468) [0x7fcdd6718468]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0xa8a799) [0x7fcdd6718799]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so: taichi::lang::Program::compile(taichi::lang::Kernel&, taichi::lang::OffloadedStmt*)
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so: taichi::lang::Kernel::compile()
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so: taichi::lang::Kernel::operator()(taichi::lang::Kernel::LaunchContextBuilder&)
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x3ec55a) [0x7fcdd607a55a]
/home/lin/taichi2/python/taichi/_lib/core/taichi_core.so(+0x3ac335) [0x7fcdd603a335]
/home/lin/anaconda3/envs/test39/bin/python(+0x174714) [0x5599efe4f714]
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_MakeTpCall
/home/lin/anaconda3/envs/test39/bin/python(+0xd5df7) [0x5599efdb0df7]
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python(+0xb022e) [0x5599efd8b22e]
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_MakeTpCall
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x1828e2) [0x5599efe5d8e2]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_FastCallDictTstate
/home/lin/anaconda3/envs/test39/bin/python(+0x17e275) [0x5599efe59275]
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xfaebb) [0x5599efdd5ebb]
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xf9fdc) [0x5599efdd4fdc]
/home/lin/anaconda3/envs/test39/bin/python(+0x1ae074) [0x5599efe89074]
/home/lin/anaconda3/envs/test39/bin/python(+0xf9fdc) [0x5599efdd4fdc]
/home/lin/anaconda3/envs/test39/bin/python(+0x1828e2) [0x5599efe5d8e2]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_FastCallDictTstate
/home/lin/anaconda3/envs/test39/bin/python(+0x17e275) [0x5599efe59275]
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_MakeTpCall
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xf910d) [0x5599efdd410d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xf9fdc) [0x5599efdd4fdc]
/home/lin/anaconda3/envs/test39/bin/python(+0x1ae074) [0x5599efe89074]
/home/lin/anaconda3/envs/test39/bin/python(+0xf9fdc) [0x5599efdd4fdc]
/home/lin/anaconda3/envs/test39/bin/python(+0x1828e2) [0x5599efe5d8e2]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_FastCallDictTstate
/home/lin/anaconda3/envs/test39/bin/python(+0x17e275) [0x5599efe59275]
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xfaebb) [0x5599efdd5ebb]
/home/lin/anaconda3/envs/test39/bin/python(+0x1af0ed) [0x5599efe8a0ed]
/home/lin/anaconda3/envs/test39/bin/python(+0xfa90b) [0x5599efdd590b]
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x1828e2) [0x5599efe5d8e2]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xfaebb) [0x5599efdd5ebb]
/home/lin/anaconda3/envs/test39/bin/python(+0x1828e2) [0x5599efe5d8e2]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xfa90b) [0x5599efdd590b]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xf9fdc) [0x5599efdd4fdc]
/home/lin/anaconda3/envs/test39/bin/python(+0x1ae074) [0x5599efe89074]
/home/lin/anaconda3/envs/test39/bin/python(+0xf9fdc) [0x5599efdd4fdc]
/home/lin/anaconda3/envs/test39/bin/python(+0x1828e2) [0x5599efe5d8e2]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_FastCallDictTstate
/home/lin/anaconda3/envs/test39/bin/python(+0x17e275) [0x5599efe59275]
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_MakeTpCall
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python(+0x182b0d) [0x5599efe5db0d]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xf9fdc) [0x5599efdd4fdc]
/home/lin/anaconda3/envs/test39/bin/python(+0x1ae074) [0x5599efe89074]
/home/lin/anaconda3/envs/test39/bin/python(+0xf9fdc) [0x5599efdd4fdc]
/home/lin/anaconda3/envs/test39/bin/python(+0x1828e2) [0x5599efe5d8e2]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_FastCallDictTstate
/home/lin/anaconda3/envs/test39/bin/python(+0x17e275) [0x5599efe59275]
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_MakeTpCall
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xfaebb) [0x5599efdd5ebb]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python(+0xfaebb) [0x5599efdd5ebb]
/home/lin/anaconda3/envs/test39/bin/python: _PyFunction_Vectorcall
/home/lin/anaconda3/envs/test39/bin/python: _PyObject_Call
/home/lin/anaconda3/envs/test39/bin/python: _PyEval_EvalFrameDefault
/home/lin/
=========================================================================================================================================== short test summary info ===========================================================================================================================================
FAILED test_f16.py::test_fractal_f16[arch=vulkan-4] - RuntimeError: [spirv_ir_builder.cpp:float_immediate_number@217] Type f16 not supported.
@lin-hitonami lin-hitonami added the potential bug Something that looks like a bug but not yet confirmed label Apr 11, 2022
@lin-hitonami lin-hitonami mentioned this issue Apr 11, 2022
Merged
@lin-hitonami lin-hitonami changed the title SPIR-V does not support float immediate number of f16 type SPIR-V backend does not support float immediate number of f16 type Apr 11, 2022
@erizmr
Copy link
Contributor

erizmr commented Apr 12, 2022

cc: @bobcao3

@erizmr erizmr added the vulkan Vulkan backend label Apr 12, 2022
@bobcao3 bobcao3 mentioned this issue Apr 14, 2022
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
potential bug Something that looks like a bug but not yet confirmed vulkan Vulkan backend
Projects
None yet
Milestone
No milestone
Development

Successfully merging a pull request may close this issue.

[spirv] Add 16 bit float immediate number bobcao3/taichi
2 participants
@erizmr @lin-hitonami