From 3dfb93d8984cf758ff24ef3c525ffc44427c7ca7 Mon Sep 17 00:00:00 2001 From: Bo Qiao Date: Tue, 5 Jul 2022 17:51:58 +0800 Subject: [PATCH] Add simt functions in operators --- docs/lang/articles/basic/operator.md | 49 ++++++++++++++++++++++++++++ 1 file changed, 49 insertions(+) diff --git a/docs/lang/articles/basic/operator.md b/docs/lang/articles/basic/operator.md index b7031fa07f309..85e00f543d064 100644 --- a/docs/lang/articles/basic/operator.md +++ b/docs/lang/articles/basic/operator.md @@ -323,3 +323,52 @@ a@a # @ denotes matrix multiplication For now, determinant() and inverse() only works in Taichi-scope, and the size of the matrix must be 1x1, 2x2, 3x3 or 4x4. ::: + + +## Supported SIMT intrinsics +For CUDA backend, Taichi now supports warp-level and block-level intrinsics that +are needed for writing high-performance SIMT kernels. You can use them in Taichi +similar to the [usage in CUDA kernels](https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/). Currently, the following functions are supported: + + +| Operation | Mapped CUDA intrinsic | +| -------------------------- | ----------------------------- | +|`ti.simt.warp.all_nonzero` | `__all_sync` | +|`ti.simt.warp.any_nonzero` | `__any_sync` | +|`ti.simt.warp.unique` | `__uni_sync` | +|`ti.simt.warp.ballot` | `__ballot_sync` | +|`ti.simt.warp.shfl_sync_i32`| `__shfl_sync` | +|`ti.simt.warp.shfl_sync_f32`| `__shfl_sync` | +|`ti.simt.warp.shfl_up_i32` | `__shfl_up_sync` | +|`ti.simt.warp.shfl_up_f32` | `__shfl_up_sync` | +|`ti.simt.warp.shfl_down_i32`| `__shfl_down_sync`| +|`ti.simt.warp.shfl_down_f32`| `__shfl_down_sync`| +|`ti.simt.warp.shfl_xor_i32` | `__shfl_xor_sync` | +|`ti.simt.warp.match_any` | `__match_any_sync`| +|`ti.simt.warp.match_all` | `__match_all_sync`| +|`ti.simt.warp.active_mask` | `__activemask` | +|`ti.simt.warp.sync` | `__syncwarp` | + +Please refer to our [API docs](https://docs.taichi.graphics/api/taichi/lang/simt/warp/#module-taichi.lang.simt.warp) +for more information on each function. + +Here is an example to perform data exchange within a warp in Taichi: + + +```python +a = ti.field(dtype=ti.i32, shape=32) + +@ti.kernel +def foo(): + ti.loop_config(block_dim=32) + for i in range(32): + a[i] = ti.simt.warp.shfl_up_i32(ti.u32(0xFFFFFFFF), a[i], 1) + +for i in range(32): + a[i] = i * i + +foo() + +for i in range(1, 32): + assert a[i] == (i - 1) * (i - 1) +```