cuBLASMp Data Types

Data Types

cublasMpHandle_t

The cublasMpHandle_t structure holds the cuBLASMp library context (CUDA device state, auxiliary streams/events, and library handles). The handle is tied to the CUDA device and stream active at cublasMpCreate() time and must be destroyed with cublasMpDestroy().

cublasMpGrid_t

The cublasMpGrid_t structure stores the grid dimensions, layout, and the NCCL communicator associated with the participating ranks. The communicator must contain exactly nprow * npcol ranks, and all collective grid-management calls must be entered by that same rank set. It must be initialized and destroyed using cublasMpGridCreate() and cublasMpGridDestroy() respectively.

cublasMpMatrixDescriptor_t

The cublasMpMatrixDescriptor_t structure captures the shape and block-cyclic distribution of a distributed matrix. Most compute routines use descriptors over a non-null cublasMpGrid_t, while redistribution routines also allow grid == NULL on ranks that do not participate in that matrix’s grid. The block sizes (mb, nb) and local leading dimension (lld) are specified in elements; mb, nb, and lld must be positive, and rsrc/csrc are zero-based source coordinates whenever a grid is present. It must be initialized and destroyed using cublasMpMatrixDescriptorCreate() and cublasMpMatrixDescriptorDestroy() respectively.

cublasMpMatmulDescriptor_t

The cublasMpMatmulDescriptor_t structure captures the properties of a distributed matrix-matrix multiplication performed using cublasMpMatmul(). It must be initialized and destroyed using cublasMpMatmulDescriptorCreate() and cublasMpMatmulDescriptorDestroy() functions respectively.

Enumerators

cublasMpStatus_t

The type is used for function status returns. All cuBLASMp library functions return their status, which can have the following values.

Value	Meaning
CUBLASMP_STATUS_SUCCESS	The operation completed successfully.
CUBLASMP_STATUS_NOT_INITIALIZED	The cuBLASMp library was not initialized.
CUBLASMP_STATUS_ALLOCATION_FAILED	Resource allocation failed inside the cuBLASMp library.
CUBLASMP_STATUS_INVALID_VALUE	An unsupported value or parameter was passed to the function.
CUBLASMP_STATUS_ARCHITECTURE_MISMATCH	The function requires a feature absent from the device architecture.
CUBLASMP_STATUS_EXECUTION_FAILED	The GPU program failed to execute.
CUBLASMP_STATUS_INTERNAL_ERROR	An internal cuBLASMp operation failed.
CUBLASMP_STATUS_NOT_SUPPORTED	The functionality requested is not supported.

cublasMpEmulationStrategy_t

cublasMpEmulationStrategy_t enumerated type is used in cublasMpSetEmulationStrategy() to choose how to leverage floating point emulation algorithms.

Value	Meaning
CUBLASMP_EMULATION_STRATEGY_DEFAULT	This is the default emulation strategy and is equivalent to CUBLASMP_EMULATION_STRATEGY_PERFORMANT.
CUBLASMP_EMULATION_STRATEGY_PERFORMANT	A strategy which utilizes emulation whenever it provides a performance benefit.
CUBLASMP_EMULATION_STRATEGY_EAGER	A strategy which utilizes emulation whenever possible.

Note

The emulation strategy must be set after handle creation using cublasMpCreate() and before performing computations.

cublasMpGridLayout_t

Describes the ordering of the grid of processes.

Value	Meaning
CUBLASMP_GRID_LAYOUT_COL_MAJOR	The grid of processes will be accessed in column-major ordering.
CUBLASMP_GRID_LAYOUT_ROW_MAJOR	The grid of processes will be accessed in row-major ordering.

cublasMpLoggerCallback_t

Function pointer type for cuBLASMp logging callbacks. The callback function is called by the library to report logging information. The callback can be set using cublasMpLoggerSetCallback().

The callback function takes the following parameters:

Parameter	Type	Description
logLevel	int	The severity level of the log message
functionName	const char*	Name of the function generating the log message
message	const char*	The actual log message text

cublasMpMatmulDescriptorAttribute_t

Attributes of cublasMpMatmulDescriptor_t that can be set using cublasMpMatmulDescriptorSetAttribute() and queried using cublasMpMatmulDescriptorGetAttribute().

