[SIMT] Add activemask warp intrinsics (#4918)

* add activemask warp intrinsic

add test function call

del extra space

unit-test print->assert

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

python/taichi/lang/simt/warp.py

@@ -98,8 +98,9 @@ def match_all():
 
 
 def active_mask():
-    # TODO
-    pass
+    return expr.Expr(
+        _ti_core.insert_internal_func_call("cuda_active_mask",
+                                           expr.make_expr_group(), False))
 
 
 def sync(mask):

tests/python/test_simt.py

@@ -251,8 +251,18 @@ def test_match_all():
 
 @test_utils.test(arch=ti.cuda)
 def test_active_mask():
-    # TODO
-    pass
+    a = ti.field(dtype=ti.u32, shape=32)
+
+    @ti.kernel
+    def foo():
+        ti.loop_config(block_dim=16)
+        for i in range(32):
+            a[i] = ti.simt.warp.active_mask()
+
+    foo()
+
+    for i in range(32):
+        assert a[i] == 65535
 
 
 @test_utils.test(arch=ti.cuda)