CUDA 驱动程序 API :: CUDA 工具包文档

6.30. 张量映射对象管理

本节介绍低级别 CUDA 驱动程序应用程序编程接口的张量映射对象管理功能。张量核心 API 仅在计算能力为 9.0 或更高的设备上受支持。

函数

CUresult cuTensorMapEncodeIm2col ( CUtensorMap* tensorMap, CUtensorMapDataType tensorDataType, cuuint32_t tensorRank, void* globalAddress, const cuuint64_t* globalDim, const cuuint64_t* globalStrides, const int* pixelBoxLowerCorner, const int* pixelBoxUpperCorner, cuuint32_t channelsPerPixel, cuuint32_t pixelsPerColumn, const cuuint32_t* elementStrides, CUtensorMapInterleave interleave, CUtensorMapSwizzle swizzle, CUtensorMapL2promotion l2Promotion, CUtensorMapFloatOOBfill oobFill ): 创建表示 im2col 内存区域的张量映射描述符对象。
CUresult cuTensorMapEncodeIm2colWide ( CUtensorMap* tensorMap, CUtensorMapDataType tensorDataType, cuuint32_t tensorRank, void* globalAddress, const cuuint64_t* globalDim, const cuuint64_t* globalStrides, int pixelBoxLowerCornerWidth, int pixelBoxUpperCornerWidth, cuuint32_t channelsPerPixel, cuuint32_t pixelsPerColumn, const cuuint32_t* elementStrides, CUtensorMapInterleave interleave, CUtensorMapIm2ColWideMode mode, CUtensorMapSwizzle swizzle, CUtensorMapL2promotion l2Promotion, CUtensorMapFloatOOBfill oobFill ): 创建表示 im2col 内存区域的张量映射描述符对象，但元素仅沿 W 维度加载。
CUresult cuTensorMapEncodeTiled ( CUtensorMap* tensorMap, CUtensorMapDataType tensorDataType, cuuint32_t tensorRank, void* globalAddress, const cuuint64_t* globalDim, const cuuint64_t* globalStrides, const cuuint32_t* boxDim, const cuuint32_t* elementStrides, CUtensorMapInterleave interleave, CUtensorMapSwizzle swizzle, CUtensorMapL2promotion l2Promotion, CUtensorMapFloatOOBfill oobFill ): 创建表示平铺内存区域的张量映射描述符对象。
CUresult cuTensorMapReplaceAddress ( CUtensorMap* tensorMap, void* globalAddress ): 使用更新的全局地址修改现有张量映射描述符。

函数

CUresult cuTensorMapEncodeIm2col ( CUtensorMap* tensorMap, CUtensorMapDataType tensorDataType, cuuint32_t tensorRank, void* globalAddress, const cuuint64_t* globalDim, const cuuint64_t* globalStrides, const int* pixelBoxLowerCorner, const int* pixelBoxUpperCorner, cuuint32_t channelsPerPixel, cuuint32_t pixelsPerColumn, const cuuint32_t* elementStrides, CUtensorMapInterleave interleave, CUtensorMapSwizzle swizzle, CUtensorMapL2promotion l2Promotion, CUtensorMapFloatOOBfill oobFill )

创建表示 im2col 内存区域的张量映射描述符对象。

参数

tensorMap: - 要创建的张量映射对象
tensorDataType: - 张量数据类型
tensorRank: - 张量的维度；必须至少为 3
globalAddress: - 张量描述的内存区域的起始地址
globalDim: - 包含沿每个维度的张量大小（元素数量）的数组tensorRank维度
globalStrides: - 包含沿每个维度的步幅大小（以字节为单位）的数组tensorRank- 1 维度
pixelBoxLowerCorner: - 包含下限框角 DHW 维度的数组
pixelBoxUpperCorner: - 包含上限框角 DHW 维度的数组
channelsPerPixel: - 每个像素的通道数
pixelsPerColumn: - 每列的像素数
elementStrides: - 包含每个维度的遍历步幅的数组tensorRank维度
interleave: - 张量寻址的交错布局类型
swizzle: - 共享内存内部的 Bank 交换模式
l2Promotion: - L2 提升大小
oobFill: - 指示是否将使用零或特殊的 NaN 常量来填充越界元素

返回值

CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE

描述

