triton.language.load

triton.language.load(pointer, mask=None, other=None, boundary_check=(), padding_option='', cache_modifier='', eviction_policy='', volatile=False)

返回一个数据张量，其值从由pointer定义的内存位置加载：

1. 如果pointer是单元素指针，则会加载一个标量。在这种情况下：

   • mask 和 other 也必须是标量，

   • other 被隐式类型转换为 pointer.dtype.element_ty，并且

   • boundary_check 和 padding_option 必须为空。

2. 如果pointer是一个N维指针张量，则会加载一个N维张量。在这种情况下：

   • mask 和 other 会被隐式广播到 pointer.shape 的形状

   • other 被隐式类型转换为 pointer.dtype.element_ty，并且

   • boundary_check 和 padding_option 必须为空。

3. 如果pointer是由make_block_ptr定义的块指针，则会加载一个张量。在这种情况下：

   • mask 和 other 必须为 None，并且

   • 可以指定boundary_check和padding_option来控制越界访问的行为。

Parameters:

• pointer (triton.PointerType, 或 dtype=triton.PointerType 的数据块) – 指向要加载数据的指针

• mask (Block of triton.int1, 可选参数) – 如果mask[idx]为false，则不加载地址pointer[idx]处的数据 (使用块指针时必须为None)

• other (Block, optional) – 如果 mask[idx] 为 false，则返回 other[idx]

• boundary_check (tuple of ints, optional) – 整数元组，表示应进行边界检查的维度

• padding_option – 应为 {"", "zero", "nan"} 中的一个，表示越界时使用的填充值。"" 表示未定义的值。

• cache_modifier (str, 可选参数，可选值为 {"", ".ca", ".cg", ".cv"}，其中 ".ca" 表示在所有级别缓存，".cg" 表示在全局级别缓存（在L2及以下缓存，不包括L1），".cv" 表示不缓存并重新获取。更多详情请参阅cache operator) – 修改NVIDIA PTX中的缓存选项

• eviction_policy (str, optional) – 更改NVIDIA PTX中的驱逐策略

• volatile (bool, optional) – 修改NVIDIA PTX中的volatile选项