cuda.bindings._v2.driver.tensor_map_encode_im2col#

cuda.bindings._v2.driver.tensor_map_encode_im2col(
int tensor_data_type,
uint32_t tensor_rank,
intptr_t global_address,
intptr_t global_dim,
intptr_t global_strides,
intptr_t pixel_box_lower_corner,
intptr_t pixel_box_upper_corner,
uint32_t channels_per_pixel,
uint32_t pixels_per_column,
intptr_t element_strides,
int interleave,
int swizzle,
int l2promotion,
int oob_fill,
)#

Create a tensor map descriptor object representing im2col memory region.

Creates a descriptor for Tensor Memory Access (TMA) object specified by the parameters describing a im2col memory layout 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, which specifies the number of tensor dimensions, must be 3, 4, or 5.

  • 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**.

  • pixel_box_lower_corner array specifies the coordinate offsets {D, H, W} of the bounding box from top/left/front corner. The number of offsets and their precision depend on the tensor dimensionality:.

  • When tensor_rank is 3, one signed offset within range [-32768, 32767] is supported.

  • When tensor_rank is 4, two signed offsets each within range [-128, 127] are supported.

  • When tensor_rank is 5, three offsets each within range [-16, 15] are supported.

  • pixel_box_upper_corner array specifies the coordinate offsets {D, H, W} of the bounding box from bottom/right/back corner. The number of offsets and their precision depend on the tensor dimensionality:.

  • When tensor_rank is 3, one signed offset within range [-32768, 32767] is supported.

  • When tensor_rank is 4, two signed offsets each within range [-128, 127] are supported.

  • When tensor_rank is 5, three offsets each within range [-16, 15] are supported. The bounding box specified by pixel_box_lower_corner and pixel_box_upper_corner must have non-zero area.

  • channels_per_pixel, which specifies the number of elements which must be accessed along C dimension, must be less than or equal to 256. Additionally, when tensor_data_type is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, channels_per_pixel must be 128.

  • pixels_per_column, which specifies the number of elements that must be accessed along the {N, D, H, W} dimensions, must be less than or equal to 1024.

  • 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 the element_strides array are one, boxDim specifies the number of elements to load. However, if element_strides[i] is not equal to one for some i, then TMA loads ceil( boxDim[i] / element_strides[i]) number of elements along i-th dimension. To load N elements along i-th dimension, boxDim[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 channels_per_pixel multiplied by element size in bytes 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 granularity at which L2 requests are 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; must be at least 3.

  • 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.

  • pixel_box_lower_corner (intptr_t) – Array containing DHW dimensions of lower box corner.

  • pixel_box_upper_corner (intptr_t) – Array containing DHW dimensions of upper box corner.

  • channels_per_pixel (uint32_t) – Number of channels per pixel.

  • pixels_per_column (uint32_t) – Number of pixels per column.

  • 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 will be used to fill out-of-bound elements.

Returns:

Tensor map object to create.

Return type:

CUtensorMap

See also

cuTensorMapEncodeIm2col