rocWMMA

https://rocm.docs.amd.com/projects/rocWMMA/en/latest/
Admin
rocWMMA is a C++ header library for accelerating mixed-precision matrix multiply-accumulate...
Tokens:13,086
Snippets:158
Trust Score:10
Update:5 months ago
Show doc for...
Tokens:
Raw
### Install rocWMMA on Ubuntu/Debian

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Use apt-get to install rocWMMA development files, samples, and tests on Ubuntu or Debian-based systems.

```bash
sudo apt-get update
sudo apt-get install rocwmma-dev rocwmma-samples rocwmma-tests

```

--------------------------------

### Install rocWMMA on SLES

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Use dnf to install rocWMMA development files, samples, and tests on SLES.

```bash
sudo dnf upgrade
sudo dnf install rocwmma-dev rocwmma-samples rocwmma-tests

```

--------------------------------

### Install rocWMMA on RHEL-based platforms

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Use yum to install rocWMMA development files, samples, and tests on RHEL or similar distributions.

```bash
sudo yum update
sudo yum install rocwmma-dev rocwmma-samples rocwmma-tests

```

--------------------------------

### Example of Deprecated Cooperative API Usage

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/conceptual/migration-guide.html

Illustrates the usage of deprecated cooperative load and store functions, including the manual calculation of warp count and index. Mismatched values can lead to incorrect behavior.

```cpp
// Global read (macro tile)
using GRBuffA = fragment<matrix_a, MACRO_TILE_X, ROCWMMA_N, ROCWMMA_K, InputT, DataLayoutA>;

// Local warp coordinate relative to current thread block (wg).
constexpr auto warpDims        = make_coord2d(WARPS_X, WARPS_Y);
auto           localWarpCoord  = make_coord2d(threadIdx.x / WARP_SIZE, threadIdx.y);

// WorkItems will be split up by minimum IOCount to perform either global read or local write.
// These are inputs to cooperative functions.
constexpr auto warpCount = get<0>(warpDims) * get<1>(warpDims);

// Scheduling warp order is analogous to row major priority.
// E.g. Wg = (128, 2) = 2x2 warps
// (0, 0)   (0, 1)   Share Schedule: w0 = (0, 0), w1 = (0, 1), 
// (1, 0)   (1, 1)                   w2 = (1, 0), w3 = (1, 1), count = 4
const auto warpIndex = get<0>(localWarpCoord) * get<1>(warpDims) + get<1>(localWarpCoord);

// Transfer data from global memory to local memory
GRBuffA grBuffA;
load_matrix_coop_sync<warpCount>(grBuffA, gAddrA, lda, warpIndex);
store_matrix_coop_sync<warpCount>(ldsAddr, applyDataLayout<DataLayoutLds, warpCount>(applyTranspose(grBuffA)), ldsld, warpIndex);
```

--------------------------------

### Verify ROCm Version on Ubuntu

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Use this command to check the installed ROCm version on an Ubuntu system.

```bash
apt show rocm-libs -a
```

--------------------------------

### Verify ROCm Version on RHEL

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Use this command to check the installed ROCm version on a RHEL-based system.

```bash
yum info rocm-libs
```

--------------------------------

### Basic rocWMMA Build Configuration

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Sets up the build environment for rocWMMA. Replace <build_dir> with your desired build directory.

```bash
CC=/opt/rocm/bin/amdclang CXX=/opt/rocm/bin/amdclang++ cmake -B <build_dir>
```

--------------------------------

### Build rocWMMA Library and Tests

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

After configuration, use this command to build the library and tests. Replace <nproc> with the number of parallel build processes.

```bash
cmake --build <build_dir> -- -j<nproc>

```

--------------------------------

### Download rocWMMA for ROCm 7.1.1 and Prior

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Clone the rocWMMA repository directly for ROCm versions 7.1.1 and earlier. Replace x.y with your ROCm version.

```bash
git clone -b release/rocm-rel-x.y https://github.com/ROCm/rocWMMA.git
cd rocWMMA
```

--------------------------------

### Full Fragment Support in rocWMMA

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/conceptual/migration-guide.html

Demonstrates the usage of full fragment sizes in rocWMMA, where fragment dimensions match minimum block sizes and no padding is used.

