.. meta:: :description: MI300 features for hipSPARSELt :keywords: hipSPARSELt, ROCm, MI300 .. _mi300-features: ********************************************************************** hipSPARSELt features for the Instinct MI300 series ********************************************************************** hipSPARSELt provides hardware acceleration support for sparse matrix multiplication operations on AMD Instinct™ MI300 series GPUs using SMFMA (Sparse Matrix Fused Multiply Add) matrix instructions. For hardware-accelerated sparse-dense matrix operations, the following conditions apply: * One matrix must be structurally sparse. Two out of every four elements on the K axis must be zero. * The other matrix must be dense. * The output matrix serves as the accumulation (destination) matrix. * The sparse indexes determine which two elements out of every four are non-zero, where for each index pair, ``index0 < index1`` and ``index0 != index1``. Data types and precision ======================== While the Instinct MI300 series supports multiple data types for hardware-accelerated sparse matrix operations, the hipSPARSELt library currently enables hardware acceleration for a subset of these types: .. list-table:: :header-rows: 1 * - Input/output type - Library data type - Hardware acceleration in hipSPARSELt * - float16 - ``HIP_R_16F`` - ✅ * - bfloat16 - ``HIP_R_16BF`` - ✅ * - int8 - ``HIP_R_8I`` - ✅ * - float8_fnuz - ``HIP_R_8F_E4M3_FNUZ`` - ✅ * - bfloat8_fnuz - ``HIP_R_8F_E5M2_FNUZ`` - ✅ .. note:: All floating-point operations accumulate in ``float32``, while integer operations accumulate in ``int32``. For comprehensive information on the supported data types and their characteristics, see the :ref:`Data type support ` page. Matrix format requirements ========================== For MI300 hardware acceleration, the matrices must meet the specific format requirements for the matrix type. Sparse matrix ------------- Sparse matrices must: * Follow a 2:4 structured sparsity pattern * Use a compressed sparse format * Meet the alignment requirements for the selected data type Dense matrix ------------ Dense matrices must: * Be in standard dense format * Meet the alignment requirements for the selected data type .. note:: The hipSPARSELt library supports configurations where either the first or second operand is the structured sparse matrix. The other operand must be a dense matrix. This flexibility is achieved through library-level handling that adjusts the memory layout and operation order accordingly. Matrix preparation ================== The sparse matrix must be prepared for SMFMA operations through a two-step process of pruning and compression. Understanding the pruning and compression process ------------------------------------------------- **Pruning** involves converting a dense matrix into a structured sparse format by selectively zeroing out elements according to a specific pattern. For MI300 SMFMA operations, the matrix must use the 2:4 pattern (50% sparsity) where exactly two out of every four consecutive elements along the K dimension are preserved, with the rest set to zero. **Compression** follows pruning and converts the pruned matrix into a hardware-optimized format that stores only the non-zero values and their corresponding indices. This reduces the memory requirements and enables direct use of the SMFMA hardware instructions. Pruning and compression workflow -------------------------------- The complete process involves several steps: #. Initialize the matrix descriptors and matmul descriptor. #. Determine the workspace requirements for compression. #. Allocate memory for the pruned and compressed matrices. #. Perform pruning using the 2:4 sparsity pattern. #. Verify the pruning was successful. #. Compress the pruned matrix. .. figure:: ../data/prune-workflow.svg Detailed implementation example ------------------------------- Here's a detailed implementation of the pruning and compression process: .. code-block:: cpp // Step 1: Set up matrix and matmul descriptors (see Descriptor Configuration section) // ... // Step 2: Get workspace sizes for compression size_t workspace_size, compressed_size, compress_buffer_size; CHECK_HIPSPARSELT_ERROR(hipsparseLtSpMMACompressedSize( &handle, &plan, &compressed_size, &compress_buffer_size)); // Step 3: Allocate memory // Allocate compression buffers void* d_compressed = nullptr; // Will hold the final compressed matrix void* d_compressBuffer = nullptr; // Temporary buffer for compression CHECK_HIP_ERROR(hipMalloc(&d_compressed, compressed_size)); CHECK_HIP_ERROR(hipMalloc(&d_compressBuffer, compress_buffer_size)); // Allocate temporary buffer for the pruned (but not yet compressed) matrix void* dp = nullptr; CHECK_HIP_ERROR(hipMalloc(&dp, matrix_size)); // Step 4: Prune matrix A using 2:4 sparsity pattern // The original dense matrix 'dA' is pruned into the temporary buffer 'dp' // HIPSPARSELT_PRUNE_SPMMA_STRIP selects the two largest elements in each 1x4 strip CHECK_HIPSPARSELT_ERROR(hipsparseLtSpMMAPrune( &handle, &matmul, dA, dp, HIPSPARSELT_PRUNE_SPMMA_STRIP, stream)); // Step 5: Verify pruning success // Check if the pruned matrix maintains the required 2:4 sparsity pattern bool is_valid = false; CHECK_HIPSPARSELT_ERROR(hipsparseLtSpMMAPruneCheck( &handle, &matmul, dp, &is_valid, stream)); if (!is_valid) { fprintf(stderr, "Error: Matrix pruning failed to achieve required sparsity pattern\n"); goto cleanup; } // Step 6: Compress the pruned matrix // Convert the pruned matrix 'dp' into the compressed format 'd_compressed' // This format is optimized for direct use with the MI300 SMFMA hardware instructions CHECK_HIPSPARSELT_ERROR( hipsparseLtSpMMACompress(&handle, &plan, dp, d_compressed, d_compressBuffer, stream)); // Once completed, 'd_compressed' contains the matrix in a format ready for SMFMA operations // This compressed matrix is used in place of the original dense matrix in subsequent operations // Always clean up resources when done cleanup: if (dp) CHECK_HIP_ERROR(hipFree(dp)); if (d_compressBuffer) CHECK_HIP_ERROR(hipFree(d_compressBuffer)); // Don't free d_compressed yet as it's needed for matmul operations Pruning options and parameters ------------------------------ The pruning function supports different strategies through the ``alg`` parameter: * ``HIPSPARSELT_PRUNE_SPMMA_STRIP``: Zeroes out two elements in a 1x4 strip. Non-zero elements have the maximum L1-norm value in all combinations in the strip. * ``HIPSPARSELT_PRUNE_SPMMA_TILE``: Zeroes out eight elements in a 4x4 tile. Non-zero elements have the maximum L1-norm value in all combinations in the tile. There are exactly two elements in each row and column. Performance considerations -------------------------- * Pruning is typically performed once during initialization, not in the critical computation path. * The quality of pruning affects the accuracy of matrix multiplication results. * For matrices that don't naturally fit the 2:4 pattern, consider pretraining with structured sparsity constraints. Matrix operation setup ====================== The following sections describe how to set up matrix operations with SMFMA, including descriptor configuration and execution plans. Descriptor configuration ------------------------ Matrix descriptors define the properties and memory layout for each matrix in the operation. For SMFMA: * Matrix A requires a structured descriptor with 50% sparsity (2:4 pattern). * Matrix B uses a dense descriptor. * The compute type should be set to ``FP32`` accumulation on AMD platforms. .. code-block:: cpp // Initialize sparse matrix A descriptor CHECK_HIPSPARSELT_ERROR( hipsparseLtStructuredDescriptorInit(&handle, &matA, row_a, col_a, lda, 16, HIP_R_16F, HIPSPARSE_ORDER_COL, HIPSPARSELT_SPARSITY_50_PERCENT)); // Initialize dense matrix B descriptor CHECK_HIPSPARSELT_ERROR(hipsparseLtDenseDescriptorInit( &handle, &matB, row_b, col_b, ldb, 16, HIP_R_16F, HIPSPARSE_ORDER_COL)); // Initialize compute type auto compute_type = HIPSPARSELT_COMPUTE_32F; // For AMD platforms // Initialize matmul descriptor CHECK_HIPSPARSELT_ERROR(hipsparseLtMatmulDescriptorInit( &handle, &matmul, trans_a, trans_b, &matA, &matB, &matC, &matD, compute_type)); Plan creation and execution --------------------------- Create and execute the matrix multiplication plan: .. code-block:: cpp // Initialize matmul plan CHECK_HIPSPARSELT_ERROR(hipsparseLtMatmulAlgSelectionInit( &handle, &alg_sel, &matmul, HIPSPARSELT_MATMUL_ALG_DEFAULT)); CHECK_HIPSPARSELT_ERROR(hipsparseLtMatmulPlanInit(&handle, &plan, &matmul, &alg_sel)); // Get and allocate workspace if needed size_t workspace_size = 0; CHECK_HIPSPARSELT_ERROR(hipsparseLtMatmulGetWorkspace(&handle, &plan, &workspace_size)); void* d_workspace = nullptr; if(workspace_size > 0) { CHECK_HIP_ERROR(hipMalloc(&d_workspace, workspace_size)); } // Configure and create stream int num_streams = 1; // Default to single stream hipStream_t streams[1]; CHECK_HIP_ERROR(hipStreamCreate(&streams[0])); // Execute matmul operation CHECK_HIPSPARSELT_ERROR(hipsparseLtMatmul(&handle, &plan, &alpha, d_compressed, // Compressed sparse matrix A dB, // Dense matrix B &beta, dC, dD, d_workspace, &streams[0], num_streams)); // Cleanup resources CHECK_HIP_ERROR(hipStreamDestroy(streams[0])); if (d_workspace) CHECK_HIP_ERROR(hipFree(d_workspace)); Stream usage notes ------------------ * Single stream is sufficient for most operations. * Multiple streams can improve performance when processing independent operations. * Ensure proper synchronization when using multiple streams with dependent operations. * The stream count is specified using the ``num_streams`` parameter. Supported operations ==================== hipSPARSELt with MI300 hardware acceleration supports several types of matrix operations, described in detail below. Standard matrix multiplication ------------------------------ The basic operation performs sparse-dense matrix multiplication with optional scaling: .. code-block:: cpp // D = α × op(sparse_matrix) × op(dense_matrix) + β × C // Note: Either the first or second operand can be the sparse matrix CHECK_HIPSPARSELT_ERROR(hipsparseLtMatmul(&handle, &plan, &alpha, compressed_matrix, // Compressed sparse matrix dense_matrix, // Dense matrix &beta, dC, dD, d_workspace, &streams[0], num_streams)); Batched operations ------------------ hipSPARSELt supports multiple types of batched operations, which provide support for MI300 hardware acceleration: * Broadcast mode (single sparse matrix/multiple dense matrices) A single sparse matrix is multiplied with multiple dense matrices. .. code-block:: cpp // Configure sparse matrix (constant across all batch operations) CHECK_HIPSPARSELT_ERROR(hipsparseLtMatDescSetAttribute( &handle, &sparse_mat, HIPSPARSELT_MAT_NUM_BATCHES, &batch_count, sizeof(batch_count))); CHECK_HIPSPARSELT_ERROR(hipsparseLtMatDescSetAttribute( &handle, &sparse_mat, HIPSPARSELT_MAT_BATCH_STRIDE, &zero_stride, sizeof(zero_stride))); // Configure dense matrices (different for each batch operation) CHECK_HIPSPARSELT_ERROR(hipsparseLtMatDescSetAttribute( &handle, &dense_mat, HIPSPARSELT_MAT_NUM_BATCHES, &batch_count, sizeof(batch_count))); CHECK_HIPSPARSELT_ERROR(hipsparseLtMatDescSetAttribute( &handle, &dense_mat, HIPSPARSELT_MAT_BATCH_STRIDE, &dense_stride, sizeof(dense_stride))); * Multiple sparse and dense matrices Different sparse and dense matrix pairs are multiplied in each batch operation. .. code-block:: cpp // Configure sparse matrices (different for each batch operation) CHECK_HIPSPARSELT_ERROR(hipsparseLtMatDescSetAttribute( &handle, &sparse_mat, HIPSPARSELT_MAT_NUM_BATCHES, &batch_count, sizeof(batch_count))); CHECK_HIPSPARSELT_ERROR(hipsparseLtMatDescSetAttribute( &handle, &sparse_mat, HIPSPARSELT_MAT_BATCH_STRIDE, &sparse_stride, sizeof(sparse_stride))); // Configure dense matrices (different for each batch operation) CHECK_HIPSPARSELT_ERROR(hipsparseLtMatDescSetAttribute( &handle, &dense_mat, HIPSPARSELT_MAT_NUM_BATCHES, &batch_count, sizeof(batch_count))); CHECK_HIPSPARSELT_ERROR(hipsparseLtMatDescSetAttribute( &handle, &dense_mat, HIPSPARSELT_MAT_BATCH_STRIDE, &dense_stride, sizeof(dense_stride))); * Batched bias addition This feature can be combined with matrix multiplication operations to add different bias vectors for each batch operation. The bias addition is supported as a fused operation alongside hardware-accelerated matrix multiplication. .. code-block:: cpp // Configure batched bias vectors void* bias_ptr = d_bias; // Pointer to the first bias vector size_t bias_stride = bias_vector_size * sizeof(data_type); CHECK_HIPSPARSELT_ERROR(hipsparseLtMatmulDescriptorSetAttribute( &handle, &matmul, HIPSPARSELT_MATMUL_BIAS_POINTER, &bias_ptr, sizeof(void*))); CHECK_HIPSPARSELT_ERROR(hipsparseLtMatmulDescriptorSetAttribute( &handle, &matmul, HIPSPARSELT_MATMUL_BIAS_BATCH_STRIDE, &bias_stride, sizeof(bias_stride))); Fused operations ---------------- Matrix multiplication operations can be fused with additional operations to improve performance: .. code-block:: cpp // Configure bias addition void* bias_ptr = nullptr; CHECK_HIPSPARSELT_ERROR(hipsparseLtMatmulDescriptorSetAttribute( &handle, &matmul, HIPSPARSELT_MATMUL_BIAS_POINTER, &bias_ptr, sizeof(void*))); // Configure activation function (for example, ReLU) hipsparseLtActivationType_t act_type = HIPSPARSELT_ACTIVATION_RELU; CHECK_HIPSPARSELT_ERROR(hipsparseLtMatmulDescriptorSetAttribute( &handle, &matmul, HIPSPARSELT_MATMUL_ACTIVATION_TYPE, &act_type, sizeof(act_type))); Operations are executed in the following order: 1. Sparse-dense matrix multiplication 2. Scaling by alpha/beta 3. Bias addition (if configured) 4. Activation function (if configured) This fusion helps maximize performance by reducing intermediate memory operations.