[CodeGen][CUDA] Fix issues in cuda codegen #4876
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
This patch should fix errors observed below (I did not verify as I found no complete reproducers there). My own test works fine with CUDA 10.2. https://discuss.tvm.ai/t/error-fp16-cuda-compilation-error/4586 This issue has been also reported to NVIDIA CUDA team. |
|
cc @vinx13 @ZihengJiang please help to take a look |
wpan11nv
force-pushed
the
fp16_reduction_fixes
branch
2 times, most recently
from
c31bccd
to
d3d8b0c
Compare
|
also cc @yzhliu @zxy844288792 @Hzfengsy |
- Do not emit __shared__ etc. as part of type for casting - Fix fp16 reduction kernels with compiler errors: "no operator "+" matches these operands, volatile half + volatile half This patch inserts casts to remove volatile type qualifier following volatile loads (fp16 only). CUDA fp16 library headers should add volatile member functions. - Update have_fp16 to include compute 6.1 GPUs, which do support fp16, although their fp16 throughput is low. Updated tests. Signed-off-by: Wei Pan <weip@nvidia.com>
wpan11nv
force-pushed
the
fp16_reduction_fixes
branch
from
d3d8b0c
to
44d463a
Compare
|
Updated the patch as suggested. Thanks! |
vinx13
approved these changes
Feb 16, 2020
|
Thanks @wpan11nv this is merged |
alexwong
pushed a commit
to alexwong/tvm
that referenced
this pull request
Feb 26, 2020
- Do not emit __shared__ etc. as part of type for casting - Fix fp16 reduction kernels with compiler errors: "no operator "+" matches these operands, volatile half + volatile half This patch inserts casts to remove volatile type qualifier following volatile loads (fp16 only). CUDA fp16 library headers should add volatile member functions. - Update have_fp16 to include compute 6.1 GPUs, which do support fp16, although their fp16 throughput is low. Updated tests. Signed-off-by: Wei Pan <weip@nvidia.com>
alexwong
pushed a commit
to alexwong/tvm
that referenced
this pull request
Feb 28, 2020
- Do not emit __shared__ etc. as part of type for casting - Fix fp16 reduction kernels with compiler errors: "no operator "+" matches these operands, volatile half + volatile half This patch inserts casts to remove volatile type qualifier following volatile loads (fp16 only). CUDA fp16 library headers should add volatile member functions. - Update have_fp16 to include compute 6.1 GPUs, which do support fp16, although their fp16 throughput is low. Updated tests. Signed-off-by: Wei Pan <weip@nvidia.com>
zhiics
pushed a commit
to neo-ai/tvm
that referenced
this pull request
Mar 2, 2020
- Do not emit __shared__ etc. as part of type for casting - Fix fp16 reduction kernels with compiler errors: "no operator "+" matches these operands, volatile half + volatile half This patch inserts casts to remove volatile type qualifier following volatile loads (fp16 only). CUDA fp16 library headers should add volatile member functions. - Update have_fp16 to include compute 6.1 GPUs, which do support fp16, although their fp16 throughput is low. Updated tests. Signed-off-by: Wei Pan <weip@nvidia.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Do not emit shared etc. as part of type for casting
Fix fp16 reduction kernels with compiler errors:
"no operator "+" matches these operands, volatile half + volatile half
This patch inserts casts to remove volatile type qualifier following
volatile loads (fp16 only). CUDA fp16 library headers should add
volatile member functions.
Signed-off-by: Wei Pan weip@nvidia.com
Thanks for contributing to TVM! Please refer to guideline https://docs.tvm.ai/contribute/ for useful information and tips. After the pull request is submitted, please request code reviews from Reviewers by @ them in the pull request thread.