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 2021/03/02 03:57:31 UTC
[GitHub] [tvm] kongroo opened a new pull request #7561: [Codegen][CUDA] Fix: cuda codegen vectorize cast
kongroo opened a new pull request #7561:
URL: https://github.com/apache/tvm/pull/7561
Data types such as float32x8 and int32x8 are not supported in CUDA, which will result in errors like "Cannot convert type float32x8 to CUDA type!" in code generation. I tried to fix this by storing 2 32-bits values in 1 64-bits value.
Could you please help review this fix? @Laurawly
----------------------------------------------------------------
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] [tvm] kongroo commented on a change in pull request #7561: [Codegen][CUDA] Fix: cuda codegen vectorize cast
Posted by GitBox <gi...@apache.org>.
kongroo commented on a change in pull request #7561:
URL: https://github.com/apache/tvm/pull/7561#discussion_r585348709
##########
File path: tests/python/unittest/test_target_codegen_cuda.py
##########
@@ -511,8 +511,8 @@ def check(t0, t1):
# schedule
s = tvm.te.create_schedule(C.op)
- ob, ib = s[C].split(s[C].op.axis[0], nparts=32)
- _, iib = s[C].split(ib, factor=4)
+ ob, ib = s[C].split(s[C].op.axis[0], nparts=n // factor)
Review comment:
changed
----------------------------------------------------------------
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] [tvm] Laurawly commented on a change in pull request #7561: [Codegen][CUDA] Fix: cuda codegen vectorize cast
Posted by GitBox <gi...@apache.org>.
Laurawly commented on a change in pull request #7561:
URL: https://github.com/apache/tvm/pull/7561#discussion_r585287887
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -238,12 +253,53 @@ void CodeGenCUDA::PrintType(DataType t, std::ostream& os) { // NOLINT(*)
break;
}
}
- case 16:
- os << "short";
- break;
- case 32:
- os << "int";
+ case 16: {
+ if (t.is_scalar()) {
+ os << "short";
+ } else if (t.lanes() <= 4) {
+ os << "short" << lanes;
+ } else if (t.lanes() <= 8) {
+ // Emit CUDA code to access int16 vector elements.
+ //
+ // short4 is stored as int2
+ //
+ // s4.x is emitted as *(short2*)(&(i2.x)).x
+ // s4.y is emitted as *(short2*)(&(i2.x)).y
+ // s4.z is emitted as *(short2*)(&(i2.y)).x
+ // s4.w is emitted as *(short2*)(&(i2.y)).y
+ //
+ ICHECK_EQ(t.lanes() % 2, 0) << "only support even lane for shorT type with lanes > 4";
+ os << "int" << t.lanes() / 2;
+ } else {
+ fail = true;
+ }
+ if (!fail) {
+ return;
+ }
Review comment:
Missing `break` here?
##########
File path: tests/python/unittest/test_target_codegen_cuda.py
##########
@@ -511,8 +511,8 @@ def check(t0, t1):
# schedule
s = tvm.te.create_schedule(C.op)
- ob, ib = s[C].split(s[C].op.axis[0], nparts=32)
- _, iib = s[C].split(ib, factor=4)
+ ob, ib = s[C].split(s[C].op.axis[0], nparts=n // factor)
Review comment:
We can also directly say `factor=factor` here.
----------------------------------------------------------------
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] [tvm] kongroo commented on a change in pull request #7561: [Codegen][CUDA] Fix: cuda codegen vectorize cast
Posted by GitBox <gi...@apache.org>.
kongroo commented on a change in pull request #7561:
URL: https://github.com/apache/tvm/pull/7561#discussion_r585348096
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -238,12 +253,53 @@ void CodeGenCUDA::PrintType(DataType t, std::ostream& os) { // NOLINT(*)
break;
}
}
- case 16:
- os << "short";
- break;
- case 32:
- os << "int";
+ case 16: {
+ if (t.is_scalar()) {
+ os << "short";
+ } else if (t.lanes() <= 4) {
+ os << "short" << lanes;
+ } else if (t.lanes() <= 8) {
+ // Emit CUDA code to access int16 vector elements.
+ //
+ // short4 is stored as int2
+ //
+ // s4.x is emitted as *(short2*)(&(i2.x)).x
+ // s4.y is emitted as *(short2*)(&(i2.x)).y
+ // s4.z is emitted as *(short2*)(&(i2.y)).x
+ // s4.w is emitted as *(short2*)(&(i2.y)).y
+ //
+ ICHECK_EQ(t.lanes() % 2, 0) << "only support even lane for shorT type with lanes > 4";
+ os << "int" << t.lanes() / 2;
+ } else {
+ fail = true;
+ }
+ if (!fail) {
+ return;
+ }
Review comment:
fixed
----------------------------------------------------------------
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] [tvm] vinx13 merged pull request #7561: [Codegen][CUDA] Fix: cuda codegen vectorize cast
Posted by GitBox <gi...@apache.org>.
vinx13 merged pull request #7561:
URL: https://github.com/apache/tvm/pull/7561
----------------------------------------------------------------
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] [tvm] kongroo commented on a change in pull request #7561: [Codegen][CUDA] Fix: cuda codegen vectorize cast
Posted by GitBox <gi...@apache.org>.
kongroo commented on a change in pull request #7561:
URL: https://github.com/apache/tvm/pull/7561#discussion_r585300157
##########
File path: tests/python/unittest/test_target_codegen_cuda.py
##########
@@ -511,8 +511,8 @@ def check(t0, t1):
# schedule
s = tvm.te.create_schedule(C.op)
- ob, ib = s[C].split(s[C].op.axis[0], nparts=32)
- _, iib = s[C].split(ib, factor=4)
+ ob, ib = s[C].split(s[C].op.axis[0], nparts=n // factor)
Review comment:
like this?
```python
ob, ib = s[C].split(s[C].op.axis[0], factor=factor)
# _, iib = s[C].split(ib, factor=factor)
s[C].vectorize(ib)
```
----------------------------------------------------------------
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] [tvm] vinx13 commented on pull request #7561: [Codegen][CUDA] Fix: cuda codegen vectorize cast
Posted by GitBox <gi...@apache.org>.
vinx13 commented on pull request #7561:
URL: https://github.com/apache/tvm/pull/7561#issuecomment-789129433
Thanks @kongroo @Laurawly
----------------------------------------------------------------
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] [tvm] Laurawly commented on a change in pull request #7561: [Codegen][CUDA] Fix: cuda codegen vectorize cast
Posted by GitBox <gi...@apache.org>.
Laurawly commented on a change in pull request #7561:
URL: https://github.com/apache/tvm/pull/7561#discussion_r585300949
##########
File path: tests/python/unittest/test_target_codegen_cuda.py
##########
@@ -511,8 +511,8 @@ def check(t0, t1):
# schedule
s = tvm.te.create_schedule(C.op)
- ob, ib = s[C].split(s[C].op.axis[0], nparts=32)
- _, iib = s[C].split(ib, factor=4)
+ ob, ib = s[C].split(s[C].op.axis[0], nparts=n // factor)
Review comment:
Yeah.
----------------------------------------------------------------
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