创建由描述 im2col 内存布局的参数指定的张量内存访问 (TMA) 对象的描述符，并在以下位置返回它tensorMap.

张量映射对象仅在计算能力为 9.0 或更高的设备上受支持。此外，张量映射对象是不透明值，因此，应仅通过 CUDA API 和 PTX 访问。

传递的参数受以下要求约束

tensorMap地址必须与 64 字节对齐。

tensorDataType必须是 CUtensorMapDataType 中的枚举，其定义为

‎    typedef enum CUtensorMapDataType_enum {
              CU_TENSOR_MAP_DATA_TYPE_UINT8 = 0,       // 1 byte
              CU_TENSOR_MAP_DATA_TYPE_UINT16,          // 2 bytes
              CU_TENSOR_MAP_DATA_TYPE_UINT32,          // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_INT32,           // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_UINT64,          // 8 bytes
              CU_TENSOR_MAP_DATA_TYPE_INT64,           // 8 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT16,         // 2 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT32,         // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT64,         // 8 bytes
              CU_TENSOR_MAP_DATA_TYPE_BFLOAT16,        // 2 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT32_FTZ,     // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_TFLOAT32,        // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_TFLOAT32_FTZ     // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,    // 4 bits
              CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B,   // 4 bits
              CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B    // 6 bits
          } CUtensorMapDataType;

CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B 将“16 x U4”打包值复制到与 8 字节对齐的内存。打包值之间没有间隙。CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 将“16 x U4”打包值复制到与 16 字节对齐的内存。每 8 字节的打包值块之间有 8 字节的间隙。CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 将“16 x U6”打包值复制到与 16 字节对齐的内存。每 12 字节的打包值块之间有 4 字节的间隙。

tensorRank，指定张量维度的数量，必须为 3、4 或 5。

globalAddress，指定描述的内存区域的起始地址，必须与 16 字节对齐。还需要满足以下要求
- 当interleave为 CU_TENSOR_MAP_INTERLEAVE_32B 时，globalAddress必须与 32 字节对齐。
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，globalAddress必须与 32 字节对齐。

globalDim数组，指定每个维度的张量大小tensorRank维度，必须为非零且小于或等于 2^32。此外，对于打包数据类型，还需要满足以下要求
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，globalDim[0] 必须是 128 的倍数。
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B 时，globalDim[0] 必须是 2 的倍数。
- 打包数据类型的维度必须反映各个 U# 值的数量。

globalStrides数组，指定较低维度的张量步幅tensorRank- 1 维，以字节为单位，必须是 16 的倍数且小于 2^40。此外，还需要满足以下要求
- 当interleave为 CU_TENSOR_MAP_INTERLEAVE_32B 时，步幅必须是 32 的倍数。
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，步幅必须是 32 的倍数。每个后续指定的维度都包含前一个维度的步幅
```
‎    globalStrides[0] = globalDim[0] * elementSizeInBytes(tensorDataType) + padding[0];
          for (i = 1; i < tensorRank - 1; i++)
              globalStrides[i] = globalStrides[i – 1] * (globalDim[i] + padding[i]);
              assert(globalStrides[i] >= globalDim[i]);
```

pixelBoxLowerCorner数组指定边界框从顶部/左侧/前角开始的坐标偏移量 {D, H, W}。偏移量的数量及其精度取决于张量维度
- 当tensorRank为 3 时，支持范围 [-32768, 32767] 内的一个有符号偏移量。
- 当tensorRank为 4 时，支持范围 [-128, 127] 内的两个有符号偏移量。
- 当tensorRank为 5 时，支持范围 [-16, 15] 内的三个偏移量。

pixelBoxUpperCorner数组指定边界框从底部/右侧/后角开始的坐标偏移量 {D, H, W}。偏移量的数量及其精度取决于张量维度
- 当tensorRank为 3 时，支持范围 [-32768, 32767] 内的一个有符号偏移量。
- 当tensorRank为 4 时，支持范围 [-128, 127] 内的两个有符号偏移量。
- 当tensorRank为 5 时，支持范围 [-16, 15] 内的三个偏移量。由以下项指定的边界框pixelBoxLowerCorner和pixelBoxUpperCorner必须具有非零面积。

channelsPerPixel，指定必须沿 C 维度访问的元素数量，必须小于或等于 256。此外，当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，channelsPerPixel必须为 128。