```cpp
// Fragment types, assuming ROCWMMA_MNK are minimum block sizes.
// These fragments will not use any padding.
using FragA = fragment<matrix_a, ROCWMMA_M, ROCWMMA_N, ROCWMMA_K, InputT, DataLayoutA>;
using FragB = fragment<matrix_b, ROCWMMA_M, ROCWMMA_N, ROCWMMA_K, InputT, DataLayoutB>;
using Accum = fragment<accumulator, ROCWMMA_M, ROCWMMA_N, ROCWMMA_K, AccumT>;

FragA fragA;
FragB fragB;
Accum accum;
fill_fragment(accum, 0);

load_matrix_sync(fragA, gAddrA, lda);
load_matrix_sync(fragB, gAddrB, ldb);
mma_sync(accum, fragA, fragB, accum);

store_matrix_sync(gResC, accum, ldc, layout_t::mem_row_major);

```

--------------------------------

### Configure rocWMMA Build with Assembly and Tests

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Use this command to configure the build with assembly code generation and tests enabled. Replace `<build_dir>` with your desired build directory.

```bash
CC=/opt/rocm/bin/amdclang CXX=/opt/rocm/bin/amdclang++ cmake -B <build_dir> . -DROCWMMA_BUILD_ASSEMBLY=ON -DROCWMMA_BUILD_TESTS=ON
```

--------------------------------

### Partial Fragment Support in rocWMMA

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/conceptual/migration-guide.html

Illustrates the usage of partial fragment sizes in rocWMMA, where fragment dimensions are smaller than minimum block sizes and padding is applied internally.

```cpp
// Fragment types, which are partial fragments.
// These fragments will use padding to minimum block sizes internally.
// Note: The dimensions (2, 3, 1) are smaller than BlockMNK, creating partial fragments
using FragA = fragment<matrix_a, 2, 3, 1, InputT, DataLayoutA>;
using FragB = fragment<matrix_b, 2, 3, 1, InputT, DataLayoutB>;
using Accum = fragment<accumulator, 2, 3, 1, AccumT>;

FragA fragA;
FragB fragB;
Accum accum;
fill_fragment(accum, 0);

load_matrix_sync(fragA, gAddrA, lda);
load_matrix_sync(fragB, gAddrB, ldb);
mma_sync(accum, fragA, fragB, accum);

store_matrix_sync(gResC, accum, ldc, layout_t::mem_row_major);

```

--------------------------------

### Build rocWMMA in Debug Mode

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Configures the build for debugging. Replace <build_dir> with your desired build directory.

```bash
CC=/opt/rocm/bin/amdclang CXX=/opt/rocm/bin/amdclang++ cmake -B <build_dir> . -DCMAKE_BUILD_TYPE=Debug
```

--------------------------------

### Configure rocWMMA Build with CMake

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Use this command to configure the build directory for rocWMMA. Set ROCcWMMA_BUILD_TESTS to OFF and ROCcWMMA_BUILD_SAMPLES to ON. Ensure the correct compilers are specified.

```bash
CC=/opt/rocm/bin/amdclang CXX=/opt/rocm/bin/amdclang++ cmake -B <build_dir> . -DROCWMMA_BUILD_TESTS=OFF -DROCWMMA_BUILD_SAMPLES=ON
```

--------------------------------

### Run Emulation Smoke Test

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Use this command to execute a subset of ROCWMMA test cases for emulators. Specify the build directory and the desired test mode (smoke, regression, or extended).

```bash
rtest.py --install_dir <build_dir> --emulation smoke
```

--------------------------------

### Build rocWMMA Library Only

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Configure and build only the rocWMMA library, disabling tests and samples. Assumes default Release mode.

```bash
CC=/opt/rocm/bin/amdclang CXX=/opt/rocm/bin/amdclang++ cmake -B <build_dir> . -DROCWMMA_BUILD_TESTS=OFF -DROCWMMA_BUILD_SAMPLES=OFF
```

--------------------------------

### Configure rocWMMA Build with Tests

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Use this command to configure the build directory with CMake, specifying the compiler and enabling test suite compilation.

```bash
CC=/opt/rocm/bin/amdclang CXX=/opt/rocm/bin/amdclang++ cmake -B <build_dir> . -DROCWMMA_BUILD_TESTS=ON

```

--------------------------------

### Download rocWMMA for ROCm 7.2.0+

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Clone the rocWMMA project from the rocm-libraries monorepo for ROCm versions 7.2.0 and later. Replace x.y with your ROCm version.