CUDA-version-specific attribute behavior is selected from the cuBLAS library loaded at runtime, which may have a different minor version than the toolkit used to build cuBLASMp.

Value (CUBLASMP_MATMUL_ DESCRIPTOR_ATTRIBUTE_*)	Meaning	Type	Value (complete name)
TRANSA	Indicates which operation needs to be performed with the dense matrix A.	cublasOperation_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_TRANSA
TRANSB	Indicates which operation needs to be performed with the dense matrix B.	cublasOperation_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_TRANSB
COMPUTE_TYPE	Indicates the compute type of the matrix multiplication.	cublasComputeType_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_COMPUTE_TYPE
ALGO_TYPE	Selects the algorithm type. DEFAULT lets cuBLASMp choose an implementation and its fallbacks. An explicitly selected algorithm returns CUBLASMP_STATUS_NOT_SUPPORTED when its requirements are not met.	cublasMpMatmulAlgoType_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_ALGO_TYPE
COMMUNICATION_SM_COUNT	Indicates the number of SMs to be used for communication.	int	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_COMMUNICATION_SM_COUNT
EPILOGUE	Specifies the epilogue operation to be performed after matrix multiplication (e.g., activation functions, bias addition).	cublasMpMatmulEpilogue_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE
BIAS_POINTER	Bias or bias gradient vector pointer in the device memory. Input vector with length that matches the number of rows of matrix D when one of the following epilogues is used: CUBLASMP_MATMUL_EPILOGUE_BIAS, CUBLASMP_MATMUL_EPILOGUE_RELU_BIAS, CUBLASMP_MATMUL_EPILOGUE_RELU_AUX_BIAS, CUBLASMP_MATMUL_EPILOGUE_GELU_BIAS, CUBLASMP_MATMUL_EPILOGUE_GELU_AUX_BIAS. Output vector with length that matches the number of rows of matrix D when one of the following epilogues is used: CUBLASMP_MATMUL_EPILOGUE_DRELU_BGRAD, CUBLASMP_MATMUL_EPILOGUE_DGELU_BGRAD, CUBLASMP_MATMUL_EPILOGUE_BGRADA. Output vector with length that matches the number of columns of matrix D when one of the following epilogues is used: CUBLASMP_MATMUL_EPILOGUE_BGRADB. Bias vector elements are the same type as the datatype of matrix D.	void*	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_BIAS_POINTER
BIAS_BATCH_STRIDE	Stride in elements between bias values in batched operations. Currently, batched matmul is not supported and this parameter has to be to 0 (default value).	int64_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_BIAS_BATCH_STRIDE
BIAS_DATA_TYPE	Data type of the bias vector (e.g., CUDA_R_16F, CUDA_R_32F).	cudaDataType_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_BIAS_DATA_TYPE
EPILOGUE_AUX_POINTER	Pointer for epilogue auxiliary buffer. If CUBLASMP_MATMUL_EPILOGUE_RELU_AUX or CUBLASMP_MATMUL_EPILOGUE_RELU_AUX_BIAS epilogue is used - output vector for ReLu bit-mask in forward pass; if CUBLASMP_MATMUL_EPILOGUE_DRELU or CUBLASMP_MATMUL_EPILOGUE_DRELU_BGRAD epilogue is used - input vector for ReLu bit-mask in backward pass; if CUBLASMP_MATMUL_EPILOGUE_GELU_AUX_BIAS epilogue is used - output of GELU input matrix in forward pass; if CUBLASMP_MATMUL_EPILOGUE_DGELU or CUBLASMP_MATMUL_EPILOGUE_DGELU_BGRAD epilogue is used - input of GELU input matrix for backward pass. For the data type, see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_DATA_TYPE. Requires setting the CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_LD attribute.	void*	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER
EPILOGUE_AUX_LD	Leading dimension for epilogue auxiliary buffer. If CUBLASMP_MATMUL_EPILOGUE_RELU_AUX, CUBLASMP_MATMUL_EPILOGUE_RELU_AUX_BIAS, CUBLASMP_MATMUL_EPILOGUE_DRELU, or CUBLASMP_MATMUL_EPILOGUE_DRELU_BGRAD epilogue is used - ReLu bit-mask matrix leading dimension in elements (i.e. bits). Must be divisible by 128 and be no less than the number of rows in the output matrix. If CUBLASMP_MATMUL_EPILOGUE_GELU_AUX_BIAS, CUBLASMP_MATMUL_EPILOGUE_DGELU, or CUBLASMP_MATMUL_EPILOGUE_DGELU_BGRAD epilogue is used - GELU input matrix leading dimension in elements. Must be divisible by 8 and be no less than the number of rows in the output matrix.	int64_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_LD
EPILOGUE_AUX_BATCH_STRIDE	Stride in bytes between auxiliary output buffers in batched epilogue operations. Currently, batched matmul is not supported and this parameter has to be to 0 (default value).	int64_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_BATCH_STRIDE
EPILOGUE_AUX_DATA_TYPE	The type of the data that will be stored in CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER. ReLU bit-mask epilogues use CUDA_R_8U storage and this attribute may remain unset or be set to CUDA_R_8U. For other epilogues, an unset attribute follows the cuBLASLt default. With CUDA 13 or newer, FP8/FP4 matmuls use the matrix C element type (CType). With CUDA 12.x, FP8/FP4 matmuls use CUDA_R_16BF, except that CUDA 12.9 uses CUDA_R_32F for non-ReLU/DReLU epilogues when DType is CUDA_R_32F. Other matmuls use the matrix D element type (DType). A non-bitmask auxiliary epilogue returns CUBLASMP_STATUS_NOT_SUPPORTED when this resolves to CUDA_R_16BF while CType is CUDA_R_16F; set an explicitly supported auxiliary type to use that matrix-type combination. Explicitly supported types depend on the CUDA version and matmul type combination; see the cuBLASLt CUBLASLT_MATMUL_DESC_EPILOGUE_AUX_DATA_TYPE restrictions.	cudaDataType_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_DATA_TYPE
EPILOGUE_AUX_SCALE_POINTER	Device pointer to the scaling factor value to convert results from compute type data range to storage data range in the auxiliary matrix that is set via CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER. The scaling factor value must have the same type as the compute type. If not specified, the scaling factor is assumed to be 1.	void*	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_SCALE_POINTER
EPILOGUE_AUX_AMAX_POINTER	Device pointer to receive the maximum absolute auxiliary value before applying EPILOGUE_AUX_SCALE_POINTER and converting to the auxiliary storage data type. The computed value has the same type as the compute type. If not specified, the maximum absolute value is not computed.	void*	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_AMAX_POINTER
EPILOGUE_AUX_SCALE_MODE	Scaling mode that defines how the matrix scaling factor for the auxiliary matrix is interpreted. Default value: CUBLASMP_MATMUL_MATRIX_SCALE_SCALAR_FP32.	cublasMpMatmulMatrixScale_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_SCALE_MODE
A_SCALE_POINTER	Device pointer to the scale factor value that converts data in matrix A to the compute data type range. The scaling factor must have the same type as the compute type. Matrix scaling is supported only when compute type is CUBLAS_COMPUTE_32F and the output matrix type is not complex. If not specified, the scaling factor is assumed to be 1.	const void*	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_A_SCALE_POINTER
B_SCALE_POINTER	Equivalent to CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_A_SCALE_POINTER for matrix B.	const void*	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_B_SCALE_POINTER
C_SCALE_POINTER	Equivalent to CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_A_SCALE_POINTER for matrix C.	const void*	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_C_SCALE_POINTER
D_SCALE_POINTER	Equivalent to CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_A_SCALE_POINTER for matrix D.	const void*	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_D_SCALE_POINTER
A_SCALE_MODE	Scaling mode that defines how the matrix scaling factor for matrix A is interpreted. Default value: CUBLASMP_MATMUL_MATRIX_SCALE_SCALAR_FP32.	cublasMpMatmulMatrixScale_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_A_SCALE_MODE
B_SCALE_MODE	Scaling mode for matrix B. For ReduceScatter/AllReduce with transB == CUBLAS_OP_N, explicit SPLIT_P2P (ReduceScatter only) and SPLIT_MULTICAST requests use independently padded, chunk-packed cuBLASLt B scales for VEC16/VEC32/VEC128. DEFAULT and NO_OVERLAP retain the full rank-local scale layout. BLK128x128 still requires 128-column-aligned chunks; deprecated ATOMIC_MULTICAST uses one full-tensor GEMM without B-scale column-chunk restrictions. Default value: CUBLASMP_MATMUL_MATRIX_SCALE_SCALAR_FP32.	cublasMpMatmulMatrixScale_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_B_SCALE_MODE
C_SCALE_MODE	Scaling mode that defines how the matrix scaling factor for matrix C is interpreted. Default value: CUBLASMP_MATMUL_MATRIX_SCALE_SCALAR_FP32.	cublasMpMatmulMatrixScale_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_C_SCALE_MODE
D_SCALE_MODE	Scaling mode that defines how the matrix scaling factor for matrix D is interpreted. Default value: CUBLASMP_MATMUL_MATRIX_SCALE_SCALAR_FP32.	cublasMpMatmulMatrixScale_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_D_SCALE_MODE
AMAX_D_POINTER	Device pointer to the memory location that on completion will be set to the maximum of absolute values in the output matrix. The computed value has the same type as the compute type. If not specified, the maximum absolute value is not computed.	void*	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_AMAX_D_POINTER
D_OUT_SCALE_POINTER	Device pointer to the scale factors that are used to convert data in matrix D to the compute data type range. The scaling factor value type is defined by the scaling mode (see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_D_OUT_SCALE_MODE). If set for an unsupported matrix data, scale, scale mode, and compute type combination, or missing for a supported combination, then calling cublasMpMatmul() will return CUBLASMP_STATUS_INVALID_VALUE. Default value: NULL.	void*	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_D_OUT_SCALE_POINTER
D_OUT_SCALE_MODE	Scaling mode that defines how the output matrix scaling factor for matrix D is interpreted. Default value: 0. See cublasMpMatmulMatrixScale_t for supported values.	cublasMpMatmulMatrixScale_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_D_OUT_SCALE_MODE
COMMUNICATION_TYPE	Data type used for inter-GPU communication in distributed matmul operations. When specified, GEMM output is produced in this type before the collective operation and post-processing (beta*C, epilogue, d_scale, d_out_scale) is applied after the reduce. Must satisfy: d_type <= communication_type <= scale_type (by element size). Supported by all Matmul + ReduceScatter and Matmul + AllReduce algorithm variants. FP8 communication can round or overflow before final output conversion and is not equivalent to a compute-type reduction followed by one cast. Finite saturation on overflow is not guaranteed across algorithms. Choose a wider communication type if partials or their reduction can exceed the FP8 range. Default value: scale type derived from compute type.	cudaDataType_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_COMMUNICATION_TYPE
BIAS_RESULT_SCHEME	Controls how epilogue output buffers (bias gradients) are handled after matrix multiplication. See cublasMpResultScheme_t for supported values. When set to FULL, performs AllReduce on the bias gradient output buffer to produce the full (reduced) result on all ranks. This is useful for Matmul + ReduceScatter algorithm where each rank computes partial sums of the bias gradient. Applicable to bias gradient epilogues (CUBLASMP_MATMUL_EPILOGUE_DRELU_BGRAD, CUBLASMP_MATMUL_EPILOGUE_DGELU_BGRAD, CUBLASMP_MATMUL_EPILOGUE_BGRADA, CUBLASMP_MATMUL_EPILOGUE_BGRADB) when used with Matmul + ReduceScatter algorithm. Default value: CUBLASMP_RESULT_SCHEME_DEFAULT.	cublasMpResultScheme_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_BIAS_RESULT_SCHEME
POINTER_MODE	Pointer mode of the alpha and beta scalars passed to cublasMpMatmul(). When set to CUBLASMP_POINTER_MODE_DEVICE, alpha and beta are dereferenced on the stream from device memory, avoiding a device-to-host copy and synchronization when the scalars are produced by a preceding GPU kernel. Applies to both alpha and beta. Default value: CUBLASMP_POINTER_MODE_HOST.	cublasMpPointerMode_t	CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_POINTER_MODE