pixelsPerColumn，指定必须沿 {N, D, H, W} 维度访问的元素数量，必须小于或等于 1024。

elementStrides数组，指定沿每个维度的迭代步长tensorRank维度，必须为非零且小于或等于 8。请注意，当interleave为 CU_TENSOR_MAP_INTERLEAVE_NONE 时，此数组的第一个元素将被忽略，因为 TMA 不支持维度零的步幅。当elementStrides数组的所有元素都为一时，boxDim指定要加载的元素数量。但是，如果elementStrides[i] 对于某些i不等于一时，则 TMA 沿第 i 个维度加载 ceil(boxDim[i] /elementStrides[i]) 个元素。要沿第 i 个维度加载 N 个元素，boxDim[i] 必须设置为 N *elementStrides[i]。

interleave指定类型为 CUtensorMapInterleave 的交错布局，其定义为
```
‎    typedef enum CUtensorMapInterleave_enum {
              CU_TENSOR_MAP_INTERLEAVE_NONE = 0,
              CU_TENSOR_MAP_INTERLEAVE_16B,
              CU_TENSOR_MAP_INTERLEAVE_32B
          } CUtensorMapInterleave;
```
TMA 支持交错布局，例如 NC/8HWC8（其中 C8 在内存中使用 16 字节，假设每个通道 2 字节）或 NC/16HWC16（其中 C16 使用 32 字节）。当interleave为 CU_TENSOR_MAP_INTERLEAVE_NONE 且swizzle不是 CU_TENSOR_MAP_SWIZZLE_NONE 时，边界框内部维度（计算为channelsPerPixel乘以从tensorDataType派生的元素大小（以字节为单位））必须小于或等于交换大小。
- CU_TENSOR_MAP_SWIZZLE_32B 要求边界框内部维度小于或等于 32。
- CU_TENSOR_MAP_SWIZZLE_64B 要求边界框内部维度小于或等于 64。
- CU_TENSOR_MAP_SWIZZLE_128B* 要求边界框内部维度小于或等于 128。此外，tensorDataTypeCU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 需要interleave为 CU_TENSOR_MAP_INTERLEAVE_NONE。

swizzle，指定共享内存 Bank 交换模式，必须是 CUtensorMapSwizzle 类型，其定义为
```
‎    typedef enum CUtensorMapSwizzle_enum {
              CU_TENSOR_MAP_SWIZZLE_NONE = 0,
              CU_TENSOR_MAP_SWIZZLE_32B,                   // Swizzle 16B chunks within 32B  span
              CU_TENSOR_MAP_SWIZZLE_64B,                   // Swizzle 16B chunks within 64B  span
              CU_TENSOR_MAP_SWIZZLE_128B,                  // Swizzle 16B chunks within 128B span
              CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B,         // Swizzle 32B chunks within 128B span
              CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B_FLIP_8B, // Swizzle 32B chunks within 128B span, additionally swap lower 8B with upper 8B within each 16B for every alternate row
              CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B          // Swizzle 64B chunks within 128B span
          } CUtensorMapSwizzle;
```
数据以特定顺序组织在全局内存中；但是，这可能与应用程序在共享内存中访问数据的顺序不匹配。这种数据组织上的差异可能会在访问共享内存时导致 Bank 冲突。为了避免此问题，可以将数据加载到共享内存中，并在共享内存 Bank 之间进行混洗。当interleave为 CU_TENSOR_MAP_INTERLEAVE_32B 时，swizzle必须为 CU_TENSOR_MAP_SWIZZLE_32B。其他交错模式可以具有任何交换模式。当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 时，仅支持以下交换模式
- CU_TENSOR_MAP_SWIZZLE_NONE（加载和存储）
- CU_TENSOR_MAP_SWIZZLE_128B（加载和存储）
- CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B（加载和存储）
- CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B（仅存储）当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，仅支持以下交换模式
- CU_TENSOR_MAP_SWIZZLE_NONE（仅加载）
- CU_TENSOR_MAP_SWIZZLE_128B（仅加载）
- CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B（仅加载）

l2Promotion指定 L2 获取大小，指示 L2 请求从 DRAM 填充的字节粒度。它必须是 CUtensorMapL2promotion 类型，其定义为