```bash
git clone -b release/rocm-rel-x.y https://github.com/ROCm/rocm-libraries.git
cd projects/rocwmma
```

--------------------------------

### Build rocWMMA Targeting Specific GPU

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Configures the build to target the gfx908 GPU with nack support. Replace <build_dir> with your desired build directory.

```bash
CC=/opt/rocm/bin/amdclang CXX=/opt/rocm/bin/amdclang++ cmake -B <build_dir> . -DGPU_TARGETS=gfx908:xnack-
```

--------------------------------

### Build rocWMMA Library

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Compiles the rocWMMA library after CMake configuration. Use -j<nproc> to specify the number of parallel build processes.

```bash
cmake --build <build_dir> -- -j<nproc>
```

--------------------------------

### Simplified Cooperative Fragment Usage in ROCWMMA 2.0

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/conceptual/migration-guide.html

Demonstrates the simplified usage of cooperative fragment operations in ROCWMMA 2.0. It utilizes a cooperative fragment scheduler to manage data distribution, eliminating the need for manual warp count and index calculations.

```cpp
// Global read (macro tile)
// Distribute segments of macro tile data between waves of the thread block in
// row major order.
using CoopScheduler = fragment_scheduler::coop_row_major_2d<TBLOCK_X, TBLOCK_Y>;
```

--------------------------------

### Fused Multiply-Add with Partial Fragments

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/conceptual/migration-guide.html

Shows that fused multiply-add operations remain valid for partial fragments, as padded elements are treated as zero.

```cpp
// Fused multiply-add still valid for partials as padded elements are 0
for(int i = 0; i < frag.num_elements; i++)
{
   frag.x[i] = frag.x[i] * (alpha + 1);
}


```

--------------------------------

### rocwmma::load_matrix_sync

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Loads the entire fragment from a data pointer. This function has two overloads, one that infers the data layout and one that allows manual specification of the data layout.

```APIDOC
## rocwmma::load_matrix_sync (infer layout)

### Description
Loads the entire fragment from the data pointer according to its matrix and data layout contexts. Data pointer may point to either local or global memory.

### Parameters
* **frag** (FragT&) - Fragment of type MatrixT with its associated block sizes, data type and layout
* **data** (const DataT*) - Data pointer to global or local memory
* **ldm** (uint32_t) - Leading dimension size

### Template Parameters
* **FragT** – Opaque fragment type
* **DataT** – Datatype
```

```APIDOC
## rocwmma::load_matrix_sync (manual layout)

### Description
Loads the entire fragment from the data pointer according to its matrix layout and data layout contexts. Data pointer may point to either local or global memory. This overload provides manual selection of data layout of the incoming memory pointer, which will be transformed to conform to the data layout of the fragment.

### Parameters
* **frag** (FragT&) - Fragment of type MatrixT with its associated block sizes, data type and layout
* **data** (const DataT*) - Data pointer to global/local memory
* **ldm** (uint32_t) - Leading dimension size
* **layout** (layout_t) - Data layout

### Template Parameters
* **FragT** – Opaque fragment type
* **DataT** – Datatype
```

--------------------------------

### rocwmma::store_matrix_sync

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Stores the entire fragment to a data pointer. This function has two overloads, one that infers the data layout and one that allows manual specification of the data layout.

```APIDOC
## rocwmma::store_matrix_sync (infer layout)

### Description
Stores the entire fragment to the data pointer according to its matrix and data layouts. Data pointer may point to either local or global memory.

### Parameters
* **data** (DataT*) - Data pointer to global/local memory
* **frag** (const FragT&) - Fragment of type MatrixT with its associated block sizes, data type and layout
* **ldm** (uint32_t) - Leading dimension size

### Template Parameters
* **FragT** – Opaque fragment type
* **DataT** – Datatype
```

```APIDOC
## rocwmma::store_matrix_sync (manual layout)

### Description
Stores the entire fragment to the data pointer according to its matrix layout and data layout contexts. Data pointer may point to either local or global memory. This overload provides manual selection of data layout of the outgoing memory pointer, which the data layout of the fragment will be transformed to.

### Parameters
* **data** (DataT*) - Data pointer to global/local memory
* **frag** (const FragT&) - Fragment of type MatrixT with its associated block sizes, data type and layout
* **ldm** (uint32_t) - Leading dimension size
* **layout** (layout_t) - Data layout

### Template Parameters
* **FragT** – Opaque fragment type
* **DataT** – Datatype
```

