cuda.bindings._v2.driver.tensor_map_encode_tiled#

cuda.bindings._v2.driver.tensor_map_encode_tiled(
int tensor_data_type,
uint32_t tensor_rank,
intptr_t global_address,
intptr_t global_dim,
intptr_t global_strides,
intptr_t box_dim,
intptr_t element_strides,
int interleave,
int swizzle,
int l2promotion,
int oob_fill,
)#

Create a tensor map descriptor object representing tiled memory region.

Creates a descriptor for Tensor Memory Access (TMA) object specified by the parameters describing a tiled region and returns it in tensor_map. Tensor map objects are only supported on devices of compute capability 9.0 or higher. Additionally, a tensor map object is an opaque value, and, as such, should only be accessed through CUDA APIs and PTX. The parameters passed are bound to the following requirements:.

  • tensor_map address must be aligned to 64 bytes.

  • tensor_data_type has to be an enum from CUtensor_mapDataType which is defined as:.

  • **View CUDA Toolkit Documentation for a C++ code example**.

  • CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B copies ‘16 x U4’ packed values to memory aligned as 8 bytes. There are no gaps between packed values. CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B copies ‘16 x U4’ packed values to memory aligned as 16 bytes. There are 8 byte gaps between every 8 byte chunk of packed values. CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B copies ‘16 x U6’ packed values to memory aligned as 16 bytes. There are 4 byte gaps between every 12 byte chunk of packed values.

  • tensor_rank must be non-zero and less than or equal to the maximum supported dimensionality of 5. If interleave is not CU_TENSOR_MAP_INTERLEAVE_NONE, then tensor_rank must additionally be greater than or equal to 3.

  • global_address, which specifies the starting address of the memory region described, must be 16 byte aligned. The following requirements need to also be met:.

  • When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, global_address must be 32 byte aligned.

  • When tensor_data_type is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, global_address must be 32 byte aligned.

global_dim array, which specifies tensor size of each of the tensor_rank dimensions, must be non-zero and less than or equal to 2^32. Additionally, the following requirements need to be met for the packed data types:.

  • When tensor_data_type is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, global_dim[0] must be a multiple of 128.

  • When tensor_data_type is CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B, global_dim[0] must be a multiple of 2.

  • Dimension for the packed data types must reflect the number of individual U# values.

global_strides array, which specifies tensor stride of each of the lower tensor_rank - 1 dimensions in bytes, must be a multiple of 16 and less than 2^40. Additionally, the following requirements need to be met:.

  • When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, the strides must be a multiple of 32.

  • When tensor_data_type is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, the strides must be a multiple of 32. Each following dimension specified includes previous dimension stride:.

  • **View CUDA Toolkit Documentation for a C++ code example**.

box_dim array, which specifies number of elements to be traversed along each of the tensor_rank dimensions, must be non-zero and less than or equal to 256. Additionally, the following requirements need to be met:.

  • When interleave is CU_TENSOR_MAP_INTERLEAVE_NONE, { box_dim[0] * elementSizeInBytes( tensor_data_type ) } must be a multiple of 16 bytes.

  • When tensor_data_type is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, box_dim[0] must be 128.