‎    typedef enum CUtensorMapL2promotion_enum {
              CU_TENSOR_MAP_L2_PROMOTION_NONE = 0,
              CU_TENSOR_MAP_L2_PROMOTION_L2_64B,
              CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
              CU_TENSOR_MAP_L2_PROMOTION_L2_256B
          } CUtensorMapL2promotion;

oobFill，指示是否应使用零或特殊的 NaN 常量来填充越界元素，必须是 CUtensorMapFloatOOBfill 类型，其定义为
```
‎    typedef enum CUtensorMapFloatOOBfill_enum {
              CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE = 0,
              CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA
          } CUtensorMapFloatOOBfill;
```
请注意，仅当tensorDataType表示浮点数据类型，并且当tensorDataType不是 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B、CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 和 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 时，才能使用 CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA。

另请参阅

cuTensorMapEncodeTiled, cuTensorMapEncodeIm2colWide, cuTensorMapReplaceAddress

CUresult cuTensorMapEncodeIm2colWide ( CUtensorMap* tensorMap, CUtensorMapDataType tensorDataType, cuuint32_t tensorRank, void* globalAddress, const cuuint64_t* globalDim, const cuuint64_t* globalStrides, int pixelBoxLowerCornerWidth, int pixelBoxUpperCornerWidth, cuuint32_t channelsPerPixel, cuuint32_t pixelsPerColumn, const cuuint32_t* elementStrides, CUtensorMapInterleave interleave, CUtensorMapIm2ColWideMode mode, CUtensorMapSwizzle swizzle, CUtensorMapL2promotion l2Promotion, CUtensorMapFloatOOBfill oobFill )

创建表示 im2col 内存区域的张量映射描述符对象，但元素仅沿 W 维度加载。

参数

tensorMap: - 要创建的张量映射对象
tensorDataType: - 张量数据类型
tensorRank: - 张量的维度；必须至少为 3
globalAddress: - 张量描述的内存区域的起始地址
globalDim: - 包含沿每个维度的张量大小（元素数量）的数组tensorRank维度
globalStrides: - 包含沿每个维度的步幅大小（以字节为单位）的数组tensorRank- 1 维度
pixelBoxLowerCornerWidth: - 左侧框角的宽度偏移量
pixelBoxUpperCornerWidth: - 右侧框角的宽度偏移量
channelsPerPixel: - 每个像素的通道数
pixelsPerColumn: - 每列的像素数
elementStrides: - 包含每个维度的遍历步幅的数组tensorRank维度
interleave: - 张量寻址的交错布局类型
mode: - W 或 W128 模式
swizzle: - 共享内存内部的 Bank 交换模式
l2Promotion: - L2 提升大小
oobFill: - 指示是否将使用零或特殊的 NaN 常量来填充越界元素

返回值

CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE

描述

创建由描述 im2col 内存布局的参数指定的张量内存访问 (TMA) 对象的描述符，其中行始终沿 W 维度加载，并在以下位置返回它tensorMap。这假设内存中的张量布局为 NDHWC、NHWC 或 NWC。

此 API 仅在计算能力为 10.0 或更高的设备上受支持。此外，张量映射对象是不透明值，因此，应仅通过 CUDA API 和 PTX 访问。

传递的参数受以下要求约束

tensorMap地址必须与 64 字节对齐。

tensorDataType必须是 CUtensorMapDataType 中的枚举，其定义为

‎    typedef enum CUtensorMapDataType_enum {
              CU_TENSOR_MAP_DATA_TYPE_UINT8 = 0,       // 1 byte
              CU_TENSOR_MAP_DATA_TYPE_UINT16,          // 2 bytes
              CU_TENSOR_MAP_DATA_TYPE_UINT32,          // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_INT32,           // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_UINT64,          // 8 bytes
              CU_TENSOR_MAP_DATA_TYPE_INT64,           // 8 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT16,         // 2 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT32,         // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT64,         // 8 bytes
              CU_TENSOR_MAP_DATA_TYPE_BFLOAT16,        // 2 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT32_FTZ,     // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_TFLOAT32,        // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_TFLOAT32_FTZ     // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,    // 4 bits
              CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B,   // 4 bits
              CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B    // 6 bits
          } CUtensorMapDataType;

tensorRank，指定张量维度的数量，必须为 3、4 或 5。

