网格与块规格#
grid
,又称循环中的核函数#
当使用 jax.experimental.pallas.pallas_call()
时,核函数会根据 pallas_call
的 grid
参数的指定,在不同的输入上多次执行。从概念上讲:
pl.pallas_call(some_kernel, grid=(n,))(...)
映射到
for i in range(n):
some_kernel(...)
网格可以泛化为多维的,对应于嵌套循环。例如:
pl.pallas_call(some_kernel, grid=(n, m))(...)
等价于
for i in range(n):
for j in range(m):
some_kernel(...)
这泛化到任何整数元组(长度为 d
的网格将对应 d
个嵌套循环)。核函数会执行 prod(grid)
次。默认的网格值 ()
会导致一次核函数调用。这些调用中的每一个都称为一个“程序”。为了访问核函数当前正在执行哪个程序(即网格的哪个元素),我们使用 jax.experimental.pallas.program_id()
。例如,对于调用 (1, 2)
,program_id(axis=0)
返回 1
,而 program_id(axis=1)
返回 2
。您还可以使用 jax.experimental.pallas.num_programs()
获取给定轴的网格大小。
有关使用此 API 的简单核函数,请参见 网格示例。
BlockSpec
,又称如何分块输入#
结合 grid
参数,我们需要向 Pallas 提供如何为每次调用切分输入的信息。具体来说,我们需要提供一个从*循环迭代*到*要操作的输入和输出块*的映射。这通过 jax.experimental.pallas.BlockSpec
对象提供。
在我们深入了解 BlockSpec
的细节之前,您可能希望回顾 Pallas 快速入门中的 块规格示例。
通过 in_specs
和 out_specs
将 BlockSpec
提供给 pallas_call
,每个输入和输出各一个。
首先,我们讨论当 indexing_mode == pl.Blocked()
时 BlockSpec
的语义。
非正式地讲,BlockSpec
的 index_map
将调用索引(数量与 grid
元组的长度相同)作为参数,并返回**块索引**(整体数组的每个轴一个块索引)。然后,每个块索引会乘以 block_shape
中对应的轴大小,以获取对应数组轴上的实际元素索引。
注意
并非所有块形状都受支持。
在 TPU 上,仅支持秩至少为 1 的块。此外,您的块形状的最后两个维度必须等于整体数组的相应维度,或者分别可被 8 和 128 整除。对于秩为 1 的块,块维度必须等于数组维度,或者可被
128 * (32 / bitwidth(dtype))
整除。在 GPU 上,当使用 Mosaic GPU 后端时,块的大小不受限制。但是,由于硬件限制,最小数组维度的大小必须是 16 字节的倍数。例如,如果输入是
jnp.float16
,则它必须是 8 的倍数。在 GPU 上,当使用 Triton 后端时,块本身的大小不受限制,但每个操作(包括加载或存储)必须在大小为 2 的幂的数组上进行操作。
如果块形状不能将整体形状均匀划分,那么每个轴上的最后一次迭代仍会收到对 block_shape
块的引用,但超出边界的元素会在输入时进行填充,在输出时被丢弃。填充值未指定,您应假定它们是垃圾数据。在 interpret=True
模式下,我们对浮点值使用 NaN 进行填充,以便用户有机会发现访问越界元素,但此行为不应依赖。请注意,每个块中至少有一个元素必须在界限内。
更准确地说,输入 x
(形状为 x_shape
)每个轴的切片计算方式如下面的函数 slice_for_invocation
所示
>>> import jax
>>> from jax.experimental import pallas as pl
>>> def slices_for_invocation(x_shape: tuple[int, ...],
... x_spec: pl.BlockSpec,
... grid: tuple[int, ...],
... invocation_indices: tuple[int, ...]) -> tuple[slice, ...]:
... assert len(invocation_indices) == len(grid)
... assert all(0 <= i < grid_size for i, grid_size in zip(invocation_indices, grid))
... block_indices = x_spec.index_map(*invocation_indices)
... assert len(x_shape) == len(x_spec.block_shape) == len(block_indices)
... elem_indices = []
... for x_size, block_size, block_idx in zip(x_shape, x_spec.block_shape, block_indices):
... start_idx = block_idx * block_size
... # At least one element of the block must be within bounds
... assert start_idx < x_size
... elem_indices.append(slice(start_idx, start_idx + block_size))
... return elem_indices
例如
>>> slices_for_invocation(x_shape=(100, 100),
... x_spec = pl.BlockSpec((10, 20), lambda i, j: (i, j)),
... grid = (10, 5),
... invocation_indices = (2, 4))
[slice(20, 30, None), slice(80, 100, None)]
>>> # Same shape of the array and blocks, but we iterate over each block 4 times
>>> slices_for_invocation(x_shape=(100, 100),
... x_spec = pl.BlockSpec((10, 20), lambda i, j, k: (i, j)),
... grid = (10, 5, 4),
... invocation_indices = (2, 4, 0))
[slice(20, 30, None), slice(80, 100, None)]
>>> # An example when the block is partially out-of-bounds in the 2nd axis.
>>> slices_for_invocation(x_shape=(100, 90),
... x_spec = pl.BlockSpec((10, 20), lambda i, j: (i, j)),
... grid = (10, 5),
... invocation_indices = (2, 4))
[slice(20, 30, None), slice(80, 100, None)]
下面定义的函数 show_program_ids
使用 Pallas 来显示调用索引。iota_2D_kernel
将每个输出块填充一个十进制数,其中第一位数字表示第一个轴上的调用索引,第二位数字表示第二个轴上的调用索引。
>>> def show_program_ids(x_shape, block_shape, grid,
... index_map=lambda i, j: (i, j)):
... def program_ids_kernel(o_ref): # Fill the output block with 10*program_id(1) + program_id(0)
... axes = 0
... for axis in range(len(grid)):
... axes += pl.program_id(axis) * 10**(len(grid) - 1 - axis)
... o_ref[...] = jnp.full(o_ref.shape, axes)
... res = pl.pallas_call(program_ids_kernel,
... out_shape=jax.ShapeDtypeStruct(x_shape, dtype=np.int32),
... grid=grid,
... in_specs=[],
... out_specs=pl.BlockSpec(block_shape, index_map),
... interpret=True)()
... print(res)
例如
>>> show_program_ids(x_shape=(8, 6), block_shape=(2, 3), grid=(4, 2),
... index_map=lambda i, j: (i, j))
[[ 0 0 0 1 1 1]
[ 0 0 0 1 1 1]
[10 10 10 11 11 11]
[10 10 10 11 11 11]
[20 20 20 21 21 21]
[20 20 20 21 21 21]
[30 30 30 31 31 31]
[30 30 30 31 31 31]]
>>> # An example with out-of-bounds accesses
>>> show_program_ids(x_shape=(7, 5), block_shape=(2, 3), grid=(4, 2),
... index_map=lambda i, j: (i, j))
[[ 0 0 0 1 1]
[ 0 0 0 1 1]
[10 10 10 11 11]
[10 10 10 11 11]
[20 20 20 21 21]
[20 20 20 21 21]
[30 30 30 31 31]]
>>> # It is allowed for the shape to be smaller than block_shape
>>> show_program_ids(x_shape=(1, 2), block_shape=(2, 3), grid=(1, 1),
... index_map=lambda i, j: (i, j))
[[0 0]]
当多个调用写入输出数组的相同元素时,结果是平台相关的。
在下面的示例中,我们有一个 3D 网格,其中最后一个网格维度未用于块选择(index_map=lambda i, j, k: (i, j)
)。因此,我们迭代同一个输出块 10 次。下面显示的输出是在 CPU 上使用 interpret=True
模式生成的,该模式目前按顺序执行调用。在 TPU 上,程序以并行和顺序结合的方式执行,并且此函数会生成所示的输出。参见 值得注意的属性和限制。
>>> show_program_ids(x_shape=(8, 6), block_shape=(2, 3), grid=(4, 2, 10),
... index_map=lambda i, j, k: (i, j))
[[ 9 9 9 19 19 19]
[ 9 9 9 19 19 19]
[109 109 109 119 119 119]
[109 109 109 119 119 119]
[209 209 209 219 219 219]
[209 209 209 219 219 219]
[309 309 309 319 319 319]
[309 309 309 319 319 319]]
None
值作为维度值出现在 block_shape
中时,其行为与值 1
相同,不同之处在于相应的块轴会被压缩(您也可以传入 pl.Squeezed()
而非 None
)。在下面的示例中,请注意当块形状被指定为 (None, 2)
时(主导维度被压缩),o_ref
的形状是 (2,)。
>>> def kernel(o_ref):
... assert o_ref.shape == (2,)
... o_ref[...] = jnp.full((2,), 10 * pl.program_id(1) + pl.program_id(0))
>>> pl.pallas_call(kernel,
... jax.ShapeDtypeStruct((3, 4), dtype=np.int32),
... out_specs=pl.BlockSpec((None, 2), lambda i, j: (i, j)),
... grid=(3, 2), interpret=True)()
Array([[ 0, 0, 10, 10],
[ 1, 1, 11, 11],
[ 2, 2, 12, 12]], dtype=int32)
当我们构造一个 BlockSpec
时,我们可以对 block_shape
参数使用 None
值,在这种情况下,整体数组的形状被用作 block_shape
。如果我们将 None
值用于 index_map
参数,则会使用一个默认的索引映射函数,该函数返回一个由零组成的元组:index_map=lambda *invocation_indices: (0,) * len(block_shape)
。
>>> show_program_ids(x_shape=(4, 4), block_shape=None, grid=(2, 3),
... index_map=None)
[[12 12 12 12]
[12 12 12 12]
[12 12 12 12]
[12 12 12 12]]
>>> show_program_ids(x_shape=(4, 4), block_shape=(4, 4), grid=(2, 3),
... index_map=None)
[[12 12 12 12]
[12 12 12 12]
[12 12 12 12]
[12 12 12 12]]
“元素”索引模式#
上述行为适用于默认的“分块”索引模式。当 block_shape
元组中使用整数时,例如 (4, 8)
,它等效于传入 pl.Blocked(block_size)
对象,例如 (pl.Blocked(4), pl.Blocked(8))
。分块索引模式意味着 index_map
返回的索引是*块索引*。我们可以传入 pl.Blocked
以外的对象来改变 index_map
的语义,最值得注意的是 pl.Element(block_size)
。当使用 pl.Element
索引模式时,索引映射函数返回的值直接用作数组索引,而无需先按块大小进行缩放。当使用 pl.Element
模式时,您可以将数组的虚拟填充指定为维度的低-高填充元组:其行为就像整体数组在输入时进行了填充。与当块形状不能整除整体数组形状时分块索引模式的填充值类似,元素模式的填充值不作任何保证。
目前,Element
模式仅在 TPU 上支持。
>>> # element without padding
>>> show_program_ids(x_shape=(8, 6), block_shape=(pl.Element(2), pl.Element(3)),
... grid=(4, 2),
... index_map=lambda i, j: (2*i, 3*j))
[[ 0 0 0 1 1 1]
[ 0 0 0 1 1 1]
[10 10 10 11 11 11]
[10 10 10 11 11 11]
[20 20 20 21 21 21]
[20 20 20 21 21 21]
[30 30 30 31 31 31]
[30 30 30 31 31 31]]
>>> # element, first pad the array with 1 row and 2 columns.
>>> show_program_ids(x_shape=(7, 7),
... block_shape=(pl.Element(2, (1, 0)),
... pl.Element(3, (2, 0))),
... grid=(4, 3),
... index_map=lambda i, j: (2*i, 3*j))
[[ 0 1 1 1 2 2 2]
[10 11 11 11 12 12 12]
[10 11 11 11 12 12 12]
[20 21 21 21 22 22 22]
[20 21 21 21 22 22 22]
[30 31 31 31 32 32 32]
[30 31 31 31 32 32 32]]