本文档介绍了 HLO 编号分析,可让您对 HLO 运算符符号化地计算编号映射。索引映射是一个函数,用于将 一个张量的索引映射到另一个张量的索引,例如HLO 的索引 为 HLO 指令输入的索引,反之亦然。
示例
对于从 tensor<20xf32>
到 tensor<10x20x30xf32>
的广播
p0 = f32[20] parameter(0)
bc0 = f32[10, 20, 30] broadcast(p0), dimensions={1}
从输出到输入的编号映射为 (i, j, k) -> (j)
(对于 i in
[0, 10]
、j in [0, 20]
和 k in [0, 30]
)。
设计初衷
XLA GPU 使用多种定制解决方案来推理合并、运算数 利用率和平铺方案(详见下文)。编制索引分析的目标是为此类用例提供可重复使用的组件。索引分析 基于 MLIR 的仿射地图基础架构构建,并添加了 HLO 语义。
合并
当我们知道读取输入的哪些元素/slice 来计算输出的元素时,对于非琐碎情况,推理内存合并就变得可行。
操作数利用率
XLA 中的操作数利用率表示指令中每个输入的 (假设其输出已被完全使用)。目前,利用率也不 针对一般情况计算的。借助索引分析,您可以精确计算利用率。
平铺:
功能块/切片是张量的一个超矩形子集,由偏移量、大小和步长参数化。图块传播是一种计算 使用操作本身的图块参数的操作的生产方/使用方。那里 已经是 库 它会针对 softmax 和 dot 执行。图块传播可以变得更加通用, 如果通过索引映射来表示,则是可靠的。
函数和域
编制索引映射的函数为 f(x) = f(d, r, rt),用于将张量 A
的多重索引 d 映射到张量 B
的元素/范围。参数r表示
张量 B
中存在的维度,但张量 A
中不存在的维度。通过
参数 rt 指的是运行时值,如收集操作索引。
例如,如果我们从 tensor<2x4x8x16xf32>
归约为
tensor<4x8xf32>
,那么从 2D 输出到 4D 输入的索引映射为
(d0, d1) -> (r0, d0, d1, r1)
,其中 d_i
是
对应于输出张量的索引。范围变量 r_j
编码
多个值,即要计算输出的 (d0, d1)
元素,我们需要
输入的 (r0, d0, d1, r1)
元素,其中 r0 in [0, 1]
和
r1 in [0, 15]
。
此映射可以根据 HLO 指令的属性构建,也可以将未融合指令的映射组合起来,以获取融合的编号。映射还具有一个域,用于指定映射适用于张量的哪些元素。
f(x) s.t.
lb <= g(x) <= ub
由于我们希望尽可能减少重新计算,因此需要一个用于 计算。XLA 已经依赖于 MLIR,因此我们使用 mlir::AffineMap 而不是编写另一个符号算术库。
典型的 AffineMap
如下所示:
(d0)[s0, s1] -> (s0 + 5, d0 * 2, s1 * 3 + 50)
AffineMap
有两种类型的参数:维度和符号。通过
dimensions 对应于维度变量 d,symbols 对应于
范围变量 r 和 RT 变量 rt。“AffineMap
”不包含
关于维度范围的元数据,因此我们必须提供此类数据
发布。
struct Interval {
int64_t lower;
int64_t upper;
};
// Dimension variable represents a dimension of a tensor or a GPU grid.
struct DimVar {
Interval bounds;
};
// RangeVar variable represents a range of values, e.g. to compute a single
// element of the reduction's result we need a range of values from the input
// tensor.
struct RangeVar {
Interval range;
};
// RTVar represents a runtime value, e.g. a dynamic offset in
// HLO dynamic-update-slice op.
struct RTVar {
Interval feasible_values;
const HloInstruction* hlo;
// This is a map from the iteration space of the corresponding indexing map to
// the iteration space of `hlo`. It shows what element of `hlo` we need to
// extract to get the runtime value for the RTVar.
mlir::AffineMap map;
};
class IndexingMap {
mlir::AffineMap affine_map_;
std::vector<DimVar> dim_vars_;
std::vector<RangeVar> range_vars_;
std::vector<RTVar> rt_vars_;
llvm::DenseMap<mlir::AffineExpr, Interval> constraints_;
};
dim_vars_
对编号映射的维度变量 d 编码包含边界盒约束条件,这些约束条件通常与转置、求和、逐元素求积等运算的输出张量的形状一致,但也有一些例外情况,例如 HloConcatenateInstruction。
range_vars_
会对 r 参数可以采用的可能值进行编码。
rt_vars_
会在运行时存储关联的 hlo 指令及其访问模式和可行值。例如,偏移值是动态的
对于 1D HloDynamicSliceInstruction
。相应的 RTVar
将有一个 HloInstruction*
,用于生成具有 (d0) -> ()
访问模式的秩为 0 的张量,因为对于输出的每个元素,我们都会从偏移量张量中提取相同的元素来计算输入的索引。我们还可以假设 slice 的偏移量始终介于 0
和 tensor_size - slice_size - 1
之间。
我们通过示例来了解上述所有内容的实际含义。
为未融合操作编制索引
Elementwise
对于元素级操作,索引映射是一种标识。
p0 = f32[10, 20] parameter(0)
p1 = f32[10, 20] parameter(1)
add = f32[10, 20] add(p0, p1)
输入映射的输出如下:
- 输出 ->input_i:
(d0, d1) -> (d0, d1)
domain:
d0 in [0, 9]
d1 in [0, 19]
输入到输出映射
- input_i -> output:
(d0, d1) -> (d0, d1)
domain:
d0 in [0, 9]
d1 in [0, 19]
广播
广播是指在绘制地图时, 在将输入映射到输出时添加该输出。
p0 = f32[20] parameter(0)
bc0 = f32[10, 20, 30] broadcast(p0), dimensions={1}
输入映射的输出如下:
(d0, d1, d2) -> (d1)
domain:
d0 in [0, 9]
d1 in [0, 19]
d2 in [0, 29]
输入到输出映射
(d0)[s0, s1] -> (s0, d0, s1)
domain:
d0 in [0, 19]
s0 in [0, 9]
s1 in [0, 29]
请注意,现在输入到输出的右侧有 s
映射。这些符号表示值范围。例如,在本例中,索引为 d0
的输入的每个元素都会映射到输出的 10x1x30 切片。
Constant 和 Iota
幸运的是,它们没有任何输入参数,因此无需计算编入索引的内容。
DynamicSlice
DynamicSlice 与 Slice 类似,但偏移是动态的。
src = s32[2,2,258] parameter(0)
of1 = s32[] parameter(1)
of2 = s32[] parameter(2)
of3 = s32[] parameter(3)
ds = dynamic-slice(s32[2,2,258] src, s32[] of1, s32[] of2, s32[] of3), dynamic_slice_sizes={1, 2, 32}
src
的输出到输入映射:
(d0, d1, d2)[s0, s1, s2] -> (d0 + s0, d1 + s1, d2 + s2)
domain:
d0 in [0, 0]
d1 in [0, 1]
d2 in [0, 31]
s0 in [0, 1]
hlo: of1 = s32[] parameter(1)
(d0, d1, d2) -> ()
s1 in [0, 0]
hlo: of2 = s32[] parameter(2)
(d0, d1, d2) -> ()
s2 in [0, 226]
hlo: of3 = s32[] parameter(3)
(d0, d1, d2) -> ()
请注意,现在右侧显示了输入到输出的映射。这些是表示运行时值的符号。例如,在本例中,对于输出中索引为 d0, d1, d2
的每个元素,我们都会访问 slice 偏移量 of1
、of2
和 of3
来计算输入的索引。运行时变量的间隔是假设整个 slice 都保持在边界内而得出的。
of1
、of2
和 of3
的输入映射的输出:
(d0, d1, d2) -> ()
domain:
d0 in [0, 0]
d1 in [0, 1]
d2 in [0, 31]
DynamicUpdateSlice
src = s32[20,30] parameter(0)
upd = s32[5,10] parameter(1)
of1 = s32[] parameter(2)
of2 = s32[] parameter(3)
dus = s32[20,30] dynamic-update-slice(
s32[20,30] src, s32[5,10] upd, s32[] of1, s32[] of2)
src
的输入映射的输出非常简单。通过调整权重,
将域限制为未更新的索引,但目前将映射编入索引
不支持不等式约束。
(d0, d1) -> (d0, d1)
domain:
d0 in [0, 19]
d1 in [0, 29]
upd
的输入映射的输出:
(d0, d1)[s0, s1] -> (d0 - s0, d1 - s1)
domain:
d0 in [0, 19]
d1 in [0, 29]
s0 in [0, 15]
hlo: of1 = s32[] parameter(2)
(d0, d1) -> ()
s1 in [0, 20]
hlo: of2 = s32[] parameter(3)
(d0, d1) -> ()
请注意,现在右侧显示了输入到输出的映射。这些是表示运行时值的符号。例如,在本例中,对于索引为 d0, d1
的输出的每个元素,我们都会访问 slice 偏移量 of1
和 of2
来计算输入的索引。间隔时间
通过假设整个 Slice 保留在
范围。
of1
和 of2
的输出到输入映射:
(d0, d1) -> ()
domain:
d0 in [0, 19]
d1 in [0, 29]
收集
仅支持简化的收集。请参阅 [gather_simplifier]。https://github.com/openxla/xla/blob/main/xla/hlo/transforms/simplifiers/gather_simplifier.h。
operand = f32[33,76,70] parameter(0)
indices = s32[1806,2] parameter(1)
gather = f32[1806,7,8,4] gather(operand, indices),
offset_dims={1,2,3},
collapsed_slice_dims={},
start_index_map={0,1},
index_vector_dim=1,
slice_sizes={7,8,4}
operand
的输出到输入映射:
(d0, d1, d2, d3)[s0, s1] -> (d1 + s0, d2 + s1, d3)
domain:
d0 in [0, 1805]
d1 in [0, 6]
d2 in [0, 7]
d3 in [0, 3]
s0 in [0, 26]
hlo: indices = s32[1806,2]{1,0} parameter(1)
(d0, d1, d2, d3) -> (d0, 0)
s1 in [0, 68]
hlo: indices = s32[1806,2]{1,0} parameter(1)
(d0, d1, d2, d3) -> (d0, 1)
请注意,现在右侧显示了输入到输出的映射。这些是表示运行时值的符号。例如,在本例中,对于索引为 d0, d1, d2, d3
的输出的每个元素,我们都会从 indices
张量中提取元素 (d0, 0) 和 (d0, 1)。
indices
的输出到输入映射:
(d0, d1, d2, d3)[s0] -> (d0, s0)
domain:
d0 in [0, 1805]
d1 in [0, 6]
d2 in [0, 7]
d3 in [0, 3]
s0 in [0, 1]
范围变量 s0
表明,我们需要 indices
张量的整个行 (d0, *) 才能计算输出的元素。
转置
转置的编号映射是输入/输出维度的排列。
p0 = f32[3, 12288, 6, 128] parameter(0)
transpose = f32[3, 6, 128, 12288] transpose(p0), dimensions={0, 2, 3, 1}
输出到输入映射:
(d0, d1, d2, d3) -> (d0, d3, d1, d2)
domain:
d0 in [0, 2]
d1 in [0, 5]
d2 in [0, 127]
d3 in [0, 12287]
输入到输出映射:
(d0, d1, d2, d3) -> (d0, d2, d3, d1)
domain:
d0 in [0, 2]
d1 in [0, 12287]
d2 in [0, 5]
d3 in [0, 127]
反向
反向编制索引映射会将还原的维度更改为 upper_bound(d_i) -
d_i
:
p0 = f32[1, 17, 9, 9] parameter(0)
reverse = f32[1, 17, 9, 9] reverse(p0), dimensions={1, 2}
输出到输入映射:
(d0, d1, d2, d3) -> (d0, -d1 + 16, -d2 + 8, d3)
domain:
d0 in [0, 0]
d1 in [0, 16]
d2 in [0, 8]
d3 in [0, 8]
输入到输出映射:
(d0, d1, d2, d3) -> (d0, -d1 + 16, -d2 + 8, d3)
domain:
d0 in [0, 0]
d1 in [0, 16]
d2 in [0, 8]
d3 in [0, 8]
(可变参数)Reduce
可变参数化缩减有多个输入和多个初始化,从输出到输入的映射会添加缩减后的维度。所以,它的行为方式与广播相反, 在某种意义上是个不错的选择。
p0 = f32[256,10] parameter(0)
p0_init = f32[] constant(-inf)
p1 = s32[256,10] parameter(1)
p1_init = s32[] constant(0)
reduce = (f32[10], s32[10]) reduce(p0, p1, p0_init, p1_init),
dimensions={0}, to_apply=max
输入映射的输出如下:
- output -> input_j:
(d0)[s0] -> (s0, d0)
domain:
d0 in [0, 9]
s0 in [0, 255]
- output -> init_j:
(d0) -> ()
domain:
d0 in [0, 9]
输出的输入映射:
- input_i -> output_j:
(d0, d1) -> (d1)
domain:
d0 in [0, 255]
d1 in [0, 9]
- init_i -> output_j:
()[s0] -> (s0)
domain:
s0 in [0, 9]
其中 i, j = 0, ... INPUT_COUNT。
Slice
从切片的输出到输入进行索引会形成一个阶梯式索引映射, 对输出的每个元素都有效。从输入映射到输出的过程如下: 限制在输入中元素的步长范围。
p0 = f32[10, 20, 50] parameter(0)
slice = f32[5, 3, 25] slice(f32[10, 20, 50] p0),
slice={[5:10:1], [3:20:7], [0:50:2]}
输出到输入映射:
(d0, d1, d2) -> (d0 + 5, d1 * 7 + 3, d2 * 2)
domain:
d0 in [0, 4]
d1 in [0, 2]
d2 in [0, 24]
输出映射的输入:
(d0, d1, d2) -> (d0 - 5, (d1 - 3) floordiv 7, d2 floordiv 2)
domain:
d0 in [5, 9]
d1 in [3, 17]
d2 in [0, 48]
(d1 - 3) mod 7 in [0, 0]
d2 mod 2 in [0, 0]
调整形状
重塑有不同的形式。
收起形状
这是一个“线性化”从 N-D 变形为 1D。
p0 = f32[4,8] parameter(0)
reshape = f32[32] reshape(p0)
输入映射的输出如下:
(d0) -> (d0 floordiv 8, d0 mod 8)
domain:
d0 in [0, 31]
输入到输出映射:
(d0, d1) -> (d0 * 8 + d1)
domain:
d0 in [0, 3]
d1 in [0, 7]
展开形状
这是一个反向“收缩形状”运算,用于将一维输入重塑为 N 维输出。
p0 = f32[32] parameter(0)
reshape = f32[4, 8] reshape(p0)
输出到输入映射:
(d0, d1) -> (d0 * 8 + d1)
domain:
d0 in [0, 3]
d1 in [0, 7]
输出映射的输入:
(d0) -> (d0 floordiv 8, d0 mod 8)
domain:
d0 in [0, 31]
常规重塑
这些是重塑操作,不能表示为一次展开或 收起形状。它们只能表示为 2 个或更多 展开或收起形状。
示例 1:线性化-去线性化。
p0 = f32[4,8] parameter(0)
reshape = f32[2, 4, 4] reshape(p0)
这种重塑可以表示为将 tensor<4x8xf32>
的形状收缩为 tensor<32xf32>
,然后将形状扩展为 tensor<2x4x4xf32>
的组合。
输入映射的输出如下:
(d0, d1, d2) -> (d0 * 2 + d1 floordiv 2, d2 + (d1 mod 2) * 4)
domain:
d0 in [0, 1]
d1 in [0, 3]
d2 in [0, 3]
输出映射的输入:
(d0, d1) -> (d0 floordiv 2, d1 floordiv 4 + (d0 mod 2) * 2, d1 mod 4)
domain:
d0 in [0, 3]
d1 in [0, 7]
示例 2:展开和收起的子形状
p0 = f32[4, 8, 12] parameter(0)
reshape = f32[32, 3, 4] reshape(p0)
这种重塑可以表示为两个重塑的组合。第一个
将最外层的维度 tensor<4x8x12xf32>
收起为 tensor<32x12xf32>
第二个则将最内层的维度 tensor<32x12xf32>
扩展为
tensor<32x3x4xf32>
。
输出到输入映射:
(d0, d1, d2) -> (d0 floordiv 8, d0 mod 8, d1 * 4 + d2)
domain:
d0 in [0, 31]
d1 in [0, 2]
d2 in [0, 3]
输出映射的输入:
(d0, d1, d2) -> (d0 * 8 + d1, d2 floordiv 4, d2 mod 4)
domain:
d0 in [0, 3]
d1 in [0, 7]
d2 in [0, 11]
Bitcast
一个位图操作可以表示为 transpose-reshape-transpose 序列。 因此,其编入索引映射只是此序列的编入索引映射的组合。
串联
concat 的输出到输入映射适用于所有输入,但使用 非重叠域,即一次仅使用其中一个输入。
p0 = f32[2, 5, 7] parameter(0)
p1 = f32[2, 11, 7] parameter(1)
p2 = f32[2, 17, 7] parameter(2)
ROOT concat = f32[2, 33, 7] concatenate(f32[2, 5, 7] p0, f32[2, 11, 7] p1, f32[2, 17, 7] p2), dimensions={1}
输入到输出的映射如下所示:
- 输出 ->输入 1:
(d0, d1, d2) -> (d0, d1, d2)
domain:
d0 in [0, 1]
d1 in [0, 4]
d2 in [0, 6]
- 输出 -> 输入 2:
(d0, d1, d2) -> (d0, d1 - 5, d2)
domain:
d0 in [0, 1]
d1 in [5, 15]
d2 in [0, 6]
- 输出 ->输入 3:
(d0, d1, d2) -> (d0, d1 - 16, d2)
domain:
d0 in [0, 1]
d1 in [16, 32]
d2 in [0, 6]
输出的输入映射:
- 输入 1 ->output:
(d0, d1, d2) -> (d0, d1, d2)
domain:
d0 in [0, 1]
d1 in [0, 4]
d2 in [0, 6]
- 输入 2 ->output:
(d0, d1, d2) -> (d0, d1 + 5, d2)
domain:
d0 in [0, 1]
d1 in [0, 10]
d2 in [0, 6]
- 输入 3 -> 输出:
(d0, d1, d2) -> (d0, d1 + 16, d2)
domain:
d0 in [0, 1]
d1 in [0, 16]
d2 in [0, 6]
圆点
点的索引映射与缩减的映射非常相似。
p0 = f32[4, 128, 256] parameter(0)
p1 = f32[4, 256, 64] parameter(1)
dot = f32[4, 128, 64] dot(p0, p1),
lhs_batch_dims={0}, rhs_batch_dims={0},
lhs_contracting_dims={2}, rhs_contracting_dims={1}
输出到输入的映射如下:
- output -> input_1:
(d0, d1, d2)[s0] -> (d0, d1, s0)
domain:
d0 in [0, 3]
d1 in [0, 127]
d2 in [0, 63]
s0 in [0, 255]
- output -> input_2:
(d0, d1, d2)[s0] -> (d0, s0, d2)
domain:
d0 in [0, 3]
d1 in [0, 127]
d2 in [0, 63]
s0 in [0, 255]
输出的输入映射:
- 输入_1 ->output:
(d0, d1, d2)[s0] -> (d0, d1, s0)
domain:
d0 in [0, 3]
d1 in [0, 127]
d2 in [0, 255]
s0 in [0, 63]
- 输入_2 ->output:
(d0, d1, d2)[s0] -> (d0, s0, d1)
domain:
d0 in [0, 3]
d1 in [0, 255]
d2 in [0, 63]
s0 in [0, 127]
Pad
PadOp 的索引编制与 SliceOp 索引编制相反。
p0 = f32[4, 4] parameter(0)
p1 = f32[] parameter(1)
pad = f32[12, 16] pad(p0, p1), padding=1_4_1x4_8_0
填充配置 1_4_1x4_8_0
表示 lowPad_highPad_interiorPad_dim_0 x lowPad_highPad_interiorPad_dim_1
。
输出到输入映射:
- 输出 ->输入:
(d0, d1) -> ((d0 - 1) floordiv 2, d1 - 4)
domain:
d0 in [1, 7]
d1 in [4, 7]
(d0 - 1) mod 2 in [0, 0]
- 输出 ->init:
(d0, d1) -> ()
domain:
d0 in [0, 11]
d1 in [0, 15]
ReduceWindow
XLA 中的 ReduceWindow 也会执行填充。因此,编制索引映射可以计算为不进行任何填充的 ReduceWindow 编制索引和 PadOp 编制索引的组合。
c_inf = f32[] constant(-inf)
p0 = f32[1024, 514] parameter(0)
reduce-window = f32[1024, 3] reduce-window(p0, c_inf),
window={size=1x512 pad=0_0x0_0}, to_apply=max
输出到输入映射:
- 输出 ->输入:
(d0, d1)[s0] -> (d0, d1 + s0)
domain:
d0 in [0, 1023]
d1 in [0, 2]
s0 in [0, 511]
- output -> init:
(d0, d1) -> ()
domain:
d0 in [0, 1023]
d1 in [0, 2]
为 Fusion 建立地图索引
融合操作的索引映射是 集群。某些输入可能会以不同的访问模式多次读取。
一个输入,多个编入映射
下面是 p0 + transpose(p0)
的示例。
f {
p0 = f32[1000, 1000] parameter(0)
transpose_p0 = f32[1000, 1000]{0, 1} transpose(p0), dimensions={1, 0}
ROOT a0 = f32[1000, 1000] add(p0, transpose_p0)
}
p0
的输出到输入编入索引映射将是 (d0, d1) -> (d0, d1)
和 (d0, d1) -> (d1, d0)
。这意味着要计算一个元素,
我们可能需要将输入参数读取两次。
一个输入,去重后的索引映射
在某些情况下,索引映射实际上是相同的,尽管 并非显而易见
f {
p0 = f32[20, 10, 50] parameter(0)
lhs_transpose_1 = f32[10, 20, 50] transpose(p0), dimensions={1, 0, 2}
lhs_e = f32[10, 20, 50] exponential(lhs_transpose_1)
lhs_transpose_2 = f32[10, 50, 20] transpose(lhs_e), dimensions={0, 2, 1}
rhs_transpose_1 = f32[50, 10, 20] transpose(p0), dimensions={2, 1, 0}
rhs_log = f32[50, 10, 20] exponential(rhs_transpose_1)
rhs_transpose_2 = f32[10, 50, 20] transpose(rhs_log), dimensions={1, 0, 2}
ROOT add = f32[10, 50, 20] add(lhs_transpose_2, rhs_transpose_2)
}
在本例中,p0
的输出到输入编制索引映射只是 (d0, d1, d2) -> (d2, d0, d1)
。
Softmax
软最大值的 parameter 0
的输出到输入编号映射:
(d0, d1, d2)[s0] -> (d0, d1, s0)
domain:
d0 in [0, 1]
d1 in [0, 64]
d2 in [0, 124]
s0 in [0, 124]
和
(d0, d1, d2) -> (d0, d1, d2)
domain:
d0 in [0, 1]
d1 in [0, 64]
d2 in [0, 124]
其中 s0
是指输入的最内层。
为地图简化程序编制索引
上游 mlir::AffineMap
的默认简化器无法使任何
对维度/符号范围做出的假设。因此,它无法高效地使用 mod
和 div
简化表达式。
我们可以利用对仿射映射中子表达式的下限和上限的了解,进一步简化这些表达式。
简化程序可以重写以下表达式。
[0, 6] x [0, 14]
中 d 的“(d0, d1) -> (d0 + d1 floordiv 16, d1 mod 16)
”变为“(d0, d1) -> (d0, d1)
”di in [0, 9]
的(d0, d1, d2) -> ((100d0 + 10d1 + d2) floorDiv 100, ((100d0 + 10d1 + d2) mod 100) floordiv 10, d2 mod 10)
变为(d0, d1, d2) -> (d0, d1, d2)
。- “
d_i in [0, 9]
”的“(d0, d1, d2) -> ((16d0 + 4d1 + d2) floordiv 8, (16d0 + 4d1 + d2) mod 8)
”会变为“(d0, d1, d2) -> (2d0 + (4d1 + d2) floordiv 8,(4d1 + d2) mod 8)
”。 [0, 9] x [0, 10]
中 d 的(d0, d1) -> (-(-11d0 - d1 + 109) floordiv 11 + 9)
变为(d0, d1) -> (d0)
。
将地图编入索引简化器可以让我们了解 HLO 中的变形会相互抵消。
p0 = f32[10, 10, 10] parameter(0)
reshape1 = f32[50, 20] reshape(p0)
reshape2 = f32[10, 10, 10] reshape(reshape1)
在对地图进行整合并对其进行简化之后,我们将
(d0, d1, d2) -> (d0, d1, d2)
。
对映射进行索引编制简化后,限制条件也会随之简化。
- 类型为
lower_bound <= affine_expr (floordiv, +, -, *) constant <= upper_bound
的约束条件会重写为updated_lower_bound <= affine_expr <= updated_upped_bound
。 - 始终满足的约束条件,例如
d0 + s0 in [0, 20]
排除了d0 in [0, 5]
和s0 in [1, 3]
。 - 约束条件中的仿射表达式会优化为上文中的编入索引仿射映射。
如需查看更多示例,请参阅 indexing_map_test.cc。