cublasMpPointerMode_t

An enumerated type used to specify whether the alpha and beta scalars passed to cublasMpMatmul() reside in host or device memory. Set via the CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_POINTER_MODE attribute. Mirrors the semantics of cuBLASLt’s pointer mode and applies to both alpha and beta.

Value	Meaning
CUBLASMP_POINTER_MODE_HOST	alpha and beta point to host memory and are read on the host. This is the default.
CUBLASMP_POINTER_MODE_DEVICE	alpha and beta point to device memory and are dereferenced on the stream. When C is NULL, beta is ignored. Otherwise C must point to valid device storage, but its values do not affect the result when beta == 0. The CUBLASMP_MATMUL_ALGO_TYPE_FUSED Matmul + AllReduce algorithm does not support this mode; both the workspace query and execution return CUBLASMP_STATUS_NOT_SUPPORTED.

cublasMpMatmulAlgoType_t

Matrix-matrix multiplication algorithm types. DEFAULT lets cuBLASMp choose a supported implementation and use its fallbacks. An explicitly selected algorithm is required and returns CUBLASMP_STATUS_NOT_SUPPORTED when unavailable.

Value	Meaning
CUBLASMP_MATMUL_ALGO_TYPE_DEFAULT	Default algorithm.
CUBLASMP_MATMUL_ALGO_TYPE_SPLIT_P2P	Use split matmul with p2p communication.
CUBLASMP_MATMUL_ALGO_TYPE_SPLIT_MULTICAST	Use split matmul with multicast communication. Only valid for Matmul + ReduceScatter and Matmul + AllReduce on Compute Capability 9.0 (Hopper) and newer GPUs connected with an NVSwitch. Each internal reduction chunk must occupy a multiple of 16 bytes in the communication type.
CUBLASMP_MATMUL_ALGO_TYPE_ATOMIC_MULTICAST	Use atomic matmul with multicast communication. Only valid for Matmul + ReduceScatter on Compute Capability 9.0 (Hopper) and newer GPUs connected with an NVSwitch. Each internal reduction chunk must occupy a multiple of 16 bytes in the communication type. [DEPRECATED]
CUBLASMP_MATMUL_ALGO_TYPE_NO_OVERLAP	Perform computation and communication sequentially without overlap. This algorithm may perform better on non-NVLINK multi-node systems.
CUBLASMP_MATMUL_ALGO_TYPE_FUSED	Use fused GEMM+AllReduce. Only valid for GEMM+AllReduce on Compute Capability 9.0 (Hopper), 10.0, and 10.3 (datacenter Blackwell) GPUs connected with an NVSwitch. Requires NCCL symmetric memory with multimem support and a local output row count that is 16-byte aligned for the algorithm’s communication type. Supports real FP16/BF16 and FP8 precision combinations available to the fused algorithm. FP4 fused combinations require Compute Capability 10.0 or 10.3.

