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/08/30 17:30:39 UTC
[GitHub] [tvm] nverke opened a new pull request, #12654: [Hexagon] Create test examples to show parallelization
nverke opened a new pull request, #12654:
URL: https://github.com/apache/tvm/pull/12654
### Background:
These tests show one way to attain speedup for hexagon workloads through parallelism. Additionally, they show the speedup attained by various different workloads in the best configuration found (outer loop split into 4). This will hopefully give users an understanding of what to do and what to expect when parallelizing hexagon workloads.
Here are some tables of the runtimes on various workloads and HDKs.
#### HVX runtimes on 8gen1 HDK
<google-sheets-html-origin><meta name="sheets-original-selection" data-sheets-range="G27:K51"><meta name="sheets-banding" data-sheets-range="G27:K51" data-sheets-banding-properties="{"1":{"1":{"4":{"1":2,"2":12434877}},"2":{"4":{"1":2,"2":16777215}},"3":1,"4":{"4":{"1":2,"2":15987699}},"5":1}}"><style type="text/css"><!--td {border: 1px solid #ccc;}br {mso-data-placement:same-cell;}--></style>
Operation | # Operations | Single Thread (ms) | Parallel (ms) | Speedup
-- | -- | -- | -- | --
vrmpy | 128 | 0.00112 | 0.002242 | 0.5x
vrmpy | 256 | 0.001797 | 0.002505 | 0.72x
vrmpy | 512 | 0.003095 | 0.00309 | 1.0x
vrmpy | 1024 | 0.005732 | 0.00422 | 1.36x
vrmpy | 2048 | 0.011118 | 0.006515 | 1.71x
vrmpy | 4096 | 0.140888 | 0.056102 | 2.51x
vrmpy | 8192 | 0.265677 | 0.100795 | 2.64x
vrmpy | 16384 | 0.877999 | 0.286085 | 3.07x
vmpy | 128 | 0.001371 | 0.002348 | 0.58x
vmpy | 256 | 0.002236 | 0.002721 | 0.82x
vmpy | 512 | 0.003988 | 0.003511 | 1.14x
vmpy | 1024 | 0.007496 | 0.00504 | 1.49x
vmpy | 2048 | 0.018755 | 0.012325 | 1.52x
vmpy | 4096 | 0.172784 | 0.066596 | 2.59x
vmpy | 8192 | 0.357293 | 0.12787 | 2.79x
vmpy | 16384 | 1.223959 | 0.367546 | 3.33x
vadd | 128 | 0.001349 | 0.002355 | 0.57x
vadd | 256 | 0.002194 | 0.002685 | 0.82x
vadd | 512 | 0.003968 | 0.003543 | 1.12x
vadd | 1024 | 0.007491 | 0.005024 | 1.49x
vadd | 2048 | 0.018481 | 0.012329 | 1.5x
vadd | 4096 | 0.172362 | 0.067368 | 2.56x
vadd | 8192 | 0.353322 | 0.130838 | 2.7x
vadd | 16384 | 1.215925 | 0.368648 | 3.3x
#### HVX runtimes on 888 HDK
<google-sheets-html-origin><meta name="sheets-original-selection" data-sheets-range="G1:K25"><meta name="sheets-banding" data-sheets-range="G1:K25" data-sheets-banding-properties="{"1":{"1":{"4":{"1":2,"2":12434877}},"2":{"4":{"1":2,"2":16777215}},"3":1,"4":{"4":{"1":2,"2":15987699}},"5":1}}"><style type="text/css"><!--td {border: 1px solid #ccc;}br {mso-data-placement:same-cell;}--></style>
Operation | # Operations | Single Thread (ms) | Parallel (ms) | Speedup
-- | -- | -- | -- | --
vrmpy | 128 | 0.001157 | 0.002261 | 0.51x
vrmpy | 256 | 0.00188 | 0.003152 | 0.6x
vrmpy | 512 | 0.003201 | 0.00379 | 0.84x
vrmpy | 1024 | 0.005942 | 0.005273 | 1.13x
vrmpy | 2048 | 0.011579 | 0.009932 | 1.17x
vrmpy | 4096 | 0.258623 | 0.083141 | 3.11x
vrmpy | 8192 | 0.509286 | 0.172014 | 2.96x
vrmpy | 16384 | 0.979687 | 0.333655 | 2.94x
vmpy | 128 | 0.00142 | 0.002941 | 0.48x
vmpy | 256 | 0.002289 | 0.003374 | 0.68x
vmpy | 512 | 0.004115 | 0.004397 | 0.94x
vmpy | 1024 | 0.007751 | 0.006372 | 1.22x
vmpy | 2048 | 0.025335 | 0.01625 | 1.56x
vmpy | 4096 | 0.355102 | 0.116926 | 3.04x
vmpy | 8192 | 0.62289 | 0.214158 | 2.91x
vmpy | 16384 | 1.239413 | 0.40535 | 3.06x
vadd | 128 | 0.00141 | 0.002982 | 0.47x
vadd | 256 | 0.002307 | 0.003473 | 0.66x
vadd | 512 | 0.004115 | 0.004371 | 0.94x
vadd | 1024 | 0.007767 | 0.006365 | 1.22x
vadd | 2048 | 0.025241 | 0.01564 | 1.61x
vadd | 4096 | 0.357479 | 0.11579 | 3.09x
vadd | 8192 | 0.646734 | 0.218949 | 2.95x
vadd | 16384 | 1.243888 | 0.400593 | 3.11x
#### Scalar runtimes on 8gen1 HDK
<google-sheets-html-origin><meta name="sheets-original-selection" data-sheets-range="A27:E51"><meta name="sheets-banding" data-sheets-range="A27:E51" data-sheets-banding-properties="{"1":{"1":{"4":{"1":2,"2":12434877}},"2":{"4":{"1":2,"2":16777215}},"3":1,"4":{"4":{"1":2,"2":15987699}},"5":1}}"><style type="text/css"><!--td {border: 1px solid #ccc;}br {mso-data-placement:same-cell;}--></style>
Operation | # Operations | Single Thread (ms) | Parallel (ms) | Speedup
-- | -- | -- | -- | --
add | 128 | 0.000961 | 0.002217 | 0.43x
add | 256 | 0.001477 | 0.00277 | 0.53x
add | 512 | 0.00255 | 0.003356 | 0.76x
add | 1024 | 0.005959 | 0.005371 | 1.11x
add | 2048 | 0.015845 | 0.008645 | 1.83x
add | 4096 | 0.030957 | 0.014782 | 2.09x
add | 8192 | 0.058772 | 0.027117 | 2.17x
add | 16384 | 0.115732 | 0.050325 | 2.3x
multiply | 128 | 0.00144 | 0.002561 | 0.56x
multiply | 256 | 0.002457 | 0.003177 | 0.77x
multiply | 512 | 0.004462 | 0.004707 | 0.95x
multiply | 1024 | 0.009583 | 0.007716 | 1.24x
multiply | 2048 | 0.022285 | 0.013335 | 1.67x
multiply | 4096 | 0.043851 | 0.023989 | 1.83x
multiply | 8192 | 0.085892 | 0.047104 | 1.82x
multiply | 16384 | 0.152309 | 0.086548 | 1.76x
sub | 128 | 0.000981 | 0.002077 | 0.47x
sub | 256 | 0.001482 | 0.002512 | 0.59x
sub | 512 | 0.002549 | 0.003345 | 0.76x
sub | 1024 | 0.006042 | 0.005565 | 1.09x
sub | 2048 | 0.015846 | 0.008672 | 1.83x
sub | 4096 | 0.030947 | 0.014812 | 2.09x
sub | 8192 | 0.059008 | 0.027133 | 2.17x
sub | 16384 | 0.115776 | 0.049741 | 2.33x
#### Scalar runtimes on 888 HDK
<google-sheets-html-origin><meta name="sheets-original-selection" data-sheets-range="A1:E25"><meta name="sheets-banding" data-sheets-range="A1:E25" data-sheets-banding-properties="{"1":{"1":{"4":{"1":2,"2":12434877}},"2":{"4":{"1":2,"2":16777215}},"3":1,"4":{"4":{"1":2,"2":15987699}},"5":1}}"><style type="text/css"><!--td {border: 1px solid #ccc;}br {mso-data-placement:same-cell;}--></style>
Operation | # Operations | Single Thread (ms) | Parallel (ms) | Speedup
-- | -- | -- | -- | --
add | 128 | 0.001019 | 0.002461 | 0.41x
add | 256 | 0.001564 | 0.002714 | 0.58x
add | 512 | 0.002643 | 0.003411 | 0.77x
add | 1024 | 0.006376 | 0.005649 | 1.13x
add | 2048 | 0.016413 | 0.009117 | 1.8x
add | 4096 | 0.032022 | 0.015408 | 2.08x
add | 8192 | 0.061622 | 0.028349 | 2.17x
add | 16384 | 0.114904 | 0.051258 | 2.24x
multiply | 128 | 0.001515 | 0.002529 | 0.6x
multiply | 256 | 0.002533 | 0.003264 | 0.78x
multiply | 512 | 0.00461 | 0.004697 | 0.98x
multiply | 1024 | 0.010024 | 0.008115 | 1.24x
multiply | 2048 | 0.02271 | 0.016956 | 1.34x
multiply | 4096 | 0.044663 | 0.025626 | 1.74x
multiply | 8192 | 0.08593 | 0.050227 | 1.71x
multiply | 16384 | 0.156369 | 0.082328 | 1.9x
sub | 128 | 0.001005 | 0.002229 | 0.45x
sub | 256 | 0.001567 | 0.002598 | 0.6x
sub | 512 | 0.002646 | 0.003399 | 0.78x
sub | 1024 | 0.006413 | 0.006804 | 0.94x
sub | 2048 | 0.016408 | 0.011015 | 1.49x
sub | 4096 | 0.032029 | 0.017999 | 1.78x
sub | 8192 | 0.061626 | 0.028435 | 2.17x
sub | 16384 | 0.118424 | 0.052493 | 2.26x
--
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] nverke commented on pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
nverke commented on PR #12654:
URL: https://github.com/apache/tvm/pull/12654#issuecomment-1244063443
@tvm-bot rerun
--
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] nverke commented on a diff in pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
nverke commented on code in PR #12654:
URL: https://github.com/apache/tvm/pull/12654#discussion_r968659243
##########
tests/python/contrib/test_hexagon/test_parallel_hvx.py:
##########
@@ -0,0 +1,255 @@
+# 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.
+
+"""
+Test parallelizing HVX workloads and compare them to single thread examples.
+"""
+import numpy as np
+import tvm
+
+from tvm.script import tir as T
+from numpy.random import default_rng
+
+TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n"
+
+
+def get_vrmpy_shape_dtypes(operations):
+ return ((operations, 128), "uint8", (operations, 128), "uint8", (operations, 32), "int32")
+
+
+def get_vmpy_vadd_shape_dtype(operations):
+ return ((operations, 128), "uint8", (operations, 128), "uint8", (operations, 128), "int16")
+
+
+def vmpy_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int16")
+ for n in range(shape[0]):
+ for i in range(0, 128, 2):
+ expected[n, i // 2] = np.int16(a[n, i]) * np.int16(b[n, i])
+ for i in range(1, 128, 2):
+ expected[n, i // 2 + 64] = np.int16(a[n, i]) * np.int16(b[n, i])
+ return expected
+
+
+def vadd_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int16")
+ for n in range(shape[0]):
+ for i in range(0, 128, 2):
+ expected[n, i // 2] = np.int16(a[n, i]) + np.int16(b[n, i])
+ for i in range(1, 128, 2):
+ expected[n, i // 2 + 64] = np.int16(a[n, i]) + np.int16(b[n, i])
+ return expected
+
+
+def vrmpy_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int32")
+ for n in range(shape[0]):
+ for i in range(32):
+ for r in range(4):
+ expected[n, i] = expected[n, i] + np.uint32(a[n, i * 4 + r]) * np.uint32(
+ b[n, i * 4 + r]
+ )
+ return expected
+
+
+def get_vmpy_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 128], dtype="int16")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 128)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vmpybusv.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int16x128",
+ )
+
+ return operator
+
+
+def get_vadd_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 128], dtype="int16")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 128)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vaddubh.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int16x128",
+ )
+
+ return operator
+
+
+def get_vrmpy_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 32], dtype="int32")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 32)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vrmpyubv.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int32x32",
+ )
+
+ return operator
+
+
+def evaluate(hexagon_session, shape_dtypes, expected_producer, sch):
+ a_shape, a_dtype, b_shape, b_dtype, c_shape, c_dtype = shape_dtypes
+
+ target_hexagon = tvm.target.hexagon("v68")
+ func_tir = tvm.build(
+ sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon)
+ )
+ module = hexagon_session.load_module(func_tir)
+
+ rng = default_rng()
+ a = rng.integers(0, 16, a_shape, dtype=a_dtype)
+ b = rng.integers(0, 16, b_shape, dtype=b_dtype)
+ c = np.zeros(c_shape, dtype=c_dtype)
+
+ a_hexagon = tvm.runtime.ndarray.array(a, device=hexagon_session.device)
+ b_hexagon = tvm.runtime.ndarray.array(b, device=hexagon_session.device)
+ c_hexagon = tvm.runtime.ndarray.array(c, device=hexagon_session.device)
+
+ timer = module.time_evaluator("__tvm_main__", hexagon_session.device, number=100, repeat=10)
+ runtime = timer(a_hexagon, b_hexagon, c_hexagon)
+ tvm.testing.assert_allclose(c_hexagon.asnumpy(), expected_producer(c_shape, a, b))
+
+ return round(runtime.mean * 1000, 6)
+
+
+class TestMatMulVec:
+
+ operations = tvm.testing.parameter(
+ 128,
+ 256,
+ 512,
+ 1024, # Single thread runs faster since L2 cache can handle the entire request quickly
+ 2048,
+ 4096, # Significant performance degredation once the inputs and outputs cannot all fit in L2
Review Comment:
I added a little discussion to the description
--
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] nverke commented on a diff in pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
nverke commented on code in PR #12654:
URL: https://github.com/apache/tvm/pull/12654#discussion_r968913330
##########
tests/python/contrib/test_hexagon/test_parallel_hvx.py:
##########
@@ -0,0 +1,255 @@
+# 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.
+
+"""
+Test parallelizing HVX workloads and compare them to single thread examples.
+"""
+import numpy as np
+import tvm
+
+from tvm.script import tir as T
+from numpy.random import default_rng
+
+TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n"
+
+
+def get_vrmpy_shape_dtypes(operations):
+ return ((operations, 128), "uint8", (operations, 128), "uint8", (operations, 32), "int32")
+
+
+def get_vmpy_vadd_shape_dtype(operations):
+ return ((operations, 128), "uint8", (operations, 128), "uint8", (operations, 128), "int16")
+
+
+def vmpy_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int16")
+ for n in range(shape[0]):
+ for i in range(0, 128, 2):
+ expected[n, i // 2] = np.int16(a[n, i]) * np.int16(b[n, i])
+ for i in range(1, 128, 2):
+ expected[n, i // 2 + 64] = np.int16(a[n, i]) * np.int16(b[n, i])
+ return expected
+
+
+def vadd_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int16")
+ for n in range(shape[0]):
+ for i in range(0, 128, 2):
+ expected[n, i // 2] = np.int16(a[n, i]) + np.int16(b[n, i])
+ for i in range(1, 128, 2):
+ expected[n, i // 2 + 64] = np.int16(a[n, i]) + np.int16(b[n, i])
+ return expected
+
+
+def vrmpy_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int32")
+ for n in range(shape[0]):
+ for i in range(32):
+ for r in range(4):
+ expected[n, i] = expected[n, i] + np.uint32(a[n, i * 4 + r]) * np.uint32(
+ b[n, i * 4 + r]
+ )
+ return expected
+
+
+def get_vmpy_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 128], dtype="int16")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 128)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vmpybusv.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int16x128",
+ )
+
+ return operator
+
+
+def get_vadd_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 128], dtype="int16")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 128)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vaddubh.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int16x128",
+ )
+
+ return operator
+
+
+def get_vrmpy_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 32], dtype="int32")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 32)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vrmpyubv.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int32x32",
+ )
+
+ return operator
+
+
+def evaluate(hexagon_session, shape_dtypes, expected_producer, sch):
+ a_shape, a_dtype, b_shape, b_dtype, c_shape, c_dtype = shape_dtypes
+
+ target_hexagon = tvm.target.hexagon("v68")
+ func_tir = tvm.build(
+ sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon)
+ )
+ module = hexagon_session.load_module(func_tir)
+
+ rng = default_rng()
+ a = rng.integers(0, 16, a_shape, dtype=a_dtype)
+ b = rng.integers(0, 16, b_shape, dtype=b_dtype)
+ c = np.zeros(c_shape, dtype=c_dtype)
+
+ a_hexagon = tvm.runtime.ndarray.array(a, device=hexagon_session.device)
+ b_hexagon = tvm.runtime.ndarray.array(b, device=hexagon_session.device)
+ c_hexagon = tvm.runtime.ndarray.array(c, device=hexagon_session.device)
+
+ timer = module.time_evaluator("__tvm_main__", hexagon_session.device, number=100, repeat=10)
+ runtime = timer(a_hexagon, b_hexagon, c_hexagon)
+ tvm.testing.assert_allclose(c_hexagon.asnumpy(), expected_producer(c_shape, a, b))
+
+ return round(runtime.mean * 1000, 6)
+
+
+class TestMatMulVec:
+
+ operations = tvm.testing.parameter(
+ 128,
+ 256,
+ 512,
+ 1024, # Single thread runs faster since L2 cache can handle the entire request quickly
+ 2048,
+ 4096, # Significant performance degredation once the inputs and outputs cannot all fit in L2
+ 8192,
+ 16384,
+ )
+
+ # Experimentally best split factor but all multiples of 4 perform pretty well.
+ # This is because there are 4 HVX untis available on the device and pipelining
+ # works best with parallels of the number of available HVX.
+ split_factor = tvm.testing.parameter(4)
Review Comment:
Good suggestion! Ill see if I can do this!
--
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] nverke commented on a diff in pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
nverke commented on code in PR #12654:
URL: https://github.com/apache/tvm/pull/12654#discussion_r968656484
##########
tests/python/contrib/test_hexagon/test_parallel_scalar.py:
##########
@@ -0,0 +1,178 @@
+# 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.
+
+""" Test parallelism for multiple different scalar workloads. """
+
+import numpy as np
+import tvm
+
+from tvm.script import tir as T
+from numpy.random import default_rng
+
+TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n"
+
+
+def get_add_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] + B[vn]
+
+ return operator
+
+
+def get_multiply_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] * B[vn]
+
+ return operator
+
+
+def get_sub_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] - B[vn]
+
+ return operator
+
+
+def evaluate(hexagon_session, operations, expected, sch):
+ shape = operations
+ dtype = "float64"
Review Comment:
Yeah, thats the idea, I know that hvx does not support float64
--
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] nverke commented on pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
nverke commented on PR #12654:
URL: https://github.com/apache/tvm/pull/12654#issuecomment-1245697544
@adstraw I believe I made the code more clear as to which are producers but let me know what you think!
--
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] nverke commented on pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
nverke commented on PR #12654:
URL: https://github.com/apache/tvm/pull/12654#issuecomment-1244044726
@tvm-bot rerun
--
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] tmoreau89 commented on pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on PR #12654:
URL: https://github.com/apache/tvm/pull/12654#issuecomment-1246773027
@nverke PR is good to go minus merge conflicts that require resolution.
--
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] tmoreau89 merged pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
tmoreau89 merged PR #12654:
URL: https://github.com/apache/tvm/pull/12654
--
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] tmoreau89 commented on pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on PR #12654:
URL: https://github.com/apache/tvm/pull/12654#issuecomment-1251296243
Thanks @nverke and @adstraw - the PR has been merged
--
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] adstraw commented on a diff in pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
adstraw commented on code in PR #12654:
URL: https://github.com/apache/tvm/pull/12654#discussion_r968897620
##########
tests/python/contrib/test_hexagon/test_parallel_scalar.py:
##########
@@ -0,0 +1,178 @@
+# 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.
+
+""" Test parallelism for multiple different scalar workloads. """
+
+import numpy as np
+import tvm
+
+from tvm.script import tir as T
+from numpy.random import default_rng
+
+TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n"
+
+
+def get_add_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] + B[vn]
+
+ return operator
+
+
+def get_multiply_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] * B[vn]
+
+ return operator
+
+
+def get_sub_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] - B[vn]
+
+ return operator
+
+
+def evaluate(hexagon_session, operations, expected, sch):
+ shape = operations
+ dtype = "float64"
+
+ target_hexagon = tvm.target.hexagon("v68")
+ func_tir = tvm.build(
+ sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon)
+ )
+ module = hexagon_session.load_module(func_tir)
+
+ rng = default_rng()
+ a = rng.random(shape, dtype=dtype)
+ b = rng.random(shape, dtype=dtype)
+ c = np.zeros(shape, dtype=dtype)
+
+ a_hexagon = tvm.runtime.ndarray.array(a, device=hexagon_session.device)
+ b_hexagon = tvm.runtime.ndarray.array(b, device=hexagon_session.device)
+ c_hexagon = tvm.runtime.ndarray.array(c, device=hexagon_session.device)
+
+ timer = module.time_evaluator("__tvm_main__", hexagon_session.device, number=100, repeat=10)
+ runtime = timer(a_hexagon, b_hexagon, c_hexagon)
+
+ tvm.testing.assert_allclose(c_hexagon.asnumpy(), expected(a, b))
+
+ return round(runtime.mean * 1000, 6)
+
+
+class TestMatMulVec:
+
+ operations = tvm.testing.parameter(
+ 128,
+ 256,
+ 512,
+ 1024, # Single thread runs faster since L2 cache can handle the entire request quickly
+ 2048,
+ 4096, # Significant performance degredation once the inputs and outputs cannot all fit in L2
+ 8192,
+ 16384,
+ )
+
+ split_factor = tvm.testing.parameter(4)
Review Comment:
Same comment; can we parameterize the op type to avoid code duplication?
##########
tests/python/contrib/test_hexagon/test_parallel_hvx.py:
##########
@@ -0,0 +1,255 @@
+# 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.
+
+"""
+Test parallelizing HVX workloads and compare them to single thread examples.
+"""
+import numpy as np
+import tvm
+
+from tvm.script import tir as T
+from numpy.random import default_rng
+
+TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n"
+
+
+def get_vrmpy_shape_dtypes(operations):
+ return ((operations, 128), "uint8", (operations, 128), "uint8", (operations, 32), "int32")
+
+
+def get_vmpy_vadd_shape_dtype(operations):
+ return ((operations, 128), "uint8", (operations, 128), "uint8", (operations, 128), "int16")
+
+
+def vmpy_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int16")
+ for n in range(shape[0]):
+ for i in range(0, 128, 2):
+ expected[n, i // 2] = np.int16(a[n, i]) * np.int16(b[n, i])
+ for i in range(1, 128, 2):
+ expected[n, i // 2 + 64] = np.int16(a[n, i]) * np.int16(b[n, i])
+ return expected
+
+
+def vadd_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int16")
+ for n in range(shape[0]):
+ for i in range(0, 128, 2):
+ expected[n, i // 2] = np.int16(a[n, i]) + np.int16(b[n, i])
+ for i in range(1, 128, 2):
+ expected[n, i // 2 + 64] = np.int16(a[n, i]) + np.int16(b[n, i])
+ return expected
+
+
+def vrmpy_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int32")
+ for n in range(shape[0]):
+ for i in range(32):
+ for r in range(4):
+ expected[n, i] = expected[n, i] + np.uint32(a[n, i * 4 + r]) * np.uint32(
+ b[n, i * 4 + r]
+ )
+ return expected
+
+
+def get_vmpy_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 128], dtype="int16")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 128)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vmpybusv.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int16x128",
+ )
+
+ return operator
+
+
+def get_vadd_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 128], dtype="int16")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 128)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vaddubh.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int16x128",
+ )
+
+ return operator
+
+
+def get_vrmpy_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 32], dtype="int32")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 32)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vrmpyubv.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int32x32",
+ )
+
+ return operator
+
+
+def evaluate(hexagon_session, shape_dtypes, expected_producer, sch):
+ a_shape, a_dtype, b_shape, b_dtype, c_shape, c_dtype = shape_dtypes
+
+ target_hexagon = tvm.target.hexagon("v68")
+ func_tir = tvm.build(
+ sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon)
+ )
+ module = hexagon_session.load_module(func_tir)
+
+ rng = default_rng()
+ a = rng.integers(0, 16, a_shape, dtype=a_dtype)
+ b = rng.integers(0, 16, b_shape, dtype=b_dtype)
+ c = np.zeros(c_shape, dtype=c_dtype)
+
+ a_hexagon = tvm.runtime.ndarray.array(a, device=hexagon_session.device)
+ b_hexagon = tvm.runtime.ndarray.array(b, device=hexagon_session.device)
+ c_hexagon = tvm.runtime.ndarray.array(c, device=hexagon_session.device)
+
+ timer = module.time_evaluator("__tvm_main__", hexagon_session.device, number=100, repeat=10)
+ runtime = timer(a_hexagon, b_hexagon, c_hexagon)
+ tvm.testing.assert_allclose(c_hexagon.asnumpy(), expected_producer(c_shape, a, b))
+
+ return round(runtime.mean * 1000, 6)
+
+
+class TestMatMulVec:
+
+ operations = tvm.testing.parameter(
+ 128,
+ 256,
+ 512,
+ 1024, # Single thread runs faster since L2 cache can handle the entire request quickly
+ 2048,
+ 4096, # Significant performance degredation once the inputs and outputs cannot all fit in L2
+ 8192,
+ 16384,
+ )
+
+ # Experimentally best split factor but all multiples of 4 perform pretty well.
+ # This is because there are 4 HVX untis available on the device and pipelining
+ # works best with parallels of the number of available HVX.
+ split_factor = tvm.testing.parameter(4)
Review Comment:
Can you find a way to parameterize the operation type (vrmpy, vmpy, vadd) and create just one test case? There is a lot of duplicated code in the 3 test cases below?
--
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] nverke commented on a diff in pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
nverke commented on code in PR #12654:
URL: https://github.com/apache/tvm/pull/12654#discussion_r969881212
##########
tests/python/contrib/test_hexagon/test_parallel_scalar.py:
##########
@@ -0,0 +1,178 @@
+# 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.
+
+""" Test parallelism for multiple different scalar workloads. """
+
+import numpy as np
+import tvm
+
+from tvm.script import tir as T
+from numpy.random import default_rng
+
+TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n"
+
+
+def get_add_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] + B[vn]
+
+ return operator
+
+
+def get_multiply_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] * B[vn]
+
+ return operator
+
+
+def get_sub_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] - B[vn]
+
+ return operator
+
+
+def evaluate(hexagon_session, operations, expected, sch):
+ shape = operations
+ dtype = "float64"
+
+ target_hexagon = tvm.target.hexagon("v68")
+ func_tir = tvm.build(
+ sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon)
+ )
+ module = hexagon_session.load_module(func_tir)
+
+ rng = default_rng()
+ a = rng.random(shape, dtype=dtype)
+ b = rng.random(shape, dtype=dtype)
+ c = np.zeros(shape, dtype=dtype)
+
+ a_hexagon = tvm.runtime.ndarray.array(a, device=hexagon_session.device)
+ b_hexagon = tvm.runtime.ndarray.array(b, device=hexagon_session.device)
+ c_hexagon = tvm.runtime.ndarray.array(c, device=hexagon_session.device)
+
+ timer = module.time_evaluator("__tvm_main__", hexagon_session.device, number=100, repeat=10)
+ runtime = timer(a_hexagon, b_hexagon, c_hexagon)
+
+ tvm.testing.assert_allclose(c_hexagon.asnumpy(), expected(a, b))
+
+ return round(runtime.mean * 1000, 6)
+
+
+class TestMatMulVec:
+
+ operations = tvm.testing.parameter(
+ 128,
+ 256,
+ 512,
+ 1024, # Single thread runs faster since L2 cache can handle the entire request quickly
+ 2048,
+ 4096, # Significant performance degredation once the inputs and outputs cannot all fit in L2
+ 8192,
+ 16384,
+ )
+
+ split_factor = tvm.testing.parameter(4)
Review Comment:
Made the changes
##########
tests/python/contrib/test_hexagon/test_parallel_hvx.py:
##########
@@ -0,0 +1,255 @@
+# 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.
+
+"""
+Test parallelizing HVX workloads and compare them to single thread examples.
+"""
+import numpy as np
+import tvm
+
+from tvm.script import tir as T
+from numpy.random import default_rng
+
+TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n"
+
+
+def get_vrmpy_shape_dtypes(operations):
+ return ((operations, 128), "uint8", (operations, 128), "uint8", (operations, 32), "int32")
+
+
+def get_vmpy_vadd_shape_dtype(operations):
+ return ((operations, 128), "uint8", (operations, 128), "uint8", (operations, 128), "int16")
+
+
+def vmpy_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int16")
+ for n in range(shape[0]):
+ for i in range(0, 128, 2):
+ expected[n, i // 2] = np.int16(a[n, i]) * np.int16(b[n, i])
+ for i in range(1, 128, 2):
+ expected[n, i // 2 + 64] = np.int16(a[n, i]) * np.int16(b[n, i])
+ return expected
+
+
+def vadd_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int16")
+ for n in range(shape[0]):
+ for i in range(0, 128, 2):
+ expected[n, i // 2] = np.int16(a[n, i]) + np.int16(b[n, i])
+ for i in range(1, 128, 2):
+ expected[n, i // 2 + 64] = np.int16(a[n, i]) + np.int16(b[n, i])
+ return expected
+
+
+def vrmpy_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int32")
+ for n in range(shape[0]):
+ for i in range(32):
+ for r in range(4):
+ expected[n, i] = expected[n, i] + np.uint32(a[n, i * 4 + r]) * np.uint32(
+ b[n, i * 4 + r]
+ )
+ return expected
+
+
+def get_vmpy_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 128], dtype="int16")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 128)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vmpybusv.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int16x128",
+ )
+
+ return operator
+
+
+def get_vadd_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 128], dtype="int16")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 128)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vaddubh.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int16x128",
+ )
+
+ return operator
+
+
+def get_vrmpy_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 32], dtype="int32")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 32)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vrmpyubv.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int32x32",
+ )
+
+ return operator
+
+
+def evaluate(hexagon_session, shape_dtypes, expected_producer, sch):
+ a_shape, a_dtype, b_shape, b_dtype, c_shape, c_dtype = shape_dtypes
+
+ target_hexagon = tvm.target.hexagon("v68")
+ func_tir = tvm.build(
+ sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon)
+ )
+ module = hexagon_session.load_module(func_tir)
+
+ rng = default_rng()
+ a = rng.integers(0, 16, a_shape, dtype=a_dtype)
+ b = rng.integers(0, 16, b_shape, dtype=b_dtype)
+ c = np.zeros(c_shape, dtype=c_dtype)
+
+ a_hexagon = tvm.runtime.ndarray.array(a, device=hexagon_session.device)
+ b_hexagon = tvm.runtime.ndarray.array(b, device=hexagon_session.device)
+ c_hexagon = tvm.runtime.ndarray.array(c, device=hexagon_session.device)
+
+ timer = module.time_evaluator("__tvm_main__", hexagon_session.device, number=100, repeat=10)
+ runtime = timer(a_hexagon, b_hexagon, c_hexagon)
+ tvm.testing.assert_allclose(c_hexagon.asnumpy(), expected_producer(c_shape, a, b))
+
+ return round(runtime.mean * 1000, 6)
+
+
+class TestMatMulVec:
+
+ operations = tvm.testing.parameter(
+ 128,
+ 256,
+ 512,
+ 1024, # Single thread runs faster since L2 cache can handle the entire request quickly
+ 2048,
+ 4096, # Significant performance degredation once the inputs and outputs cannot all fit in L2
+ 8192,
+ 16384,
+ )
+
+ # Experimentally best split factor but all multiples of 4 perform pretty well.
+ # This is because there are 4 HVX untis available on the device and pipelining
+ # works best with parallels of the number of available HVX.
+ split_factor = tvm.testing.parameter(4)
Review Comment:
Made the changes
--
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] csullivan commented on a diff in pull request #12654: [Hexagon] Create test examples to show parallelization
Posted by GitBox <gi...@apache.org>.
csullivan commented on code in PR #12654:
URL: https://github.com/apache/tvm/pull/12654#discussion_r961852894
##########
tests/python/contrib/test_hexagon/test_parallel_scalar.py:
##########
@@ -0,0 +1,178 @@
+# 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.
+
+""" Test parallelism for multiple different scalar workloads. """
+
+import numpy as np
+import tvm
+
+from tvm.script import tir as T
+from numpy.random import default_rng
+
+TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n"
+
+
+def get_add_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] + B[vn]
+
+ return operator
+
+
+def get_multiply_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] * B[vn]
+
+ return operator
+
+
+def get_sub_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations], dtype="float64")
+ B = T.match_buffer(b, [operations], dtype="float64")
+ C = T.match_buffer(c, [operations], dtype="float64")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn] = A[vn] - B[vn]
+
+ return operator
+
+
+def evaluate(hexagon_session, operations, expected, sch):
+ shape = operations
+ dtype = "float64"
Review Comment:
Why choose `float64` here? Is the goal to try to avoid vectorization?
##########
tests/python/contrib/test_hexagon/test_parallel_hvx.py:
##########
@@ -0,0 +1,255 @@
+# 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.
+
+"""
+Test parallelizing HVX workloads and compare them to single thread examples.
+"""
+import numpy as np
+import tvm
+
+from tvm.script import tir as T
+from numpy.random import default_rng
+
+TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n"
+
+
+def get_vrmpy_shape_dtypes(operations):
+ return ((operations, 128), "uint8", (operations, 128), "uint8", (operations, 32), "int32")
+
+
+def get_vmpy_vadd_shape_dtype(operations):
+ return ((operations, 128), "uint8", (operations, 128), "uint8", (operations, 128), "int16")
+
+
+def vmpy_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int16")
+ for n in range(shape[0]):
+ for i in range(0, 128, 2):
+ expected[n, i // 2] = np.int16(a[n, i]) * np.int16(b[n, i])
+ for i in range(1, 128, 2):
+ expected[n, i // 2 + 64] = np.int16(a[n, i]) * np.int16(b[n, i])
+ return expected
+
+
+def vadd_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int16")
+ for n in range(shape[0]):
+ for i in range(0, 128, 2):
+ expected[n, i // 2] = np.int16(a[n, i]) + np.int16(b[n, i])
+ for i in range(1, 128, 2):
+ expected[n, i // 2 + 64] = np.int16(a[n, i]) + np.int16(b[n, i])
+ return expected
+
+
+def vrmpy_expected_producer(shape, a, b):
+ expected = np.zeros(shape, dtype="int32")
+ for n in range(shape[0]):
+ for i in range(32):
+ for r in range(4):
+ expected[n, i] = expected[n, i] + np.uint32(a[n, i * 4 + r]) * np.uint32(
+ b[n, i * 4 + r]
+ )
+ return expected
+
+
+def get_vmpy_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 128], dtype="int16")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 128)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vmpybusv.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int16x128",
+ )
+
+ return operator
+
+
+def get_vadd_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 128], dtype="int16")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 128)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vaddubh.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int16x128",
+ )
+
+ return operator
+
+
+def get_vrmpy_operator(operations):
+ @T.prim_func
+ def operator(a: T.handle, b: T.handle, c: T.handle) -> None:
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ A = T.match_buffer(a, [operations, 128], dtype="uint8")
+ B = T.match_buffer(b, [operations, 128], dtype="uint8")
+ C = T.match_buffer(c, [operations, 32], dtype="int32")
+ for n in T.grid(operations):
+ with T.block("C"):
+ vn = T.axis.remap("S", [n])
+ C[vn, T.ramp(0, 1, 32)] = T.call_llvm_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vrmpyubv.128B"),
+ T.uint32(2),
+ T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ T.reinterpret(B[vn, T.ramp(0, 1, 128)], dtype="int32x32"),
+ dtype="int32x32",
+ )
+
+ return operator
+
+
+def evaluate(hexagon_session, shape_dtypes, expected_producer, sch):
+ a_shape, a_dtype, b_shape, b_dtype, c_shape, c_dtype = shape_dtypes
+
+ target_hexagon = tvm.target.hexagon("v68")
+ func_tir = tvm.build(
+ sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon)
+ )
+ module = hexagon_session.load_module(func_tir)
+
+ rng = default_rng()
+ a = rng.integers(0, 16, a_shape, dtype=a_dtype)
+ b = rng.integers(0, 16, b_shape, dtype=b_dtype)
+ c = np.zeros(c_shape, dtype=c_dtype)
+
+ a_hexagon = tvm.runtime.ndarray.array(a, device=hexagon_session.device)
+ b_hexagon = tvm.runtime.ndarray.array(b, device=hexagon_session.device)
+ c_hexagon = tvm.runtime.ndarray.array(c, device=hexagon_session.device)
+
+ timer = module.time_evaluator("__tvm_main__", hexagon_session.device, number=100, repeat=10)
+ runtime = timer(a_hexagon, b_hexagon, c_hexagon)
+ tvm.testing.assert_allclose(c_hexagon.asnumpy(), expected_producer(c_shape, a, b))
+
+ return round(runtime.mean * 1000, 6)
+
+
+class TestMatMulVec:
+
+ operations = tvm.testing.parameter(
+ 128,
+ 256,
+ 512,
+ 1024, # Single thread runs faster since L2 cache can handle the entire request quickly
+ 2048,
+ 4096, # Significant performance degredation once the inputs and outputs cannot all fit in L2
Review Comment:
I think these are insightful comments that should go in the PR description so that we can make sure to bring them in to the commit description when this PR lands.
--
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