[Lang] Subnormal floating point numbers are flushed to zero by pybind

[sjwsl]

Subnormal numbers (approximately 1e-45 to 1e-38) are flushed to zero by pybind now, which causes some confusing behaviors. For example

import taichi as ti
ti.init(arch=ti.gpu)


@ti.kernel
def constant_folded_on_py() -> ti.f32:
    return 1e-10 * 1e-30


@ti.kernel
def calculated_on_device() -> ti.f32:
    x = 1e-10
    y = 1e-30
    return x * y


print(1e-10 * 1e-30)
print(constant_folded_on_py())
print(calculated_on_device())

The output (same on the CUDA backend) is

[Taichi] version 0.8.6, llvm 10.0.0, commit cb5670e5, osx, python 3.9.7
[Taichi] Starting on arch=metal
1.0000000000000001e-40
0.0
9.99994610111476e-41

The different results of 1 and 3 are fine because they use 64-bit and 32-bit multiplication respectively. The problem is that 1e-40 (a subnormal number) on the Python side cannot be passed to the C++ side.

Currently, we set FTZ (Flush subnormal numbers To Zero) on the CUDA backend but it doesn't seem to work. The existing FTZ tests do not make sense too IMHO, because 1e-45 used in the tests is actually smaller than the minimal subnormal number so it can't tell whether FTZ works... And the CPU backend seems to correctly flush subnormal numbers to zero because the output is

[Taichi] version 0.8.6, llvm 10.0.0, commit cb5670e5, linux, python 3.9.7
[Taichi] Starting on arch=x64
1.0000000000000001e-40
0.0
0.0

I think we need a discussion about how to handle these small float numbers, the basic purpose is to make the behaviors consistent on

• different backends
• Python side and C++ side