--------------------------------

### rocwmma::to_register_file

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Transforms the input fragment to a 'register file' fragment type. Register contents are directly mapped to a 2D matrix space.

```APIDOC
## rocwmma::to_register_file

### Description
Transforms the input fragment to a “register file” fragment type. Register contents are directly mapped to a 2D matrix space represented by [RegCount x WaveSize]. This transform is a geometry reinterpretation.

### Parameters
* **frag** (FragmentT) - Source fragment of type MatrixT with its associated block sizes, data type and layout.

### Template Parameters
* **FragT** - The source incoming fragment type.

### Returns
Target fragment after transformation.
```

--------------------------------

### Build rocWMMA Without rocBLAS Validation/Benchmarking

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Disables the use of rocBLAS for validation and benchmarking tests. Replace <build_dir> with your desired build directory.

```bash
CC=/opt/rocm/bin/amdclang CXX=/opt/rocm/bin/amdclang++ cmake -B <build_dir> . -DROCWMMA_VALIDATE_WITH_ROCBLAS=OFF -DROCWMMA_BENCHMARK_WITH_ROCBLAS=OFF
```

--------------------------------

### rocwmma::apply_data_layout

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Transforms the input fragment to have the desired data layout. This function allows changing the memory layout of the fragment.

```APIDOC
## rocwmma::apply_data_layout

### Description
Transforms the input fragment to have the desired data layout.

### Parameters
* **frag** (FragmentT) - Fragment of type MatrixT with its associated block sizes, data type and layout.

### Template Parameters
* **DataLayoutT** - The desired fragment data layout to apply.
* **FragT** - The incoming fragment type.

### Returns
Fragment with transformed data layout.
```

--------------------------------

### Declare Fragment with CoopScheduler

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/conceptual/migration-guide.html

Declare a fragment using the CoopScheduler template parameter to enable cooperative behavior. This simplifies matrix operations by abstracting wave count and index calculations.

```cpp
using GRBuffA = fragment<matrix_a, MACRO_TILE_X, ROCWMMA_N, ROCWMMA_K, InputT, DataLayoutA, CoopScheduler>;

// Transfer data from global memory to local memory
GRBuffA grBuffA;
load_matrix_sync(grBuffA, gAddrA, lda);
store_matrix_sync(ldsAddr, apply_data_layout<DataLayoutLds>(apply_transpose(grBuffA)), ldsld);
```

--------------------------------

### rocwmma::from_register_file

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Transforms the 'register file' fragment type to a target fragment type. Register contents are directly mapped to a 2D matrix space.

```APIDOC
## rocwmma::from_register_file

### Description
Transforms the “register file” fragment type to a target fragment type. Register contents are directly mapped to a 2D matrix space represented by [RegCount x WaveSize]. This transform is a geometry reinterpretation.

### Parameters
* **frag** (FragmentT) - Source fragment of type MatrixT with its associated block sizes, data type and layout.

### Template Parameters
* **DstFragT** - The target frag to transform to.
* **FragT** - The source incoming fragment type as register file.

### Returns
Fragment after transformation.
```

--------------------------------

### rocWMMA Layout Enumeration

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Defines the memory layout for matrices within rocWMMA.

```APIDOC
## rocWMMA enumeration layout_t

### Description
This enumeration defines the possible memory layouts for matrices when used with rocWMMA. It specifies whether the data is arranged in row-major or column-major order in memory.

### Values
- **mem_row_major**: Indicates that the matrix elements are stored in row-major order.
- **mem_col_major**: Indicates that the matrix elements are stored in column-major order.
```

--------------------------------

### Filter GoogleTest Cases

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Use GoogleTest filters to target specific test cases when running executables. Replace `<test_exe>` with the path to your test executable and `*name_filter*` with your desired filter.

```bash
<test_exe> --gtest_filter=\*name_filter\*
```

--------------------------------

### rocwmma::apply_fragment

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Transforms the input fragment to the target fragment type. This could include changing matrix context and/or changing data layout.

```APIDOC
## rocwmma::apply_fragment

### Description
Transforms the input fragment to the target fragment type. This could include changing matrix context and/or changing data layout, as long as there is a path from the source register layout to the destination register layout.

### Parameters
* **frag** (FragmentT) - Source fragment of type MatrixT with its associated block sizes, data type and layout.

### Template Parameters
* **DstFragT** - The target fragment type to transform to.
* **FragT** - The source incoming fragment type.

### Returns
Target fragment after transformation.
```

