@@ -3455,6 +3455,7 @@ def NVVM_Tcgen05MmaSmemDescOp : NVVM_Op<"tcgen05.mma_smem_desc", []> {
3455
3455
properties of multiplicand matrix in shared memory including its location
3456
3456
in the shared memory of the current CTA.
3457
3457
3458
+ ```
3458
3459
+-----------+------+------------------------------------------------------+
3459
3460
| Bit-field | Size | Description |
3460
3461
+-----------+------+------------------------------------------------------+
@@ -3477,6 +3478,7 @@ def NVVM_Tcgen05MmaSmemDescOp : NVVM_Op<"tcgen05.mma_smem_desc", []> {
3477
3478
| | | 6: 32-Byte swizzling |
3478
3479
| | | (Values 3, 5 and 7 are invalid) |
3479
3480
+-----------+------+------------------------------------------------------+
3481
+ ```
3480
3482
3481
3483
Example:
3482
3484
```mlir
@@ -3578,7 +3580,8 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMa<[100, 101]>]> {
3578
3580
elements from adjacent columns into a single 32-bit element during the load.
3579
3581
3580
3582
The following table describes the size of the vector for various combinations
3581
- of `num` and `shape` attributes
3583
+ of `num` and `shape` attributes:
3584
+ ```
3582
3585
|=====================================================================|
3583
3586
| num/shape | 16x32bx2/16x64b/32x32b | 16x128b | 16x256b |
3584
3587
|=====================================================================|
@@ -3591,6 +3594,7 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMa<[100, 101]>]> {
3591
3594
| x64 | 64 | 128 | NA |
3592
3595
| x128 | 128 | NA | NA |
3593
3596
|=====================================================================|
3597
+ ```
3594
3598
3595
3599
Example:
3596
3600
```mlir
@@ -3666,7 +3670,8 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMa<[100, 101]>]> {
3666
3670
in the register into two 16-bit elements and store them in adjacent columns.
3667
3671
3668
3672
The following table describes the size of the vector for various combinations
3669
- of `num` and `shape` attributes
3673
+ of `num` and `shape` attributes:
3674
+ ```
3670
3675
|=====================================================================|
3671
3676
| num/shape | 16x32bx2/16x64b/32x32b | 16x128b | 16x256b |
3672
3677
|=====================================================================|
@@ -3679,6 +3684,7 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMa<[100, 101]>]> {
3679
3684
| x64 | 64 | 128 | NA |
3680
3685
| x128 | 128 | NA | NA |
3681
3686
|=====================================================================|
3687
+ ```
3682
3688
3683
3689
Example:
3684
3690
```mlir
0 commit comments