nverke commented on code in PR #13005: URL: https://github.com/apache/tvm/pull/13005#discussion_r991531651
########## tests/python/contrib/test_hexagon/test_async_dma_pipeline.py: ########## @@ -0,0 +1,336 @@ +# 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 different strategies for loading data into vtcm before running HVX workloads. """ + +import numpy as np +import tvm + +from tvm.script import tir as T +from numpy.random import default_rng + + +def conv_approximation(size_a, size_w): + @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, [size_a, 128], dtype="uint8", align=128) + W = T.match_buffer(b, [size_w, 128], dtype="uint8", align=128) + C = T.match_buffer(c, [size_a, 32], dtype="int32", align=128) + for n, i in T.grid(size_a, size_w): + with T.block("C"): + vn, vi = T.axis.remap("SR", [n, i]) + T.reads(A[vn, 0:128], W[vi, 0:128], C[vn, 0:32]) + T.writes(C[vn, 0:32]) + with T.init(): + for x in T.serial(32): + C[vn, x] = 0 + C[vn, T.ramp(0, 1, 32)] = T.call_llvm_intrin( + T.llvm_lookup_intrinsic_id("llvm.hexagon.V6.vrmpyubv.acc.128B"), + T.uint32(3), + C[vn, T.ramp(0, 1, 32)], + T.reinterpret(A[vn, T.ramp(0, 1, 128)], dtype="int32x32"), + T.reinterpret(W[vi, T.ramp(0, 1, 128)], dtype="int32x32"), + dtype="int32x32", + ) + T.evaluate( + T.tvm_call_packed( + "device_api.hexagon.dma_wait", + 0, # QueueId + 0, # Wait for 0 in flight + dtype="int32", + ) + ) + + return tvm.tir.Schedule(operator) + + +def evaluate(hexagon_session, sch, a, b, size_a, expected_output, use_async_copy=0): + target_hexagon = tvm.target.hexagon("v68", link_params=True) + with tvm.transform.PassContext(config={"tir.use_async_copy": use_async_copy}): + func_tir = tvm.build( + sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon) + ) + module = hexagon_session.load_module(func_tir) + + 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( + np.zeros((size_a, 32), dtype="int32"), device=hexagon_session.device + ) + + if tvm.testing.utils.IS_IN_CI: + # These are reduced for CI + number = 1 + repeat = 1 + else: + number = 100 Review Comment: This one seems a bit circular since it would be ``` timer_number = 100 timer_repeat = 100 if tvm.testing.utils.IS_IN_CI: # These are reduced for CI number = 1 repeat = 1 else: number = timer_number repeat = timer_repeat timer = module.time_evaluator( "__tvm_main__", hexagon_session.device, number=number, repeat=repeat ) ``` Which makes for quite a bit of unnecessary code. Also this is to some extent a magic number. I just chose it from experience and not tied to any functionality. The other option is I could do this ``` if tvm.testing.utils.IS_IN_CI: # These are reduced for CI timer = module.time_evaluator( "__tvm_main__", hexagon_session.device, number=1, repeat=1 ) else: timer = module.time_evaluator( "__tvm_main__", hexagon_session.device, number=100, repeat=100 ) ``` But not sure that is better. Anyway 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: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
