Large (> 2**30) range for fails on Vulkan backend

[strongoier]

Describe the bug
When a range for has range > 2**30, it cannot work correctly under Vulkan backend.

To Reproduce

import taichi as ti

ti.init(arch=ti.vulkan, offline_cache=False)

@ti.kernel
def count() -> ti.i32:
    res = 2**30 + 1
    for i in range(res):
        res -= 1
    return res

print(count())  # -387, but it should be 0

Log/Screenshots

[Taichi] version 1.3.0, llvm 10.0.0, commit cd5a97c8, linux, python 3.8.10
[Taichi] Starting on arch=vulkan
[W 10/28/22 13:27:09.052 955659] [program.cpp:Program@160] Out-of-bound access checking is not supported on arch=vulkan
[W 10/28/22 13:27:09.064 955659] [vulkan_program.cpp:materialize_runtime@134] Enabling vulkan validation layer in debug mode
[W 10/28/22 13:27:09.066 955659] [vulkan_device_creator.cpp:vk_debug_callback@49] validation layer: 1, loader_add_to_layer_list: Explicit layer VK_LAYER_KHRONOS_validation is using an old API version 1.2 versus application requested 1.3
[W 10/28/22 13:27:09.071 955659] [vulkan_device_creator.cpp:vk_debug_callback@49] validation layer: 2, Validation Warning: [ VUID_Undefined ] Object 0: VK_NULL_HANDLE, type = VK_OBJECT_TYPE_INSTANCE; | MessageID = 0x79de34d4 | Unrecognized CreateInstance->pCreateInfo->pApplicationInfo.apiVersion number (0x00403000). Assuming VK_API_VERSION_1_2.
[I 10/28/22 13:27:09.080 955659] [vulkan_device_creator.cpp:pick_physical_device@399] Found Vulkan Device 0 (llvmpipe (LLVM 12.0.0, 256 bits))
[I 10/28/22 13:27:09.080 955659] [vulkan_device_creator.cpp:pick_physical_device@399] Found Vulkan Device 1 (NVIDIA GeForce RTX 3090)
[I 10/28/22 13:27:09.080 955659] [vulkan_device_creator.cpp:find_queue_families@148] Async compute queue 2, graphics queue 0
[I 10/28/22 13:27:09.081 955659] [vulkan_device_creator.cpp:find_queue_families@148] Async compute queue 2, graphics queue 0
[I 10/28/22 13:27:09.081 955659] [vulkan_device_creator.cpp:create_logical_device@469] Vulkan Device "NVIDIA GeForce RTX 3090" supports Vulkan 0 version 1.3.205

Additional comments
My first attempts of investigation reveals that this is not a numerical problem, as normal calculations for integers > 2**30 works fine. It seems related to the number of launched threads.