--------------------------------

### Deprecated ROCWMMA Cooperative Load/Store APIs

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/conceptual/migration-guide.html

These functions were previously used for cooperative data loading and storing in ROCWMMA. They required explicit wave count and wave index parameters.

```cpp
template <uint32_t WaveCount, typename MatrixT, uint32_t BlockM, uint32_t BlockN, uint32_t BlockK, typename DataT, typename DataLayoutT>
ROCWMMA_DEVICE void load_matrix_coop_sync(fragment<MatrixT, BlockM, BlockN, BlockK, DataT, DataLayoutT>& frag,
                                          const DataT*                                                   data,
                                          uint32_t                                                       ldm,
                                          uint32_t                                                       waveIndex);

template <uint32_t WaveCount, typename MatrixT, uint32_t BlockM, uint32_t BlockN, uint32_t BlockK, typename DataT, typename DataLayoutT>
ROCWMMA_DEVICE void store_matrix_coop_sync(DataT*                                                               data,
                                          fragment<MatrixT, BlockM, BlockN, BlockK, DataT, DataLayoutT> const& frag,
                                          uint32_t                                                             ldm,
                                          uint32_t                                                             waveIndex);
```

--------------------------------

### rocWMMA Fragment Class

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

The `fragment` class is a template class used for block-wise decomposition of matrix multiply-accumulate (mma) problems. It handles data loading, storing, mma operations, and transforms, abstracting away thread-level details for a simpler wavefront programming model.

```APIDOC
## class fragment

### Description
This class represents a fragment of data used in matrix multiply-accumulate operations. It is designed to work with a specific matrix context, block size, datatype, and memory layout. The fragment abstraction simplifies the programming model by handling internal thread-level details, allowing users to focus on block-wise decomposition.

### Template Parameters
- **MatrixT**: The context of the matrix (e.g., matrix_a, matrix_b, accumulator).
- **FragM/N/K**: The dimensions of the fragment.
- **DataT**: The datatype of the elements within the fragment (e.g., float, double).
- **DataLayoutT**: The in-memory layout of the data (e.g., `row_major`, `col_major`).
- **Scheduler**: The wave-wise scheduler to use.

### Public Types
- **IOTraits**: Input/output traits specific to the AMDGCN architecture, derived from the template parameters.

### Public Functions
- **operator[] (index)**: Accesses an unpacked element at the given index. Returns a mutable or immutable reference to the element.
- **operator*()**: Accesses the packed storage vector. Returns a mutable or immutable accessor to the storage.

### Public Members
- **union [anonymous]**: Internal data storage views, compatible with nvcuda::wmma.

### Public Static Functions
- **height()**: Returns the geometric height of the fragment.
- **width()**: Returns the geometric width of the fragment.
- **blockDim()**: Returns the leading block dimension (non-K).
- **kDim()**: Returns the K dimension.
- **size()**: Returns the size of the unpacked elements vector.

### Struct Traits
- **Public Types**:
  - **AccessT**: Represents the unpacked data access view.
  - **StorageT**: Represents the packed data storage view.
- **Public Static Attributes**:
  - **Size**: Asserts the fragment occupies at least one packed register and is equally splittable among the wave count.
```

--------------------------------

### Control Test Output and Verbosity

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/install/installation.html

Control test output redirection and verbosity using these command-line arguments. `--output_stream` redirects output to a file, and `--omit` can be used to exclude specific test results.

```bash
<test_exe> --output_stream "output.csv" --omit 1
```

--------------------------------

### rocwmma::synchronize_workgroup

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Synchronization point for all wavefronts in a workgroup. Guarantees pending reads/writes to LDS are flushed.

```APIDOC
## rocwmma::synchronize_workgroup

### Description
Synchronization point for all wavefronts in a workgroup. Guarantees pending reads / writes to LDS are flushed.

### Parameters
None
```

--------------------------------

### rocwmma::apply_transpose

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Applies the transpose transform to the input fragment. Transpose is defined as orthogonal matrix and data layout. E.g. T(fragment<matrix_a, BlockM, BlockN, BlockK, DataT, row_major>) = fragment<matrix_b, BlockN, BlockM, BlockK, DataT, col_major>

