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 2020/06/18 15:00:01 UTC
[GitHub] [incubator-tvm] zhanghaohit opened a new pull request #5842: [VTA][OpenCL] Cloud FPGA support
zhanghaohit opened a new pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842
This PR, coupled with [this](https://github.com/apache/incubator-tvm-vta/pull/9) on tvm-vta repo, is the basic implementation of RFC #4958
Some notes:
- this is just a basic version without much performance optimization. We've done some optimization, and achieved significant improvement (but this part of code is not ready; has to be organized and cleaned up further).
- there are some experimental features, which are also included in this PR, including
- sync all the instructions per model (instead of per layer)
- static auto-tune using profiling results
These codes are not well styled (using some environmental variables), and we have to think about how to format/implement nicely. But we think these features are useful so we also commit. We can discuss if we should include these codes in this PR, or leave it for other PRs.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-660356504
I'm thinking that the `static auto-tune using profiling results` can also be separated out in a standalone PR
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r443090563
##########
File path: vta/python/vta/environment.py
##########
@@ -241,7 +243,7 @@ def target_host(self):
return "llvm -target=armv7-none-linux-gnueabihf"
if self.TARGET == "ultra96":
return "llvm -target=aarch64-linux-gnu"
- if self.TARGET in ["sim", "tsim"]:
+ if self.TARGET in ["sim", "tsim", "intelfocl"]:
Review comment:
We may want to be a little more specific about the hardware target
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-646427274
CC-ing @vegaluisjose @huajsj @pasqoc
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r443337640
##########
File path: src/tir/transforms/lower_tvm_builtin.cc
##########
@@ -86,16 +86,19 @@ class BuiltinLower : public StmtExprMutator {
op = stmt.as<AllocateNode>();
// Get constant allocation bound.
int64_t nbytes = GetVectorBytes(op->dtype);
- if (device_type_.defined()) {
- if (const auto* dev_type = device_type_.as<IntImmNode>()) {
- if (dev_type->value == kDLCPU) {
- int32_t constant_size = op->constant_allocation_size();
- if (constant_size > 0 && constant_size * nbytes < runtime::kMaxStackAlloca) {
- return stmt;
- }
- }
- }
- }
+ // NOTE(zhanghao): remove special handling for kDLCPU
Review comment:
```bash
Traceback (most recent call last):
File "vta/tutorials/frontend/deploy_classification.py", line 210, in <module>
params=params, target_host=env.target_host)
File "/4pd/home/zhanghao/workspace/tvm-2/tvm/python/tvm/relay/build_module.py", line 251, in build
graph_json, mod, params = bld_mod.build(mod, target, target_host, params)
File "/4pd/home/zhanghao/workspace/tvm-2/tvm/python/tvm/relay/build_module.py", line 120, in build
self._build(mod, target, target_host)
File "tvm/_ffi/_cython/./packed_func.pxi", line 321, in tvm._ffi._cy3.core.PackedFuncBase.__call__
File "tvm/_ffi/_cython/./packed_func.pxi", line 256, in tvm._ffi._cy3.core.FuncCall
File "tvm/_ffi/_cython/./packed_func.pxi", line 245, in tvm._ffi._cy3.core.FuncCall3
File "tvm/_ffi/_cython/./base.pxi", line 160, in tvm._ffi._cy3.core.CALL
tvm._ffi.base.TVMError: Traceback (most recent call last):
[bt] (8) /4pd/home/zhanghao/workspace/tvm-2/tvm/build/libtvm.so(TVMFuncCall+0x4c) [0x7f385ac9bc1c]
[bt] (7) /4pd/home/zhanghao/workspace/tvm-2/tvm/build/libtvm.so(std::_Function_handler<void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), tvm::relay::backend::RelayBuildModule::GetFunction(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#3}>::_M_invoke(std::_Any_data const&, tvm::runtime::TVMArgs&&, tvm::runtime::TVMRetValue*&&)+0x316) [0x7f385ab2a566]
[bt] (6) /4pd/home/zhanghao/workspace/tvm-2/tvm/build/libtvm.so(tvm::relay::backend::RelayBuildModule::BuildRelay(tvm::IRModule, std::unordered_map<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, tvm::runtime::NDArray, std::hash<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > >, std::equal_to<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > >, std::allocator<std::pair<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const, tvm::runtime::NDArray> > > const&)+0xe31) [0x7f385ab29c11]
[bt] (5) /4pd/home/zhanghao/workspace/tvm-2/tvm/build/libtvm.so(tvm::build(tvm::Map<tvm::runtime::String, tvm::IRModule, void, void> const&, tvm::Target const&)+0x3c4) [0x7f385a4322d4]
[bt] (4) /4pd/home/zhanghao/workspace/tvm-2/tvm/build/libtvm.so(tvm::build(tvm::Map<tvm::Target, tvm::IRModule, void, void> const&, tvm::Target const&)+0x326) [0x7f385a4318c6]
[bt] (3) /4pd/home/zhanghao/workspace/tvm-2/tvm/build/libtvm.so(tvm::codegen::Build(tvm::IRModule, tvm::Target const&)+0x67a) [0x7f385a74f68a]
[bt] (2) /4pd/home/zhanghao/workspace/tvm-2/tvm/build/libtvm.so(+0x1277ea1) [0x7f385ac7eea1]
[bt] (1) /4pd/home/zhanghao/workspace/tvm-2/tvm/build/libtvm.so(tvm::codegen::LLVMModuleNode::Init(tvm::IRModule const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)+0x1388) [0x7f385ac82c68]
[bt] (0) /4pd/home/zhanghao/workspace/tvm-2/tvm/build/libtvm.so(+0x1276a57) [0x7f385ac7da57]
File "/4pd/home/zhanghao/workspace/tvm-2/tvm/src/target/llvm/llvm_module.cc", line 230
TVMError: LLVM module verification failed with the following errors:
Call parameter type does not match function signature!
%.sub = getelementptr inbounds [4 x <8 x float>], [4 x <8 x float>]* %3, i64 0, i64 0
i8* %34 = call i8* @VTABufferCPUPtr(i8* %17, <8 x float>* nonnull %.sub)
Call parameter type does not match function signature!
%.sub = getelementptr inbounds [8 x float], [8 x float]* %3, i64 0, i64 0
i8* %31 = call i8* @VTABufferCPUPtr(i8* %14, float* nonnull %.sub)
```
The raise error is due to the LLVM code here (lib/IR/Verifier.cpp):
```{.c++ filename="lib/IR/Verifier.cpp"}
2598 // Verify that all arguments to the call match the function type.
2599 for (unsigned i = 0, e = FTy->getNumParams(); i != e; ++i)
2600 Assert(CS.getArgument(i)->getType() == FTy->getParamType(i),
2601 "Call parameter type does not match function signature!",
2602 CS.getArgument(i), FTy->getParamType(i), I);
```
It will raise this error if the special handling for kDLCPU is there. I think it is because the signature for the AllocateNode is not consistent with the parameter? Any ideas about alternative fix?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r456670755
##########
File path: vta/python/vta/environment.py
##########
@@ -241,7 +243,7 @@ def target_host(self):
return "llvm -target=armv7-none-linux-gnueabihf"
if self.TARGET == "ultra96":
return "llvm -target=aarch64-linux-gnu"
- if self.TARGET in ["sim", "tsim"]:
+ if self.TARGET in ["sim", "tsim", "intelfocl"]:
Review comment:
Could you follow up and rename to say, "arria10"?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-659791729
@tmoreau89 @vegaluisjose @huajsj @pasqoc
Any comments are welcome. Thanks.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] liangfu edited a comment on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
liangfu edited a comment on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-659829041
Thanks for the update @zhanghaohit , please rebase upon latest master branch and resolve the conflicts.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r443090507
##########
File path: src/tir/transforms/lower_tvm_builtin.cc
##########
@@ -86,16 +86,19 @@ class BuiltinLower : public StmtExprMutator {
op = stmt.as<AllocateNode>();
// Get constant allocation bound.
int64_t nbytes = GetVectorBytes(op->dtype);
- if (device_type_.defined()) {
- if (const auto* dev_type = device_type_.as<IntImmNode>()) {
- if (dev_type->value == kDLCPU) {
- int32_t constant_size = op->constant_allocation_size();
- if (constant_size > 0 && constant_size * nbytes < runtime::kMaxStackAlloca) {
- return stmt;
- }
- }
- }
- }
+ // NOTE(zhanghao): remove special handling for kDLCPU
Review comment:
can you elaborate a little more about the LLVM parameter match error you're getting?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-650879202
> @zhanghaohit good news, the end to end tests ran successfully on VTA hardware with your changes. Given that this PR brings modifications to the ISA spec, we need to bump the `HW_VER` string from `0.0.1` to `0.0.2`. I can upload the Pynq bitstream so it's ready for use. I will need some help from @liangfu to re-generate the DE-10 bitstream as well.
Thanks @tmoreau89 for the help with the tests.
> In addition were you able to put a guide together? Which Intel FPGA have you tested this on?
Sure. What guide are you referring to? The guide on HW_VER?
We're using Intel Arria 10.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r456673094
##########
File path: python/tvm/autotvm/measure/measure_methods.py
##########
@@ -186,6 +189,16 @@ def __init__(self,
timeout=10, n_parallel=None,
number=4, repeat=3, min_repeat_ms=0, cooldown_interval=0.1,
check_correctness=False):
+ static_tune = os.getenv("TVM_STATIC_TUNE_EXPERIMENTAL")
Review comment:
Can you elaborate on what TVM_STATIC_TUNE_EXPERIMENTAL is being used for? I'm a little weary of inserting additional execution modes in AutoTVM dictated by environmental variables.
(1) we either add additional flags in AutoTVM config
(2) we clean this experimental mode if it's not really necessary for the OpenCL variant of VTA
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-648435167
Thanks @zhanghaohit I can generate the new VTA design on Pynq to verify that it still passes on the new ISA. I will let you know if this passes. Installation documentation will be highly welcome.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-647009494
> Thanks @zhanghaohit for the contribution! I've taken a quick pass but will need to make a more thorough pass next.
>
> I believe it's important to check two things:
> (1) that the new ISA changes won't break the current Pynq tests and binary. We'll need to update the bitstreams to make sure that it can work out of the box on the new ISA.
> (2) we would want to prepare some setup instructions for this Altera OpenCL variant of VTA to run end to end tests. Something like this: https://docs.tvm.ai/vta/install.html, which @liangfu for instance added instructions for Intel SoC FPGAs
Thanks @tmoreau89 for the comments.
(1) I've back-tested the Pynq tests. It works.
The only thing we have to do is to re-compile the bitstream to make it compatible with the new ISA. Yes. We have to update the bitstreams in the repo.
(2) Sorry that we forgot the document for the OpenCL VTA installation. We'll update the docs accordingly.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit edited a comment on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit edited a comment on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-659790884
> Note that I've also uploaded the Pynq bitstream given the ISA changes here: https://github.com/uwsampl/vta-distro/tree/master/bitstreams/pynq/0.0.2
>
> It would be great @liangfu if you had the bandwidth to synthesize the DE10 bistream as well and submit a pull request. Let me know if you are tight on time and may need some help.
>
> Finally, I will add an ultra96 variant in there as well. It would be great @zhanghaohit if you could also pre-generate the Arria10 bitstream and add it to the `vta-distro` repo.
Thanks @tmoreau89 for the help. For the Arria10 bitstream, because the fpga device we are currently using is from other vendors (not Intel), it is not easy to be used by others. We're trying to adapt to the fpga devices on AWS, which are more standard. After that, we can upload the aocx/bitstream.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-660355629
@zhanghaohit I took a second look at the changes in this PR, and I believe it would be best to break the PR down into separate, smaller PRs:
(1) modifications to VTA for OpenCL support (runtime, operator support etc.)
(2) quantization support for conv2d_transpose which is device agnostic
(3) modifications to device_annotation.cc
(4) once (3) is merged and approved, updating GraphPack pass to use device annotation
Let me know what you think. This will make the PRs easier to review and merge.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-662859131
> @zhanghaohit I took a second look at the changes in this PR, and I believe it would be best to break the PR down into separate, smaller PRs:
>
> (1) modifications to VTA for OpenCL support (runtime, operator support etc.)
> (2) quantization support for conv2d_transpose which is device agnostic
> (3) modifications to device_annotation.cc
> (4) once (3) is merged and approved, updating GraphPack pass to use device annotation
>
> Let me know what you think. This will make the PRs easier to review and merge.
Thanks @tmoreau89 for the suggestion. I've split into 3 small PRs here:
#6124
#6125
#6126
And remove the unnecessary code that are not very related to OpenCL support.
Thanks.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit edited a comment on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit edited a comment on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-647009494
> Thanks @zhanghaohit for the contribution! I've taken a quick pass but will need to make a more thorough pass next.
>
> I believe it's important to check two things:
> (1) that the new ISA changes won't break the current Pynq tests and binary. We'll need to update the bitstreams to make sure that it can work out of the box on the new ISA.
> (2) we would want to prepare some setup instructions for this Altera OpenCL variant of VTA to run end to end tests. Something like this: https://docs.tvm.ai/vta/install.html, which @liangfu for instance added instructions for Intel SoC FPGAs
Thanks @tmoreau89 for the comments.
(1) I've back-tested the ultra96 board (I only have ultra96 board on hand). It works.
The only thing we have to do is to re-compile the bitstream to make it compatible with the new ISA. Yes. We have to update the bitstreams in the repo.
(2) Sorry that we forgot the document for the OpenCL VTA installation. We'll update the docs accordingly.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r446732412
##########
File path: vta/python/vta/environment.py
##########
@@ -241,7 +243,7 @@ def target_host(self):
return "llvm -target=armv7-none-linux-gnueabihf"
if self.TARGET == "ultra96":
return "llvm -target=aarch64-linux-gnu"
- if self.TARGET in ["sim", "tsim"]:
+ if self.TARGET in ["sim", "tsim", "intelfocl"]:
Review comment:
ok. sure. we'll change the naming
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-650347598
In addition were you able to put a guide together? Which Intel FPGA have you tested this on?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r456797915
##########
File path: python/tvm/relay/op/strategy/generic.py
##########
@@ -323,8 +331,10 @@ def compute_conv2d_transpose(attrs, inputs, out_dtype):
out = topi_compute(
inputs[0], inputs[1], strides, padding, out_dtype)
output_padding = get_const_tuple(attrs.output_padding)
- out = topi.nn.pad(out, [0, 0, 0, 0],
- [0, 0, output_padding[0], output_padding[1]])
+ if output_padding[0] != 0 or output_padding[1] != 0:
Review comment:
> This will cause errors when running dcgan on vta. There are two issues here:
>
> 1. if output_padding are all 0, no need to do pad
> 2. if the size of out.shape is not 4, it will cause problems.
After rebase, this changes are not needed.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r456745434
##########
File path: python/tvm/autotvm/measure/measure_methods.py
##########
@@ -186,6 +189,16 @@ def __init__(self,
timeout=10, n_parallel=None,
number=4, repeat=3, min_repeat_ms=0, cooldown_interval=0.1,
check_correctness=False):
+ static_tune = os.getenv("TVM_STATIC_TUNE_EXPERIMENTAL")
Review comment:
`TVM_STATIC_TUNE_EXPERIMENTAL` is used for static auto tune. I think as you advised, we can split the PR into separate PRs. For static auto tune, the interface is not well nicely done. I think maybe we remove this part first.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r443090782
##########
File path: vta/tutorials/autotvm/tune_alu_vta.py
##########
@@ -0,0 +1,320 @@
+# 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.
+"""
+Auto-tuning a ALU fused op on VTA
+"""
+
+import os
+from mxnet.gluon.model_zoo import vision
+import numpy as np
+from PIL import Image
+
+import topi
+import tvm
+from tvm import te
+from tvm import rpc, autotvm, relay
+from tvm.contrib import graph_runtime, util, download
+from tvm.autotvm.measure.measure_methods import request_remote
+from tvm.autotvm.tuner import XGBTuner, GATuner, RandomTuner, GridSearchTuner
+from tvm.autotvm import record
+
+import vta
+from vta.testing import simulator
+from vta.top import graph_pack
+import copy
+
+
+#################################################################
+# Compile network
+# ---------------
+# Perform vta-specific compilation with Relay from a Gluon model
+def compile_network(env, target, model, start_pack, stop_pack, device_annot=False):
+
+ # Populate the shape and data type dictionary
+ dtype_dict = {"data": 'float32'}
+ shape_dict = {"data": (env.BATCH, 3, 224, 224)}
+
+ # Get off the shelf gluon model, and convert to relay
+ gluon_model = vision.get_model(model, pretrained=True)
+ mod, params = relay.frontend.from_mxnet(gluon_model, shape_dict)
+
+ # Update shape and type dictionary
+ shape_dict.update({k: v.shape for k, v in params.items()})
+ dtype_dict.update({k: str(v.dtype) for k, v in params.items()})
+
+ # Perform quantization in Relay
+ # Note: We set opt_level to 3 in order to fold batch norm
+ with relay.build_config(opt_level=3):
+ with relay.quantize.qconfig(global_scale=8.0, skip_conv_layers=[0]):
+ mod = relay.quantize.quantize(mod, params=params)
+
+ # Perform graph packing and constant folding for VTA target
+ if target.device_name == "vta":
+ assert env.BLOCK_IN == env.BLOCK_OUT
+ relay_prog = graph_pack(mod["main"],
+ env.BATCH,
+ env.BLOCK_OUT,
+ env.WGT_WIDTH,
+ start_name=start_pack,
+ stop_name=stop_pack,
+ device_annot=device_annot)
+
+ return relay_prog, params
+
+
+###########################################
+# Set Tuning Options
+# ------------------
+# Before tuning, we should apply some configurations.
+# Here we use an Pynq-Z1 board as an example.
+
+# Tracker host and port can be set by your environment
+tracker_host = os.environ.get("TVM_TRACKER_HOST", '0.0.0.0')
+tracker_port = int(os.environ.get("TVM_TRACKER_PORT", 9190))
+
+# Load VTA parameters from the vta/config/vta_config.json file
+env = vta.get_env()
+
+# This target is used for cross compilation. You can query it by :code:`gcc -v` on your device.
+# Set ``device=arm_cpu`` to run inference on the CPU
+# or ``device=vta`` to run inference on the FPGA.
+device = "vta"
+target = env.target if device == "vta" else env.target_vta_cpu
+
+# Name of Gluon model to compile
+# The ``start_pack`` and ``stop_pack`` labels indicate where
+# to start and end the graph packing relay pass: in other words
+# where to start and finish offloading to VTA.
+network = "resnet50_v2"
+start_pack = "nn.max_pool2d"
+stop_pack = "nn.global_avg_pool2d"
+
+# Tuning option
+log_file = "%s.alu.%s.log" % (device, network)
+tuning_option = {
+ 'log_filename': log_file,
+
+ 'tuner': 'random',
+ 'n_trial': 1000,
+ 'early_stopping': None,
+
+ 'measure_option': autotvm.measure_option(
+ builder=autotvm.LocalBuilder(n_parallel=1),
+ runner=autotvm.RPCRunner(env.TARGET,
+ host=tracker_host,
+ port=tracker_port,
+ number=5,
+ timeout=60,
+ check_correctness=True),
+ ),
+}
+
+
+def log_to_file(file_out, protocol='json'):
+ """Log the tuning records into file.
+ The rows of the log are stored in the format of autotvm.record.encode.
+ for lhs == rhs, we add an extra rhs = [] record
+
+ Parameters
+ ----------
+ file_out : str
+ The file to log to.
+ protocol: str, optional
+ The log protocol. Can be 'json' or 'pickle'
+
+ Returns
+ -------
+ callback : callable
+ Callback function to do the logging.
+ """
+ def _callback(_, inputs, results):
+ with open(file_out, "a") as f:
+ for inp, result in zip(inputs, results):
+ f.write(record.encode(inp, result, protocol) + "\n")
+
+ # we only consider task with same lhs and rhs
+ if inp.task.args[0] == inp.task.args[1]:
+ args = list(inp.task.args)
+ args[1] = (args[0][0], (), args[0][2])
+ inp_copy = copy.deepcopy(inp)
+ inp_copy.task.args = tuple(args)
+ f.write(record.encode(inp_copy, result, protocol) + "\n")
+
+ return _callback
+
+
+def tune_tasks(tasks,
+ measure_option,
+ tuner='xgb',
+ n_trial=10,
+ early_stopping=None,
+ log_filename='tuning.log',
+ use_transfer_learning=True):
+
+ # create tmp log file
+ tmp_log_file = log_filename + ".tmp"
+ if os.path.exists(tmp_log_file):
+ os.remove(tmp_log_file)
+
+ for i, tsk in enumerate(reversed(tasks)):
+ prefix = "[Task %2d/%2d] " % (i + 1, len(tasks))
+
+ # create tuner
+ if tuner == 'xgb' or tuner == 'xgb-rank':
+ tuner_obj = XGBTuner(tsk, loss_type='rank')
+ elif tuner == 'xgb_knob':
+ tuner_obj = XGBTuner(tsk, loss_type='rank', feature_type='knob')
+ elif tuner == 'ga':
+ tuner_obj = GATuner(tsk, pop_size=50)
+ elif tuner == 'random':
+ tuner_obj = RandomTuner(tsk)
+ elif tuner == 'gridsearch':
+ tuner_obj = GridSearchTuner(tsk)
+ else:
+ raise ValueError("Invalid tuner: " + tuner)
+
+ if use_transfer_learning:
+ if os.path.isfile(tmp_log_file):
+ tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file))
+
+ # do tuning
+ tsk_trial = min(n_trial, len(tsk.config_space))
+ tuner_obj.tune(n_trial=tsk_trial,
+ early_stopping=early_stopping,
+ measure_option=measure_option,
+ callbacks=[
+ autotvm.callback.progress_bar(tsk_trial, prefix=prefix),
+ log_to_file(tmp_log_file)
+ ])
+
+ # pick best records to a cache file
+ autotvm.record.pick_best(tmp_log_file, log_filename)
+ os.remove(tmp_log_file)
+
+
+########################################################################
+# Register VTA-specific tuning tasks
+def register_vta_tuning_tasks():
+ from tvm.autotvm.task import TaskExtractEnv
+
+ @tvm.te.tag_scope(tag=topi.tag.ELEMWISE)
+ def my_clip(x, a_min, a_max):
+ """Unlike topi's current clip, put min and max into two stages."""
+ const_min = tvm.tir.const(a_min, x.dtype)
+ const_max = tvm.tir.const(a_max, x.dtype)
+ x = te.compute(x.shape, lambda *i: tvm.te.min(x(*i), const_max), name="clipA")
+ x = te.compute(x.shape, lambda *i: tvm.te.max(x(*i), const_min), name="clipB")
+ return x
+
+ # init autotvm env to register VTA operator
+ TaskExtractEnv()
+
+ @autotvm.template("add.vta")
+ def _topi_add(*args, **kwargs):
+ assert not kwargs, "Do not support kwargs in template function call"
+ A, B = args[:2]
+
+ with tvm.target.vta():
+ res = vta.top.op.add_packed(*args, **kwargs)
+ res = my_clip(res, 0, 127)
+ res = topi.cast(res, "int8")
+
+ if tvm.target.Target.current().device_name == 'vta':
+ s = vta.top.op.schedule_add_packed([res])
+ else:
+ s = te.create_schedule([res.op])
+ return s, [A, B, res]
+
+ @autotvm.template("multiply.vta")
+ def _topi_multiply(*args, **kwargs):
+ assert not kwargs, "Do not support kwargs in template function call"
+ A, B = args[:2]
+
+ with tvm.target.vta():
+ res = vta.top.op.multiply_packed(*args, **kwargs)
+ res = my_clip(res, 0, 127)
+ res = topi.cast(res, "int8")
+
+ if tvm.target.Target.current().device_name == 'vta':
+ s = vta.top.op.schedule_multiply_packed([res])
+ else:
+ s = te.create_schedule([res.op])
+ return s, [A, B, res]
+
+
+########################################################################
+# Finally, we launch tuning jobs and evaluate the end-to-end performance.
+def tune_and_evaluate(tuning_opt):
+
+ if env.TARGET != "sim":
+ # Get remote from fleet node
+ remote = autotvm.measure.request_remote(env.TARGET,
+ tracker_host,
+ tracker_port,
+ timeout=10000)
+ # Reconfigure the JIT runtime and FPGA.
+ vta.reconfig_runtime(remote)
+ vta.program_fpga(remote, bitstream=None)
+ else:
+ # In simulation mode, host the RPC server locally.
+ remote = rpc.LocalSession()
+
+ # Register VTA tuning tasks
+ register_vta_tuning_tasks()
+
+ # Perform task extraction on Relay program
+ print("Extract tasks...")
+ relay_prog, params = compile_network(env, target, network, start_pack, stop_pack)
+ mod = tvm.IRModule.from_expr(relay_prog)
+ tasks = autotvm.task.extract_from_program(mod,
+ params=params,
+ ops=(relay.op.get("add"), relay.op.get("multiply"),),
+ target=target,
+ target_host=env.target_host)
+
+ # filter out non-packed alu task
+ tasks = list(filter(lambda t: len(t.args[0][1]) > 4, tasks))
+ # filter out float alu task
+ tasks = list(filter(lambda t: t.args[0][2] != "float32", tasks))
+ # filter const rhs, which will be fused with conv2d
+ # tasks = list(filter(lambda t: len(t.args[1][1]) < 1, tasks))
Review comment:
clean up this line?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r456530427
##########
File path: 3rdparty/aoclutils/opencl.cc
##########
@@ -0,0 +1,555 @@
+// Copyright (C) 2013-2018 Altera Corporation, San Jose, California, USA. All rights reserved.
Review comment:
@tqchen will this copyright work with Apache licensing rules? this will end up in 3rd party
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-653285145
Note that I've also uploaded the Pynq bitstream given the ISA changes here: https://github.com/uwsampl/vta-distro/tree/master/bitstreams/pynq/0.0.2
It would be great @liangfu if you had the bandwidth to synthesize the DE10 bistream as well and submit a pull request. Let me know if you are tight on time and may need some help.
Finally, I will add an ultra96 variant in there as well. It would be great @zhanghaohit if you could also pre-generate the Arria10 bitstream and add it to the `vta-distro` repo.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tqchen closed pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tqchen closed pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-660185671
@zhanghaohit thanks for the follow up. I recommend to echo @liangfu to rebase so the PR can be in a mergeable state. In addition, if the target is a custom FPGA board, I think we can leave it as a generic target for now. Having it ported to AWS F1 should be a great addition. I understand having setup instructions at this time would not make the most sense.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r456673553
##########
File path: python/tvm/relay/op/strategy/generic.py
##########
@@ -323,8 +331,10 @@ def compute_conv2d_transpose(attrs, inputs, out_dtype):
out = topi_compute(
inputs[0], inputs[1], strides, padding, out_dtype)
output_padding = get_const_tuple(attrs.output_padding)
- out = topi.nn.pad(out, [0, 0, 0, 0],
- [0, 0, output_padding[0], output_padding[1]])
+ if output_padding[0] != 0 or output_padding[1] != 0:
Review comment:
Which operators was this breaking on previously?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-659790884
> Note that I've also uploaded the Pynq bitstream given the ISA changes here: https://github.com/uwsampl/vta-distro/tree/master/bitstreams/pynq/0.0.2
>
> It would be great @liangfu if you had the bandwidth to synthesize the DE10 bistream as well and submit a pull request. Let me know if you are tight on time and may need some help.
>
> Finally, I will add an ultra96 variant in there as well. It would be great @zhanghaohit if you could also pre-generate the Arria10 bitstream and add it to the `vta-distro` repo.
Thanks @tmoreau89 for the help. For the Arria10 bitstream, because the fpga device we are currently using is not from other vendors (not Intel), it is not easy to be used by others. We're trying to adapt to the fpga devices on AWS, which are more standard. After that, we can upload the aocx/bitstream.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-650347427
@zhanghaohit good news, the end to end tests ran successfully on VTA hardware with your changes. Given that this PR brings modifications to the ISA spec, we need to bump the `HW_VER` string from `0.0.1` to `0.0.2`. I can upload the Pynq bitstream so it's ready for use. I will need some help from @liangfu to re-generate the DE-10 bitstream as well.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r456671336
##########
File path: vta/python/vta/transform.py
##########
@@ -381,9 +381,10 @@ def _fold_buffer_dim(buf, scope, elem_block):
def _get_2d_pattern(buf, elem_width, elem_bytes, dtype, scope, allow_fold):
elem_block = elem_bytes * 8 // elem_width
- if buf.dtype != dtype:
- raise RuntimeError("Expect buffer type to be %s instead of %s" %
- (dtype, buf.dtype))
+ # remove the checking as we have load_int8 insn
Review comment:
Please remove those lines
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] zhanghaohit commented on a change in pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
zhanghaohit commented on a change in pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#discussion_r456748692
##########
File path: python/tvm/relay/op/strategy/generic.py
##########
@@ -323,8 +331,10 @@ def compute_conv2d_transpose(attrs, inputs, out_dtype):
out = topi_compute(
inputs[0], inputs[1], strides, padding, out_dtype)
output_padding = get_const_tuple(attrs.output_padding)
- out = topi.nn.pad(out, [0, 0, 0, 0],
- [0, 0, output_padding[0], output_padding[1]])
+ if output_padding[0] != 0 or output_padding[1] != 0:
Review comment:
This will cause errors when running dcgan on vta. There are two issues here:
1. if output_padding are all 0, no need to do pad
2. if the size of out.shape is not 4, it will cause problems.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tmoreau89 commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-646918453
Thanks @zhanghaohit for the contribution! I've taken a quick pass but will need to make a more thorough pass next.
I believe it's important to check two things:
(1) that the new ISA changes won't break the current Pynq tests and binary. We'll need to update the bitstreams to make sure that it can work out of the box on the new ISA.
(2) we would want to prepare some setup instructions for this Altera OpenCL variant of VTA to run end to end tests. Something like this: https://docs.tvm.ai/vta/install.html, which @liangfu for instance added instructions for Intel SoC FPGAs
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] liangfu commented on pull request #5842: [VTA][OpenCL] Cloud FPGA support
Posted by GitBox <gi...@apache.org>.
liangfu commented on pull request #5842:
URL: https://github.com/apache/incubator-tvm/pull/5842#issuecomment-659829041
Please rebase upon latest master branch and resolve the conflicts.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org