globalAddress，指定描述的内存区域的起始地址，必须与 16 字节对齐。还需要满足以下要求
- 当interleave为 CU_TENSOR_MAP_INTERLEAVE_32B 时，globalAddress必须与 32 字节对齐。
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，globalAddress必须与 32 字节对齐。

globalDim数组，指定每个维度的张量大小tensorRank维度，必须为非零且小于或等于 2^32。此外，对于打包数据类型，还需要满足以下要求
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，globalDim[0] 必须是 128 的倍数。
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B 时，globalDim[0] 必须是 2 的倍数。
- 打包数据类型的维度必须反映各个 U# 值的数量。

globalStrides数组，指定较低维度的张量步幅tensorRank- 1 维，以字节为单位，必须是 16 的倍数且小于 2^40。此外，还需要满足以下要求
- 当interleave为 CU_TENSOR_MAP_INTERLEAVE_32B 时，步幅必须是 32 的倍数。
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，步幅必须是 32 的倍数。每个后续指定的维度都包含前一个维度的步幅
```
‎    globalStrides[0] = globalDim[0] * elementSizeInBytes(tensorDataType) + padding[0];
          for (i = 1; i < tensorRank - 1; i++)
              globalStrides[i] = globalStrides[i – 1] * (globalDim[i] + padding[i]);
              assert(globalStrides[i] >= globalDim[i]);
```

pixelBoxLowerCornerWidth指定边界框从左角开始的坐标偏移量 W。偏移量必须在范围 [-32768, 32767] 内。

pixelBoxUpperCornerWidth指定边界框从右角开始的坐标偏移量 W。偏移量必须在范围 [-32768, 32767] 内。

由以下项指定的边界框pixelBoxLowerCornerWidth和pixelBoxUpperCornerWidth必须具有非零面积。请注意，沿 D 和 H 维度的框大小始终等于 1。

channelsPerPixel，指定必须沿 C 维度访问的元素数量，必须小于或等于 256。此外，当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，channelsPerPixel必须为 128。

pixelsPerColumn，指定必须沿 W 维度访问的元素数量，必须小于或等于 1024。当mode为 CU_TENSOR_MAP_IM2COL_WIDE_MODE_W128 时，将忽略此字段。

elementStrides数组，指定沿每个维度的迭代步长tensorRank维度，必须为非零且小于或等于 8。请注意，当interleave为 CU_TENSOR_MAP_INTERLEAVE_NONE 时，此数组的第一个元素将被忽略，因为 TMA 不支持维度零的步幅。当elementStrides数组的所有元素都为一时，boxDim指定要加载的元素数量。但是，如果elementStrides[i] 对于某些i不等于一时，则 TMA 沿第 i 个维度加载 ceil(boxDim[i] /elementStrides[i]) 个元素。要沿第 i 个维度加载 N 个元素，boxDim[i] 必须设置为 N *elementStrides[i]。

interleave指定类型为 CUtensorMapInterleave 的交错布局，其定义为
```
‎    typedef enum CUtensorMapInterleave_enum {
              CU_TENSOR_MAP_INTERLEAVE_NONE = 0,
              CU_TENSOR_MAP_INTERLEAVE_16B,
              CU_TENSOR_MAP_INTERLEAVE_32B
          } CUtensorMapInterleave;
```
TMA 支持交错布局，例如 NC/8HWC8（其中 C8 在内存中使用 16 字节，假设每个通道 2 字节）或 NC/16HWC16（其中 C16 使用 32 字节）。当interleave为 CU_TENSOR_MAP_INTERLEAVE_NONE 时，边界框内部维度（计算为channelsPerPixel乘以从tensorDataType派生的元素大小（以字节为单位））必须小于或等于交换大小。
- CU_TENSOR_MAP_SWIZZLE_64B 要求边界框内部维度小于或等于 64。
- CU_TENSOR_MAP_SWIZZLE_128B* 要求边界框内部维度小于或等于 128。此外，tensorDataTypeCU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 需要interleave为 CU_TENSOR_MAP_INTERLEAVE_NONE。

mode，描述沿 W 维度加载的元素的加载，必须是以下 CUtensorMapIm2ColWideMode 类型之一
```
‎          CU_TENSOR_MAP_IM2COL_WIDE_MODE_W,
                CU_TENSOR_MAP_IM2COL_WIDE_MODE_W128
```
CU_TENSOR_MAP_IM2COL_WIDE_MODE_W 允许通过pixelsPerColumn字段指定沿 W 维度加载的元素数量。