element_strides array, which specifies the iteration step along each of the tensor_rank dimensions, must be non-zero and less than or equal to 8. Note that when interleave is CU_TENSOR_MAP_INTERLEAVE_NONE, the first element of this array is ignored since TMA doesn’t support the stride for dimension zero. When all elements of element_strides array is one, box_dim specifies the number of elements to load. However, if the element_strides[i] is not equal to one, then TMA loads ceil( box_dim[i] / element_strides[i]) number of elements along i-th dimension. To load N elements along i-th dimension, box_dim[i] must be set to N * element_strides[i].

  • interleave specifies the interleaved layout of type CUtensor_mapInterleave, which is defined as:.

  • **View CUDA Toolkit Documentation for a C++ code example**.

  • TMA supports interleaved layouts like NC/8HWC8 where C8 utilizes 16 bytes in memory assuming 2 byte per channel or NC/16HWC16 where C16 uses 32 bytes. When interleave is CU_TENSOR_MAP_INTERLEAVE_NONE and swizzle is not CU_TENSOR_MAP_SWIZZLE_NONE, the bounding box inner dimension (computed as box_dim[0] multiplied by element size derived from tensor_data_type) must be less than or equal to the swizzle size.

  • CU_TENSOR_MAP_SWIZZLE_32B requires the bounding box inner dimension to be <= 32.

  • CU_TENSOR_MAP_SWIZZLE_64B requires the bounding box inner dimension to be <= 64.

  • CU_TENSOR_MAP_SWIZZLE_128B* require the bounding box inner dimension to be <= 128. Additionally, tensor_data_type of CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B requires interleave to be CU_TENSOR_MAP_INTERLEAVE_NONE.

  • swizzle, which specifies the shared memory bank swizzling pattern, has to be of type CUtensor_mapSwizzle which is defined as:.

  • **View CUDA Toolkit Documentation for a C++ code example**.

  • Data are organized in a specific order in global memory; however, this may not match the order in which the application accesses data in shared memory. This difference in data organization may cause bank conflicts when shared memory is accessed. In order to avoid this problem, data can be loaded to shared memory with shuffling across shared memory banks. When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, swizzle must be CU_TENSOR_MAP_SWIZZLE_32B. Other interleave modes can have any swizzling pattern. When the tensor_data_type is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B, only the following swizzle modes are supported:.

  • CU_TENSOR_MAP_SWIZZLE_NONE (Load & Store).

  • CU_TENSOR_MAP_SWIZZLE_128B (Load & Store).

  • CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B (Load & Store).

  • CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B (Store only) When the tensor_data_type is CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, only the following swizzle modes are supported:.

  • CU_TENSOR_MAP_SWIZZLE_NONE (Load only).

  • CU_TENSOR_MAP_SWIZZLE_128B (Load only).

  • CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B (Load only).

  • l2promotion specifies L2 fetch size which indicates the byte granurality at which L2 requests is filled from DRAM. It must be of type CUtensor_mapL2promotion, which is defined as:.

  • **View CUDA Toolkit Documentation for a C++ code example**.

  • oob_fill, which indicates whether zero or a special NaN constant should be used to fill out-of-bound elements, must be of type CUtensor_mapFloatOOBfill which is defined as:.

  • **View CUDA Toolkit Documentation for a C++ code example**.

  • Note that CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA can only be used when tensor_data_type represents a floating-point data type, and when tensor_data_type is not CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B, CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, and CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B.

Parameters:

  • tensor_data_type (TensorMapDataType) – Tensor data type.

  • tensor_rank (uint32_t) – Dimensionality of tensor.

  • global_address (intptr_t) – Starting address of memory region described by tensor.

  • global_dim (intptr_t) – Array containing tensor size (number of elements) along each of the tensor_rank dimensions.

  • global_strides (intptr_t) – Array containing stride size (in bytes) along each of the tensor_rank - 1 dimensions.

  • box_dim (intptr_t) – Array containing traversal box size (number of elments) along each of the tensor_rank dimensions. Specifies how many elements to be traversed along each tensor dimension.

  • element_strides (intptr_t) – Array containing traversal stride in each of the tensor_rank dimensions.

  • interleave (TensorMapInterleave) – Type of interleaved layout the tensor addresses.

  • swizzle (TensorMapSwizzle) – Bank swizzling pattern inside shared memory.

  • l2promotion (TensorMapL2Promotion) – L2 promotion size.

  • oob_fill (TensorMapFloatOOBFill) – Indicate whether zero or special NaN constant must be used to fill out-of-bound elements.

Returns:

Tensor map object to create.

Return type:

CUtensorMap

See also

cuTensorMapEncodeTiled