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/12/01 18:27:10 UTC
[GitHub] [tvm] AndrewZhaoLuo opened a new pull request, #13532: [Packed Func] Pack args fp16 support
AndrewZhaoLuo opened a new pull request, #13532:
URL: https://github.com/apache/tvm/pull/13532
Needs testing
Should solve https://discuss.tvm.apache.org/t/resnet-tomixedprecision-tuning-error/14015
--
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] tqchen commented on pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
tqchen commented on PR #13532:
URL: https://github.com/apache/tvm/pull/13532#issuecomment-1334474222
I am not too sure if we should do this. Mainly because the fp16 intrinsic can be slow and software based on some of the kernels. How about considering passing fp32 as scalar argument and do live conversion in kernel when we look at GPU setting?
--
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] shingjan commented on pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
shingjan commented on PR #13532:
URL: https://github.com/apache/tvm/pull/13532#issuecomment-1352641889
Same concern here. Right now I am just cherry-picking this pr and it worked. Will circle back hopefully next week.
--
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] AndrewZhaoLuo commented on pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
AndrewZhaoLuo commented on PR #13532:
URL: https://github.com/apache/tvm/pull/13532#issuecomment-1371470841
Can confirm that this was a CSE issue which is now fixed for all the use cases I care about. masahi has a good summary in the above of what occured. I am closing for now to avoid delving deeper.
--
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] masahi commented on pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
masahi commented on PR #13532:
URL: https://github.com/apache/tvm/pull/13532#issuecomment-1353715208
Turns out, for winograd fp16, the issue was due to the following `cse_vars`
```
let cse_var_11 = (0h*2h)
let cse_var_10 = (0h*2.5h)
let cse_var_9 = (0h*1h)
let cse_var_8 = (0h*1.5h)
let cse_var_7 = (0h*0h)
let cse_var_6 = (0h*0.5h)
let cse_var_5 = (0h*-2h)
let cse_var_4 = (0h*-2.5h)
let cse_var_3 = (0h*-1h)
let cse_var_2 = (0h*-1.5h)
let cse_var_1 = (0h*-0.5h)
```
which are used by two kernels in winograd, so they are computed on the host once and passed to the winograd kernels via `SplitHostDevice`. I'm going to update the arithmetic simplifier to remove such useless much (surprised it doesn't do it already).
@AndrewZhaoLuo I wonder, for layer_norm what variables are being passed between two kernels? Useless variables like above or something meaningful?
--
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] masahi commented on pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
masahi commented on PR #13532:
URL: https://github.com/apache/tvm/pull/13532#issuecomment-1352627955
> One workaround for now is disabling CSE-TIR
Yup can confirm that disabling TIR CSE fixes this issue. Turns out the combination of CSE + `SplitHostDevice` is generating a signature like this
```
extern "C" __global__ void __launch_bounds__(256) main_kernel2(half* __restrict__ inverse, half* __restrict__ bgemm, half cse_var_12, half cse_var_10, half cse_var_3, half cse_var_9, half cse_var_5, half cse_var_8, half cse_var_15, half cse_var_7, half cse_var_6) {
```
and I'm not sure if this is a good idea.
--
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] tvm-bot commented on pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
tvm-bot commented on PR #13532:
URL: https://github.com/apache/tvm/pull/13532#issuecomment-1334177257
<!---bot-comment-->
Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @-ing them in a comment.
<sub>Generated by [tvm-bot](https://github.com/apache/tvm/blob/main/ci/README.md#github-actions)</sub>
--
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] masahi commented on pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
masahi commented on PR #13532:
URL: https://github.com/apache/tvm/pull/13532#issuecomment-1354323713
Found a bug https://github.com/apache/tvm/pull/13631
@shingjan This may fix your issue.
--
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] AndrewZhaoLuo commented on pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
AndrewZhaoLuo commented on PR #13532:
URL: https://github.com/apache/tvm/pull/13532#issuecomment-1352223089
@shingjan I will have time next week to look at this and try TQ' suggestion. One workaround for now is disabling CSE-TIR pass I believe for now.
--
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] shingjan commented on pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
shingjan commented on PR #13532:
URL: https://github.com/apache/tvm/pull/13532#issuecomment-1351946423
Hit this problem while tuning some convolutions from resnet18 on cuda w. fp16. Comment to keep in the loop of this problem.
--
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] masahi commented on pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
masahi commented on PR #13532:
URL: https://github.com/apache/tvm/pull/13532#issuecomment-1352409775
Also getting the error
```
2: tvm::runtime::CUDAModuleNode::GetFunction(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, tvm::runtime::Ob
jectPtr<tvm::runtime::Object> const&)
1: tvm::runtime::PackedFunc tvm::runtime::PackFuncVoidAddr<tvm::runtime::CUDAWrappedFunc>(tvm::runtime::CUDAWrappedFunc, std::vector<DLDataType, std:
:allocator<DLDataType> > const&)
0: tvm::runtime::detail::GetArgConvertCode(DLDataType)
File "/home/masa/projects/dev/tvm/src/runtime/cuda/../pack_args.h", line 146
File "/home/masa/projects/dev/tvm/src/runtime/library_module.cc", line 80
TVMError:
---------------------------------------------------------------
An error occurred during the execution of TVM.
For more information, please see: https://tvm.apache.org/docs/errors.html
---------------------------------------------------------------
Check failed: ret == 0 (-1 vs. 0) : TVMError: Cannot handle float16 as device function argument
```
when tuning winograd fp16 tasks from stable diffusion UNet. Normal convolution works.
--
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] AndrewZhaoLuo closed pull request #13532: [Packed Func] Pack args fp16 support
Posted by GitBox <gi...@apache.org>.
AndrewZhaoLuo closed pull request #13532: [Packed Func] Pack args fp16 support
URL: https://github.com/apache/tvm/pull/13532
--
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