swizzle，指定共享内存 Bank 交换模式，必须是以下 CUtensorMapSwizzle 模式之一（不支持其他交换模式）
```
‎    typedef enum CUtensorMapSwizzle_enum {
              CU_TENSOR_MAP_SWIZZLE_64B,                   // Swizzle 16B chunks within 64B  span
              CU_TENSOR_MAP_SWIZZLE_128B,                  // Swizzle 16B chunks within 128B span
              CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B,         // Swizzle 32B chunks within 128B span
          } CUtensorMapSwizzle;
```
数据以特定顺序组织在全局内存中；但是，这可能与应用程序在共享内存中访问数据的顺序不匹配。这种数据组织上的差异可能会在访问共享内存时导致 Bank 冲突。为了避免此问题，可以将数据加载到共享内存中，并在共享内存 Bank 之间进行混洗。当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 时，仅支持以下交换模式
- CU_TENSOR_MAP_SWIZZLE_128B（加载和存储）
- CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B（加载和存储）当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，仅支持以下交换模式
- CU_TENSOR_MAP_SWIZZLE_128B（仅加载）
- CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B（仅加载）

l2Promotion指定 L2 获取大小，指示 L2 请求从 DRAM 填充的字节粒度。它必须是 CUtensorMapL2promotion 类型，其定义为

‎    typedef enum CUtensorMapL2promotion_enum {
              CU_TENSOR_MAP_L2_PROMOTION_NONE = 0,
              CU_TENSOR_MAP_L2_PROMOTION_L2_64B,
              CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
              CU_TENSOR_MAP_L2_PROMOTION_L2_256B
          } CUtensorMapL2promotion;

oobFill，指示是否应使用零或特殊的 NaN 常量来填充越界元素，必须是 CUtensorMapFloatOOBfill 类型，其定义为
```
‎    typedef enum CUtensorMapFloatOOBfill_enum {
              CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE = 0,
              CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA
          } CUtensorMapFloatOOBfill;
```
请注意，仅当tensorDataType表示浮点数据类型，并且当tensorDataType不是 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B、CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 和 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 时，才能使用 CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA。

另请参阅

cuTensorMapEncodeTiled, cuTensorMapEncodeIm2col, cuTensorMapReplaceAddress

CUresult cuTensorMapEncodeTiled ( CUtensorMap* tensorMap, CUtensorMapDataType tensorDataType, cuuint32_t tensorRank, void* globalAddress, const cuuint64_t* globalDim, const cuuint64_t* globalStrides, const cuuint32_t* boxDim, const cuuint32_t* elementStrides, CUtensorMapInterleave interleave, CUtensorMapSwizzle swizzle, CUtensorMapL2promotion l2Promotion, CUtensorMapFloatOOBfill oobFill )

创建表示平铺内存区域的张量映射描述符对象。

参数

tensorMap: - 要创建的张量映射对象
tensorDataType: - 张量数据类型
tensorRank: - 张量的维度
globalAddress: - 张量描述的内存区域的起始地址
globalDim: - 包含沿每个维度的张量大小（元素数量）的数组tensorRank维度
globalStrides: - 包含沿每个维度的步幅大小（以字节为单位）的数组tensorRank- 1 维度
boxDim: - 包含沿每个维度遍历框大小（元素数量）的数组。tensorRank指定沿每个张量维度要遍历的元素数量。
elementStrides: - 包含每个维度的遍历步幅的数组tensorRank维度
interleave: - 张量寻址的交错布局类型
swizzle: - 共享内存内部的 Bank 交换模式
l2Promotion: - L2 提升大小
oobFill: - 指示是否必须使用零或特殊的 NaN 常量来填充越界元素

返回值

CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE

描述

为张量内存访问 (TMA) 对象创建描述符，该对象由描述平铺区域的参数指定，并在以下位置返回tensorMap.

张量映射对象仅在计算能力为 9.0 或更高的设备上受支持。此外，张量映射对象是不透明值，因此，应仅通过 CUDA API 和 PTX 访问。

传递的参数受以下要求约束

tensorMap地址必须与 64 字节对齐。

tensorDataType必须是 CUtensorMapDataType 中的枚举，其定义为