cublasMpMatmulMatrixScale_t

An enumerated type used to specify the scaling mode that defines how scaling factor pointers are interpreted.

Value	Meaning
CUBLASMP_MATMUL_MATRIX_SCALE_SCALAR_FP32	Scaling factors are single-precision scalars applied to the whole tensors. Supported with Compute Capability 8.9 and above.
CUBLASMP_MATMUL_MATRIX_SCALE_VEC16_UE4M3	One 8-bit CUDA_R_8F_UE4M3 scale per 16-element block in the innermost dimension. Uses the cuBLASLt tiled 32x4x4 swizzle, not a linear column-major array; see the layout below. Total size: roundup(ceil(m/16), 4) * roundup(n, 128) bytes per cuBLASLt tensor. Explicit RS/AR splits with non-transposed B use one B-scale tensor per column chunk. Supported with Compute Capability 10.0 and above.
CUBLASMP_MATMUL_MATRIX_SCALE_VEC32_UE8M0	One 8-bit CUDA_R_8F_UE8M0 scale per 32-element block in the innermost dimension. Uses the cuBLASLt tiled 32x4x4 swizzle, not a linear column-major array; see the layout below. Total size: roundup(ceil(m/32), 4) * roundup(n, 128) bytes per cuBLASLt tensor. Explicit RS/AR splits with non-transposed B use one B-scale tensor per column chunk. Supported with Compute Capability 10.0 and above.
CUBLASMP_MATMUL_MATRIX_SCALE_OUTER_VEC_FP32	Scaling factors are vectors of CUDA_R_32F values. This mode is only applicable to matrices A and B, in which case the vectors are expected to have M and N elements respectively, and each (i, j)-th element of product of A and B is multiplied by i-th element of A scale and j-th element of B scale. Supported with Compute Capability 9.0.
CUBLASMP_MATMUL_MATRIX_SCALE_VEC128_FP32	Scaling factors are tensors that contain a dedicated CUDA_R_32F scaling factor for each 128-element block in the innermost dimension of the corresponding data tensor. Explicit RS/AR splits with non-transposed B use one B-scale tensor per column chunk. Supported with Compute Capability 9.0.
CUBLASMP_MATMUL_MATRIX_SCALE_BLK128x128_FP32	Scaling factors are tensors that contain a dedicated CUDA_R_32F scaling factor for each 128x128-element block in the corresponding data tensor. Supported with Compute Capability 9.0.

