[cuda] warp_barrier intrinsic is never called when using ti.simt.warp.sync()

[YuCrazing]

Describe the bug
As an expected result, when we use ti.simt.warp.sync(), warp_barrier should be invoked at the backend. However, we can not find the warp_barrier intrinsic in the IR.

To Reproduce

Code:

import taichi as ti

ti.init(arch=ti.cuda, print_ir=True)

@ti.kernel
def test_kernel(arr: ti.template()):
    for I in ti.grouped(arr):
        ti.simt.warp.sync(ti.u32(0xFFFFFFFF))

arr = ti.field(ti.i32, 32)
test_kernel(arr)

IR info:

[Taichi] version 1.0.2, llvm 10.0.0, commit 6d538e17, linux, python 3.8.10
[Taichi] Starting on arch=cuda
[I 05/11/22 14:23:54.342 7682] [compile_to_offloads.cpp:operator()@22] [test_kernel_c56_0] Initial IR:
kernel {
  $0 : for @tmp0 in S2place<i32> noneblock_dim=adaptive {
    $1 = alloca @tmp1
    @tmp1 = (cast_value<i32> @tmp0)
  }
}
[I 05/11/22 14:23:54.342 7682] [compile_to_offloads.cpp:operator()@22] [test_kernel_c56_0] Lowered:
kernel {
  $0 : struct for in S1dense (bit_vectorize -1) noneblock_dim=adaptive {
    $1 = loop $0 index 0
    $2 = alloca
    $3 = cast_value<i32> $1
    $4 : local store [$2 <- $3]
  }
}

[strongoier]

The problem here is that internal func calls are expressions, and they will not be inserted into Taichi AST if they are not assigned to others. I'll help fix it.