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 可以指定来控制越界访问的行为。

参数:

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

• mask (triton.int1 的块，可选) – 如果 mask[idx] 为 false，则不要加载地址 pointer[idx] 的数据（对于块指针必须是 None）

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

• boundary_check (整数元组, 可选) – 整数元组，指示应进行边界检查的维度

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

• cache_modifier (str, 可选, 应为 {“”, “.ca”, “.cg”, “.cv”} 之一，其中 “.ca” 表示在所有级别缓存，“.cg” 表示在全局级别缓存（在 L2 及以下缓存，不在 L1），“.cv” 表示不缓存并重新获取。参见 cache operator 以获取更多详情。) – 改变 NVIDIA PTX 中的缓存选项

• eviction_policy (str, 可选) – 改变 NVIDIA PTX 中的逐出策略

• volatile (bool, 可选) – 改变 NVIDIA PTX 中的 volatile 选项