You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/04/01 02:18:15 UTC
[GitHub] [tvm] yzh119 opened a new pull request #10855: [PTX] `ldmatrix` builtin to accelerate copying data from shared memory to warp memory
yzh119 opened a new pull request #10855:
URL: https://github.com/apache/tvm/pull/10855
We already have PTX mma and mma.sp builtin support in #9909 and #10339 . However, we have not supported corresponding data movement builtins for these mma instructions, so the data movement would not be as fast as wmma.
This PR brings the `ldmatrix` builtin, which is a native PTX warp-level instruction (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-ldmatrix), and we can use it to load several (1/2/4) 8x8 matrices from shared memory to warp memory.
@vinx13 @Hzfengsy
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] yzh119 commented on a change in pull request #10855: [PTX] `ldmatrix` builtin to accelerate copying data from shared memory to warp memory
Posted by GitBox <gi...@apache.org>.
yzh119 commented on a change in pull request #10855:
URL: https://github.com/apache/tvm/pull/10855#discussion_r840181094
##########
File path: tests/python/unittest/test_tir_ptx_ldmatrix.py
##########
@@ -0,0 +1,100 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+import tvm
+from tvm.script import tir as T
+import numpy as np
+import tvm.testing
+
+
+@T.prim_func
+def ptx_ldmatrix(
+ A: T.Buffer[(16, 16), "float16"], B: T.Buffer[(16, 16), "float16"], num: T.int32, trans: T.uint8
+) -> None:
+ T.func_attr({"global_symbol": "default_function", "tir.noalias": True})
+ bx = T.env_thread("blockIdx.x")
+ tx = T.env_thread("threadIdx.x")
+ T.launch_thread(bx, 1)
+ T.launch_thread(tx, 32)
+ A_shared = T.allocate([16, 16], "float16", scope="shared")
+ A_local = T.allocate([8], "float16", scope="local")
+ for i in range(8):
+ A_shared[i * 2 + tx // 16, tx % 16] = A[i * 2 + tx // 16, tx % 16]
+ for i in range(8):
+ A_local[i] = 0
Review comment:
@vinx13 I found that if I remove these two lines, the compiler would mistakenly ignore the effect of `ptx_ldmatrix`, even if I mark the built-in as opaque. Would you mind help checking it?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao1994 merged pull request #10855: [PTX] `ldmatrix` builtin to accelerate copying data from shared memory to warp memory
Posted by GitBox <gi...@apache.org>.
junrushao1994 merged pull request #10855:
URL: https://github.com/apache/tvm/pull/10855
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org