‎    typedef enum CUtensorMapDataType_enum {
              CU_TENSOR_MAP_DATA_TYPE_UINT8 = 0,       // 1 byte
              CU_TENSOR_MAP_DATA_TYPE_UINT16,          // 2 bytes
              CU_TENSOR_MAP_DATA_TYPE_UINT32,          // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_INT32,           // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_UINT64,          // 8 bytes
              CU_TENSOR_MAP_DATA_TYPE_INT64,           // 8 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT16,         // 2 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT32,         // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT64,         // 8 bytes
              CU_TENSOR_MAP_DATA_TYPE_BFLOAT16,        // 2 bytes
              CU_TENSOR_MAP_DATA_TYPE_FLOAT32_FTZ,     // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_TFLOAT32,        // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_TFLOAT32_FTZ,    // 4 bytes
              CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,    // 4 bits
              CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B,   // 4 bits
              CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B    // 6 bits
          } CUtensorMapDataType;

tensorRank必须为非零，且小于或等于最大支持维度 5。如果interleave不是 CU_TENSOR_MAP_INTERLEAVE_NONE，则tensorRank还必须大于或等于 3。

globalAddress，指定描述的内存区域的起始地址，必须与 16 字节对齐。还需要满足以下要求
- 当interleave为 CU_TENSOR_MAP_INTERLEAVE_32B 时，globalAddress必须与 32 字节对齐。
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，globalAddress必须与 32 字节对齐。

globalDim数组，指定每个维度的张量大小tensorRank维度，必须为非零且小于或等于 2^32。此外，对于打包数据类型，还需要满足以下要求
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，globalDim[0] 必须是 128 的倍数。
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B 时，globalDim[0] 必须是 2 的倍数。
- 打包数据类型的维度必须反映各个 U# 值的数量。

globalStrides数组，指定较低维度的张量步幅tensorRank- 1 维，以字节为单位，必须是 16 的倍数且小于 2^40。此外，还需要满足以下要求
- 当interleave为 CU_TENSOR_MAP_INTERLEAVE_32B 时，步幅必须是 32 的倍数。
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，步幅必须是 32 的倍数。每个后续指定的维度都包含前一个维度的步幅
```
‎    globalStrides[0] = globalDim[0] * elementSizeInBytes(tensorDataType) + padding[0];
          for (i = 1; i < tensorRank - 1; i++)
              globalStrides[i] = globalStrides[i – 1] * (globalDim[i] + padding[i]);
              assert(globalStrides[i] >= globalDim[i]);
```

boxDim数组，该数组指定沿每个维度要遍历的元素数量，tensorRank必须为非零，且小于或等于 256。此外，还需要满足以下要求
- 当interleave为 CU_TENSOR_MAP_INTERLEAVE_NONE，{boxDim[0] * elementSizeInBytes(tensorDataType) } 必须是 16 字节的倍数。
- 当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 或 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B，boxDim[0] 必须为 128。

elementStrides数组，指定沿每个维度的迭代步长tensorRank维度，必须为非零且小于或等于 8。请注意，当interleave为 CU_TENSOR_MAP_INTERLEAVE_NONE，此数组的第一个元素将被忽略，因为 TMA 不支持维度零的步幅。当elementStrides数组的所有元素都为一时，boxDim指定要加载的元素数量。但是，如果elementStrides[i] 不等于 1，则 TMA 加载 ceil(boxDim[i] /elementStrides[i]) 个元素。要沿第 i 个维度加载 N 个元素，boxDim[i] 必须设置为 N *elementStrides[i]。

interleave指定类型为 CUtensorMapInterleave 的交错布局，其定义为
```
‎    typedef enum CUtensorMapInterleave_enum {
              CU_TENSOR_MAP_INTERLEAVE_NONE = 0,
              CU_TENSOR_MAP_INTERLEAVE_16B,
              CU_TENSOR_MAP_INTERLEAVE_32B
          } CUtensorMapInterleave;
```
TMA 支持交错布局，例如 NC/8HWC8（其中 C8 在内存中使用 16 字节，假设每个通道 2 字节）或 NC/16HWC16（其中 C16 使用 32 字节）。当interleave为 CU_TENSOR_MAP_INTERLEAVE_NONE 且swizzle不是 CU_TENSOR_MAP_SWIZZLE_NONE 时，边界框内部维度（计算为boxDim[0] 乘以从tensorDataType派生的元素大小（以字节为单位））必须小于或等于交换大小。
- CU_TENSOR_MAP_SWIZZLE_32B 要求边界框内部维度小于或等于 32。
- CU_TENSOR_MAP_SWIZZLE_64B 要求边界框内部维度小于或等于 64。
- CU_TENSOR_MAP_SWIZZLE_128B* 要求边界框内部维度小于或等于 128。此外，tensorDataTypeCU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 需要interleave为 CU_TENSOR_MAP_INTERLEAVE_NONE。

