[PATCH] D107046: [NVPTX] Add NVPTX intrinsics for CUDA PTX 6.5 ldmatrix instructions

[Steffen Larsen via Phabricator via llvm-commits]

steffenlarsen marked 2 inline comments as done and an inline comment as not done.
steffenlarsen added inline comments.


================
Comment at: llvm/lib/Target/NVPTX/NVPTXIntrinsics.td:7919
+//
+// ldmatrix.sync.aligned.m8n8[|.trans][|.shared].b16
+//
----------------
tra wrote:
> Nit: `|` is redundant here, IMO.
> I think `[something]` already reads as `something is optional.`
A similar syntax is used for `wmma.load` which is specified as

```
//
// wmma.load.[a|b|c].sync.[row|col].m16n16k16[|.global|.shared].[f16|f32]
//
```

a bit further up.

For this case there wouldn't be any ambiguities, but I think it's better to keep it consistent.


================
Comment at: llvm/lib/Target/NVPTX/NVPTXIntrinsics.td:7954-7955
+      } // space
+    } // transposed
+} // defset
 
----------------
tra wrote:
> Nit: something is off with the indentation here.
Good eye. I've fixed it.


================
Comment at: llvm/test/CodeGen/NVPTX/wmma.py:636
+"""
+  intrinsic_template = "llvm.nvvm.ldmatrix.sync.aligned.${geom}.${x}${trans}.${itype}.${pspace}"
+  instruction_template = "ldmatrix.sync.aligned.${geom}.${x}${trans}${space}.${itype}"
----------------
tra wrote:
> `x` -> `frag` ?
Yeah, `frag` might make the mental mapping a bit simpler. 😄 


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D107046/new/

https://reviews.llvm.org/D107046