For VEC16/VEC32, the following cuBLASLt tiled layout maps scale row sr (matrix row divided by 16 or 32) and stored column col to a byte offset. All divisions below are integer divisions. For chunk-packed B scales, use the dimensions and column index within each chunk and add that chunk’s byte offset separately:

ld = roundup(ceil(rows / vector_size), 4)
tile = sr / 4 + (col / 128) * (ld / 4)
offset = tile * 512 + (col % 32) * 16 + ((col / 32) % 4) * 4 + sr % 4

See cuBLASLt’s block scaling layout documentation for the tensor layout and padding requirements. A tensors and non-split B tensors use their full rank-local stored dimensions.

cublasMpMatmulEpilogue_t

An enumerated type used to specify the epilogue operation to be performed after matrix multiplication. Epilogues enable fusion of additional operations such as activation functions, bias addition, and their derivatives for training, providing better performance by avoiding separate kernel launches.

Value (CUBLASMP_MATMUL_ EPILOGUE_*)	Meaning	Supported Matmul algorithms	Value (complete name)
DEFAULT	No special postprocessing.	all	CUBLASMP_MATMUL_EPILOGUE_DEFAULT
ALLREDUCE	Performs AllReduce communication operation after matrix multiplication.	all	CUBLASMP_MATMUL_EPILOGUE_ALLREDUCE
RELU	Apply ReLU point-wise transform to the results (x := max(x, 0)).	all	CUBLASMP_MATMUL_EPILOGUE_RELU
RELU_AUX	Apply ReLU point-wise transform to the results (x := max(x, 0)). This epilogue mode produces an extra output, see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER of cublasMpMatmulDescriptorAttribute_t.	all	CUBLASMP_MATMUL_EPILOGUE_RELU_AUX
BIAS	Apply (broadcast) bias from the bias vector. Bias vector length must match matrix D rows, and it must be packed (such as stride between vector elements is 1). Bias vector is broadcast to all columns and added before applying the final postprocessing.	AllGather + Matmul, Matmul + ReduceScatter, Matmul + AllReduce	CUBLASMP_MATMUL_EPILOGUE_BIAS
RELU_BIAS	Apply bias and then ReLU transform.	AllGather + Matmul, Matmul + ReduceScatter, Matmul + AllReduce	CUBLASMP_MATMUL_EPILOGUE_RELU_BIAS
DRELU	Apply ReLU gradient to matmul output. Store ReLU gradient in the output matrix. This epilogue mode requires an extra input (ReLU bit-mask from forward pass), see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER of cublasMpMatmulDescriptorAttribute_t.	AllGather + Matmul, Matmul + ReduceScatter, Matmul + AllReduce	CUBLASMP_MATMUL_EPILOGUE_DRELU
DRELU_BGRAD	Apply independently ReLU and Bias gradients to matmul output. Store ReLU gradient in the output matrix, and Bias gradient in the bias buffer (see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_BIAS_POINTER of cublasMpMatmulDescriptorAttribute_t). This epilogue mode requires an extra input (ReLU bit-mask from forward pass), see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER of cublasMpMatmulDescriptorAttribute_t.	AllGather + Matmul, Matmul + ReduceScatter, Matmul + AllReduce	CUBLASMP_MATMUL_EPILOGUE_DRELU_BGRAD
RELU_AUX_BIAS	Apply bias and then ReLU transform. This epilogue mode produces an extra output, see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER of cublasMpMatmulDescriptorAttribute_t.	AllGather + Matmul, Matmul + ReduceScatter, Matmul + AllReduce	CUBLASMP_MATMUL_EPILOGUE_RELU_AUX_BIAS
GELU	Apply GELU point-wise transform to the results (x := GELU(x)).	all	CUBLASMP_MATMUL_EPILOGUE_GELU
GELU_AUX	Apply GELU point-wise transform to the results (x := GELU(x)). This epilogue mode outputs GELU input as a separate matrix (useful for training). See CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER of cublasMpMatmulDescriptorAttribute_t.	all	CUBLASMP_MATMUL_EPILOGUE_GELU_AUX
GELU_BIAS	Apply Bias and then GELU transform.	AllGather + Matmul, Matmul + ReduceScatter, Matmul + AllReduce	CUBLASMP_MATMUL_EPILOGUE_GELU_BIAS
GELU_AUX_BIAS	Apply Bias and then GELU transform. This epilogue mode outputs GELU input as a separate matrix (useful for training). See CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER of cublasMpMatmulDescriptorAttribute_t.	AllGather + Matmul, Matmul + ReduceScatter, Matmul + AllReduce	CUBLASMP_MATMUL_EPILOGUE_GELU_AUX_BIAS
DGELU	Apply GELU gradient to matmul output. Store GELU gradient in the output matrix. This epilogue mode requires an extra input, see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER of cublasMpMatmulDescriptorAttribute_t.	AllGather + Matmul, Matmul + ReduceScatter, Matmul + AllReduce	CUBLASMP_MATMUL_EPILOGUE_DGELU
DGELU_BGRAD	Apply independently GELU and Bias gradient to matmul output. Store GELU gradient in the output matrix, and Bias gradient in the bias buffer (see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_BIAS_POINTER of cublasMpMatmulDescriptorAttribute_t). This epilogue mode requires an extra input, see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_EPILOGUE_AUX_POINTER of cublasMpMatmulDescriptorAttribute_t.	AllGather + Matmul, Matmul + ReduceScatter, Matmul + AllReduce	CUBLASMP_MATMUL_EPILOGUE_DGELU_BGRAD
BGRADA	Apply Bias gradient to the input matrix A. The bias size corresponds to the number of rows of the matrix D. The reduction happens over the GEMM’s “k” dimension. Store Bias gradient in the bias buffer, see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_BIAS_POINTER of cublasMpMatmulDescriptorAttribute_t.	all	CUBLASMP_MATMUL_EPILOGUE_BGRADA
BGRADB	Apply Bias gradient to the input matrix B. The bias size corresponds to the number of columns of the matrix D. The reduction happens over the GEMM’s “k” dimension. Store Bias gradient in the bias buffer, see CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_BIAS_POINTER of cublasMpMatmulDescriptorAttribute_t.	all	CUBLASMP_MATMUL_EPILOGUE_BGRADB

cublasMpResultScheme_t

An enumerated type used to specify how epilogue output buffers (such as bias gradients) are handled after matrix multiplication in distributed algorithms.

Value	Meaning
CUBLASMP_RESULT_SCHEME_DEFAULT	Use default behavior. Equivalent to CUBLASMP_RESULT_SCHEME_PARTIAL.
CUBLASMP_RESULT_SCHEME_PARTIAL	Leave epilogue output as partial sums. Each rank retains its local partial result. The user is responsible for performing any required reduction externally.
CUBLASMP_RESULT_SCHEME_FULL	Perform AllReduce on the epilogue output buffer to produce the full (reduced) result on all ranks. This is useful for Matmul + ReduceScatter where each rank computes partial sums of the bias gradient.