```APIDOC
## rocwmma::apply_transpose

### Description
Applies the transpose transform to the input fragment. Transpose is defined as orthogonal matrix and data layout.

### Parameters
* **frag** (FragmentT) - The input fragment of type MatrixT with its associated block sizes, data type and layout.

### Template Parameters
* **FragT** - The incoming fragment type.

### Returns
Transposed (orthogonal) fragment.
```

--------------------------------

### rocwmma::fill_fragment

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Fills the entire fragment with the desired value. This function is templated on the fragment type (FragT) and the data type (DataT).

```APIDOC
## rocwmma::fill_fragment

### Description
Fills the entire fragment with the desired value.

### Parameters
* **frag** (FragT&) - Fragment of type MatrixT with its associated block sizes, data type and layout
* **value** (DataT) - Fill value of type DataT

### Template Parameters
* **FragT** – Opaque fragment type
* **DataT** – Datatype
```

--------------------------------

### rocwmma::mma_sync

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Performs the Multiply-Accumulate operation on fragments A, B, and C, storing the result in fragment D (D = A * B + C). The accumulator fragment C can be the same as D.

```APIDOC
## rocwmma::mma_sync

### Description
Performs the Multiply-Accumulate operation on the fragments A, B, C and D (D = A * B + C)
Note
Frag c = d is valid

### Parameters
* **d** (FragAccumOut&) - Accumulator output D
* **a** (const FragA&) - Input fragment A
* **b** (const FragB&) - Input fragment B
* **c** (FragAccumIn&) - Input accumulator fragment C

### Template Parameters
* **FragA** – Opaque fragment type for matrix A data
* **FragB** – Opaque fragment type for matrix A data
* **FragAccumIn** – Opaque fragment type for input accumulation data
* **FragAccumOut** – Opaque fragment type for output accumulation data
```

--------------------------------

### ROCWMMA Fragment Scheduler Definitions

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/conceptual/migration-guide.html

Defines fragment scheduler types, including the default non-cooperative scheduler and cooperative schedulers like `coop_row_major_2d`. These schedulers embed scheduling logic and parameters.

```cpp
namespace fragment_scheduler
{
   //! @struct default
   //! @brief The default fragment scheduler; each wave operates independently.
   using default_schedule = IOScheduler::Default;

   //! @struct coop_row_major_2d
   //! @brief  A cooperative scheduling strategy where each wave in the 2d thread block
   //! will contribute to the fragment operation in row_major grid order.
   //! All waves are scheduled in row_major order.
   //! E.g. (TBlockX, TBlockY) => 2x2 waves
   //! w0 = (0, 0),  w1 = (0, 1),
   //! w2 = (1, 0),  w3 = (1, 1)
   //! @tparam TBlockX the size of the thread-block in the X dimension
   //! @tparam TBlockY the size of the thread-block in the Y dimension
   using coop_row_major_2d = IOScheduler::RowMajor2d<TBlockX, TBlockY>;

   ...
}
```

--------------------------------

### rocWMMA API Functions

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Core functions for matrix operations within rocWMMA.

```APIDOC
## rocWMMA API functions

* `fill_fragment()`
* `load_matrix_sync()`
* `load_matrix_sync()`
* `store_matrix_sync()`
* `store_matrix_sync()`
* `mma_sync()`
* `synchronize_workgroup()`
```

--------------------------------

### Synchronous API Functions

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

rocWMMA API functions such as `load_matrix_sync`, `store_matrix_sync`, and `mma_sync` are synchronous when used with global memory. Explicit workgroup synchronization (`synchronize_workgroup`) might be required when using these functions with shared memory.

```APIDOC
## Synchronous API

rocWMMA API functions such as `load_matrix_sync`, `store_matrix_sync`, and `mma_sync` are synchronous when used with global memory. However, when you use these functions with shared memory, for example, LDS memory, explicit workgroup synchronization (`synchronize_workgroup`) might be required.
```

--------------------------------

### rocWMMA Transforms API Functions

Source: https://rocm.docs.amd.com/projects/rocWMMA/en/latest/api-reference/api-reference-guide.html

Functions for transforming matrix data and fragments within rocWMMA.

```APIDOC
## rocWMMA transforms API functions

* `apply_transpose()`
* `apply_data_layout()`
* `apply_fragment()`
* `to_register_file()`
* `from_register_file()`
```
Tokens:
Raw