swizzle，指定共享内存 Bank 交换模式，必须是 CUtensorMapSwizzle 类型，其定义为
```
‎    typedef enum CUtensorMapSwizzle_enum {
              CU_TENSOR_MAP_SWIZZLE_NONE = 0,
              CU_TENSOR_MAP_SWIZZLE_32B,                   // Swizzle 16B chunks within 32B  span
              CU_TENSOR_MAP_SWIZZLE_64B,                   // Swizzle 16B chunks within 64B  span
              CU_TENSOR_MAP_SWIZZLE_128B,                  // Swizzle 16B chunks within 128B span
              CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B,         // Swizzle 32B chunks within 128B span
              CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B_FLIP_8B, // Swizzle 32B chunks within 128B span, additionally swap lower 8B with upper 8B within each 16B for every alternate row
              CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B          // Swizzle 64B chunks within 128B span
          } CUtensorMapSwizzle;
```
数据以特定顺序组织在全局内存中；但是，这可能与应用程序在共享内存中访问数据的顺序不匹配。这种数据组织上的差异可能会在访问共享内存时导致 Bank 冲突。为了避免此问题，可以将数据加载到共享内存中，并在共享内存 Bank 之间进行混洗。当interleave为 CU_TENSOR_MAP_INTERLEAVE_32B 时，swizzle必须为 CU_TENSOR_MAP_SWIZZLE_32B。其他交错模式可以具有任何交换模式。当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 时，仅支持以下交换模式
- CU_TENSOR_MAP_SWIZZLE_NONE（加载和存储）
- CU_TENSOR_MAP_SWIZZLE_128B（加载和存储）
- CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B（加载和存储）
- CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B（仅存储）当tensorDataType为 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 时，仅支持以下交换模式
- CU_TENSOR_MAP_SWIZZLE_NONE（仅加载）
- CU_TENSOR_MAP_SWIZZLE_128B（仅加载）
- CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B（仅加载）

l2Promotion导出的元素大小。指定 L2 获取大小，指示从 DRAM 填充 L2 请求的字节粒度。它必须是 CUtensorMapL2promotion 类型，定义为

‎    typedef enum CUtensorMapL2promotion_enum {
              CU_TENSOR_MAP_L2_PROMOTION_NONE = 0,
              CU_TENSOR_MAP_L2_PROMOTION_L2_64B,
              CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
              CU_TENSOR_MAP_L2_PROMOTION_L2_256B
          } CUtensorMapL2promotion;

oobFill，指示是否应使用零或特殊的 NaN 常量来填充越界元素，必须是 CUtensorMapFloatOOBfill 类型，其定义为
```
‎    typedef enum CUtensorMapFloatOOBfill_enum {
              CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE = 0,
              CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA
          } CUtensorMapFloatOOBfill;
```
请注意，仅当tensorDataType表示浮点数据类型，并且当tensorDataType不是 CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B、CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B 和 CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B 时，才能使用 CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA。

另请参阅

cuTensorMapEncodeIm2col, cuTensorMapEncodeIm2colWide, cuTensorMapReplaceAddress

CUresult cuTensorMapReplaceAddress ( CUtensorMap* tensorMap, void* globalAddress )

使用更新的全局地址修改现有张量映射描述符。

参数

tensorMap: - 要修改的张量映射对象
globalAddress: - 张量描述的内存区域的起始地址，必须遵循之前的对齐要求

返回值

CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE

描述

修改传入的张量内存访问 (TMA) 对象的描述符tensorMap并使用更新的globalAddress.

张量映射对象仅在计算能力为 9.0 或更高的设备上受支持。此外，张量映射对象是不透明值，因此，应仅通过 CUDA API 调用访问。

另请参阅

cuTensorMapEncodeTiled, cuTensorMapEncodeIm2col cuTensorMapEncodeIm2colWide