@@ -3455,6 +3455,7 @@ def NVVM_Tcgen05MmaSmemDescOp : NVVM_Op<"tcgen05.mma_smem_desc", []> {
34553455 properties of multiplicand matrix in shared memory including its location
34563456 in the shared memory of the current CTA.
34573457
3458+ ```
34583459 +-----------+------+------------------------------------------------------+
34593460 | Bit-field | Size | Description |
34603461 +-----------+------+------------------------------------------------------+
@@ -3477,6 +3478,7 @@ def NVVM_Tcgen05MmaSmemDescOp : NVVM_Op<"tcgen05.mma_smem_desc", []> {
34773478 | | | 6: 32-Byte swizzling |
34783479 | | | (Values 3, 5 and 7 are invalid) |
34793480 +-----------+------+------------------------------------------------------+
3481+ ```
34803482
34813483 Example:
34823484 ```mlir
@@ -3578,7 +3580,8 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMa<[100, 101]>]> {
35783580 elements from adjacent columns into a single 32-bit element during the load.
35793581
35803582 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+ ```
35823585 |=====================================================================|
35833586 | num/shape | 16x32bx2/16x64b/32x32b | 16x128b | 16x256b |
35843587 |=====================================================================|
@@ -3591,6 +3594,7 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMa<[100, 101]>]> {
35913594 | x64 | 64 | 128 | NA |
35923595 | x128 | 128 | NA | NA |
35933596 |=====================================================================|
3597+ ```
35943598
35953599 Example:
35963600 ```mlir
@@ -3666,7 +3670,8 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMa<[100, 101]>]> {
36663670 in the register into two 16-bit elements and store them in adjacent columns.
36673671
36683672 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+ ```
36703675 |=====================================================================|
36713676 | num/shape | 16x32bx2/16x64b/32x32b | 16x128b | 16x256b |
36723677 |=====================================================================|
@@ -3679,6 +3684,7 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMa<[100, 101]>]> {
36793684 | x64 | 64 | 128 | NA |
36803685 | x128 | 128 | NA | NA |
36813686 |=====================================================================|
3687+ ```
36823688
36833689 Example:
36843690 ```mlir
0 commit comments