[PATCH] D38645: [NVPTX] Implemented wmma intrinsics and instructions.

Justin Lebar via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Wed Oct 11 11:35:25 PDT 2017


jlebar added a comment.

(Responding to Yuan, which also dumps my in-progress comments on the code; sorry for the noise.)

> Specifically, yours embeds the layout/memory space/ while ours treats them as constant arguments. We did this to reduce the amount of intrinsic functions the optimizer and codegen have to deal with.

How does having more intrinsic functions cause a problem for the optimizer / codegen?



================
Comment at: llvm/test/CodeGen/NVPTX/wmma.py:16
+def make_wmma_ld_ret_ty(abc, itype):
+  return "{" + ", ".join(make_wmma_slice_ty(abc, itype)) + "}"
+
----------------
I guess it's a nit, but perhaps

   "{%s}" % ", ".join(...)

?   I'd certainly find that more readable.


================
Comment at: llvm/test/CodeGen/NVPTX/wmma.py:21
+check_f16_4 = "{{" + ", *".join(["%hh[0-9]+"]*4) + "}}"
+check_f32_8 = "{{" + ", *".join(["%f[0-9]+"]*8) + "}}"
+
----------------
Nit, space around the asterisk, as with other operators.


================
Comment at: llvm/test/CodeGen/NVPTX/wmma.py:150
+declare ${ret_ty} @llvm.nvvm.wmma.mma.sync.$intrinsic_suffix(
+\t${values});
+
----------------
Is there some reason we're using `\t` rather than indenting with spaces?  That might make this easier to read...


================
Comment at: llvm/test/CodeGen/NVPTX/wmma.py:154
+define ${ret_ty} @test_wmma_mma_${function_suffix}(
+\t${values}) {
+; CHECK wmma.mma.${intrinsic_suffix} {{.*}}[%rd{{[0-9+]}}
----------------
s/values/params/?


================
Comment at: llvm/test/CodeGen/NVPTX/wmma.py:185
+    test_params["intrinsic_suffix"] = Template(suffix_template).substitute(params)
+    test_params["function_suffix"] = test_params["intrinsic_suffix"].replace(".","_")
+    test_params["ret_ty"] = make_wmma_ld_ret_ty("d", dtype)
----------------
Space after `,`


================
Comment at: llvm/test/CodeGen/NVPTX/wmma.py:188
+    # Make list of all values in a,b,c
+    values = ",\n\t".join(make_wmma_slice_args(t, abcd, prefix=abcd)
+                          for abcd, t in (("a", "f16"),("b", "f16"), ("c", ctype)))
----------------
`test_params["values"] = ...`?


================
Comment at: llvm/test/CodeGen/NVPTX/wmma.py:189
+    values = ",\n\t".join(make_wmma_slice_args(t, abcd, prefix=abcd)
+                          for abcd, t in (("a", "f16"),("b", "f16"), ("c", ctype)))
+
----------------
Missing space after `,`


https://reviews.llvm.org/D38645





More information about the llvm-commits mailing list