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/08/03 10:42:43 UTC
[GitHub] [tvm] echuraev opened a new pull request #8636: Add vectorization to cuda conv2d_nhwc schedule
echuraev opened a new pull request #8636:
URL: https://github.com/apache/tvm/pull/8636
Adding vectorization significantly improved performance. About 6-7x
boost.
Thanks for contributing to TVM! Please refer to guideline https://tvm.apache.org/docs/contribute/ for useful information and tips. After the pull request is submitted, please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @ them in the pull request thread.
--
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] echuraev commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-912803703
@jcf94, @masahi, @tmoreau89 could you please take a look one again on this 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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tmoreau89 commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-925972724
@jcf94 - please update the PR review based on latest comments and updates to the PR. I will dismiss your review in 24hrs so we can unblock this PR given that we haven't heard back from you.
--
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] jwfromm removed a comment on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
jwfromm removed a comment on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-916260974
@jcf94 can you take another look at this PR and see if @echuraev changes look good?
--
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] echuraev commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-911900705
@mbrookhart I wasn't able to reproduce the regression on NVidia 3070. As a possible solution, I can create a separate `conv2d_nhwc` schedule for OpenCL. But I saw that my fix for OpenCL issue with `global_work_size` also works for default schedule on Cuda: https://github.com/apache/tvm/pull/8636/files#diff-05fdfdcbc0bdf86e1df35950ae34877c2f9dbddab6a99ca630582547d4e7e0faL88-L89
| | Results on NVidia 3090 (1024 trials per kernel) | Results on NVidia 3070 Mobile (512 trials per kernel) |
|----------------------------|-------------------------------------------------|-------------------------------------------------------|
| Main (AutoTVM) | 0.17 ms | 0.4504 ms |
| Main (Ansor) | 0.18 ms | 0.4284 ms |
| Code from the PR (AutoTVM) | 0.26 ms | 0.4200 ms |
| Code from the PR (Ansor) | 0.19 ms | 0.4200 ms |
--
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] echuraev commented on a change in pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev commented on a change in pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#discussion_r693947325
##########
File path: python/tvm/topi/gpu/conv2d_nhwc.py
##########
@@ -85,15 +87,17 @@ def schedule_conv2d_nhwc_direct(cfg, s, Conv):
thread_yz = te.thread_axis((0, vthread_n), "vthread", name="vy")
# Schedule for output
- ni, hi, wi, fi = s[output].op.axis
- bz = s[output].fuse(hi, wi)
+ ni, _, wi, fi = s[output].op.axis
+ bz = wi
+ fi, vec = s[output].split(fi, factor=vec_factor)
+ s[output].vectorize(vec)
Review comment:
Sorry for the late reply, I was on vacation. Thank you for your question @masahi!
Today in the latest commit I added vectorization to the conv2d inner loop. The diff of generated OpenCL kernel you can see below:
```diff
__kernel void my_conv_kernel0(__global float* restrict inp, __global float* restrict w, __global float* restrict Conv2dOutput) {
float Conv2dOutput_local[4];
__local float PaddedInput_shared[24];
__local float w_shared[256];
float PaddedInput_shared_local[1];
float w_shared_local[4];
for (int yy = 0; yy < 298; ++yy) {
- for (int ff_c_init = 0; ff_c_init < 4; ++ff_c_init) {
- Conv2dOutput_local[(ff_c_init)] = 0.000000e+00f;
- }
+ vstore4(((float4)(0.000000e+00f, 0.000000e+00f, 0.000000e+00f, 0.000000e+00f)), 0, Conv2dOutput_local + 0);
for (int rc_outer = 0; rc_outer < 2; ++rc_outer) {
for (int ry = 0; ry < 3; ++ry) {
for (int rx = 0; rx < 3; ++rx) {
barrier(CLK_LOCAL_MEM_FENCE);
PaddedInput_shared[(((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))))] = inp[((((((((yy * 9600) + (ry * 9600)) + (((int)get_group_id(2)) * 32)) + (rx * 32)) + (rc_outer * 16)) + (((int)get_local_id(1)) * 4)) + ((int)get_local_id(0))))];
for (int ax2_ax3_fused_outer_outer_outer = 0; ax2_ax3_fused_outer_outer_outer < 8; ++ax2_ax3_fused_outer_outer_outer) {
vstore2(vload2(0, w + (((((((ry * 3072) + (rx * 1024)) + (rc_outer * 512)) + (ax2_ax3_fused_outer_outer_outer * 64)) + ((((((int)get_local_id(1)) * 8) + (((int)get_local_id(0)) * 2)) >> 4) * 32)) + (((int)get_group_id(0)) * 16)) + (((((int)get_local_id(1)) * 8) + (((int)get_local_id(0)) * 2)) & 15))), 0, w_shared + (((ax2_ax3_fused_outer_outer_outer * 32) + (((int)get_local_id(1)) * 8)) + (((int)get_local_id(0)) * 2)));
}
barrier(CLK_LOCAL_MEM_FENCE);
for (int rc_inner = 0; rc_inner < 16; ++rc_inner) {
if (((int)get_local_id(1)) < 1) {
PaddedInput_shared_local[(0)] = PaddedInput_shared[(((((int)get_local_id(1)) * 24) + rc_inner))];
}
for (int ax3 = 0; ax3 < 4; ++ax3) {
w_shared_local[(ax3)] = w_shared[((((rc_inner * 16) + (((int)get_local_id(0)) * 4)) + ax3))];
}
- for (int ff_c = 0; ff_c < 4; ++ff_c) {
- if (((int)get_local_id(1)) < 1) {
- Conv2dOutput_local[(ff_c)] = (Conv2dOutput_local[(ff_c)] + (PaddedInput_shared_local[(0)] * w_shared_local[(ff_c)]));
- }
+ if (((int)get_local_id(1)) < 1) {
+ vstore4((vload4(0, Conv2dOutput_local + 0) + (((float4)(PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)])) * vload4(0, w_shared_local + 0))), 0, Conv2dOutput_local + 0);
}
}
}
}
}
for (int ff_outer_inner = 0; ff_outer_inner < 2; ++ff_outer_inner) {
if (((int)get_local_id(1)) < 1) {
vstore2(vload2(0, Conv2dOutput_local + (ff_outer_inner * 2)), 0, Conv2dOutput + ((((((((int)get_local_id(1)) * 2841728) + (yy * 9536)) + (((int)get_group_id(2)) * 32)) + (((int)get_group_id(0)) * 16)) + (((int)get_local_id(0)) * 4)) + (ff_outer_inner * 2)));
}
}
}
}
```
With this vectorization, the execution time didn't change in comparison with previous code generation.
What about performance boost. First, let me share my performance numbers (the numbers are average execution time in 10 runs) which I got today on the Samsung Galaxy A71:
- Original code (w/o any changes): 1225.95 ms
- Code with bug fix (read about bug fix below): 1383.18 ms
- With vectorization: 189.36 ms
*Bug fix.* In the first commit, I also fixed one accuracy problem in OpenCL. [Here](https://github.com/apache/tvm/pull/8636/commits/0a18676916dd8656110f8a27ab0fe6be154136db#diff-194f66fb68e225fb63696eabb5aac5f3e30357c8e8bb25d33e2713eb5de322aaL88-L89) after fusing `hi` and `wi` it was possible that the value of OpenCL `global_work_size` for `clEnqueueNDRangeKernel` was too high and some values in output tensor were not calculated.
> I wonder where 6-7x perf improvement comes from?
To answer on this question, let's compare the generated OpenCL code for version with bug fix and with the latest code:
```diff
__kernel void my_conv_kernel0(__global float* restrict inp, __global float* restrict w, __global float* restrict Conv2dOutput) {
- float Conv2dOutput_local[2];
+ float Conv2dOutput_local[4];
__local float PaddedInput_shared[24];
__local float w_shared[256];
float PaddedInput_shared_local[1];
- float w_shared_local[2];
+ float w_shared_local[4];
for (int yy = 0; yy < 298; ++yy) {
- for (int ff_c_init = 0; ff_c_init < 2; ++ff_c_init) {
- Conv2dOutput_local[(ff_c_init)] = 0.000000e+00f;
- }
+ vstore4(((float4)(0.000000e+00f, 0.000000e+00f, 0.000000e+00f, 0.000000e+00f)), 0, Conv2dOutput_local + 0);
for (int rc_outer = 0; rc_outer < 2; ++rc_outer) {
for (int ry = 0; ry < 3; ++ry) {
for (int rx = 0; rx < 3; ++rx) {
barrier(CLK_LOCAL_MEM_FENCE);
PaddedInput_shared[(((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))))] = inp[((((((((yy * 9600) + (ry * 9600)) + (((int)get_group_id(2)) * 32)) + (rx * 32)) + (rc_outer * 16)) + (((int)get_local_id(1)) * 4)) + ((int)get_local_id(0))))];
- for (int ax2_ax3_fused_outer_outer = 0; ax2_ax3_fused_outer_outer < 8; ++ax2_ax3_fused_outer_outer) {
- w_shared[((((ax2_ax3_fused_outer_outer * 32) + ((((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))) >> 3) * 16)) + (((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))) & 7)))] = w[((((((((ry * 3072) + (rx * 1024)) + (rc_outer * 512)) + (ax2_ax3_fused_outer_outer * 64)) + ((((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))) >> 3) * 32)) + (((int)get_group_id(0)) * 8)) + (((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))) & 7)))];
+ for (int ax2_ax3_fused_outer_outer_outer = 0; ax2_ax3_fused_outer_outer_outer < 8; ++ax2_ax3_fused_outer_outer_outer) {
+ vstore2(vload2(0, w + (((((((ry * 3072) + (rx * 1024)) + (rc_outer * 512)) + (ax2_ax3_fused_outer_outer_outer * 64)) + ((((((int)get_local_id(1)) * 8) + (((int)get_local_id(0)) * 2)) >> 4) * 32)) + (((int)get_group_id(0)) * 16)) + (((((int)get_local_id(1)) * 8) + (((int)get_local_id(0)) * 2)) & 15))), 0, w_shared + (((ax2_ax3_fused_outer_outer_outer * 32) + (((int)get_local_id(1)) * 8)) + (((int)get_local_id(0)) * 2)));
}
barrier(CLK_LOCAL_MEM_FENCE);
for (int rc_inner = 0; rc_inner < 16; ++rc_inner) {
if (((int)get_local_id(1)) < 1) {
PaddedInput_shared_local[(0)] = PaddedInput_shared[(((((int)get_local_id(1)) * 24) + rc_inner))];
}
- for (int ax3 = 0; ax3 < 2; ++ax3) {
- w_shared_local[(ax3)] = w_shared[((((rc_inner * 16) + (((int)get_local_id(0)) * 2)) + ax3))];
+ for (int ax3 = 0; ax3 < 4; ++ax3) {
+ w_shared_local[(ax3)] = w_shared[((((rc_inner * 16) + (((int)get_local_id(0)) * 4)) + ax3))];
}
- for (int ff_c = 0; ff_c < 2; ++ff_c) {
- if (((int)get_local_id(1)) < 1) {
- Conv2dOutput_local[(ff_c)] = (Conv2dOutput_local[(ff_c)] + (PaddedInput_shared_local[(0)] * w_shared_local[(ff_c)]));
- }
+ if (((int)get_local_id(1)) < 1) {
+ vstore4((vload4(0, Conv2dOutput_local + 0) + (((float4)(PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)])) * vload4(0, w_shared_local + 0))), 0, Conv2dOutput_local + 0);
}
}
}
}
}
- for (int ff_inner = 0; ff_inner < 2; ++ff_inner) {
+ for (int ff_outer_inner = 0; ff_outer_inner < 2; ++ff_outer_inner) {
if (((int)get_local_id(1)) < 1) {
- Conv2dOutput[(((((((((int)get_local_id(1)) * 2841728) + (yy * 9536)) + (((int)get_group_id(2)) * 32)) + (((int)get_group_id(0)) * 8)) + (((int)get_local_id(0)) * 2)) + ff_inner))] = Conv2dOutput_local[(ff_inner)];
+ vstore2(vload2(0, Conv2dOutput_local + (ff_outer_inner * 2)), 0, Conv2dOutput + ((((((((int)get_local_id(1)) * 2841728) + (yy * 9536)) + (((int)get_group_id(2)) * 32)) + (((int)get_group_id(0)) * 16)) + (((int)get_local_id(0)) * 4)) + (ff_outer_inner * 2)));
}
}
}
}
```
I suppose that the performance boost is connected with decreasing memory latency. We read more data in one execution unit and store them in vector data types.
--
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] jwfromm commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
jwfromm commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-916260974
@jcf94 can you take another look at this PR and see if @echuraev changes look good?
--
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 #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-904538888
Thanks @echuraev, this is really interesting. With the new commit, I was able to generate a packed float16x2 instruction for AMD Vega or newer GPUs, which is supposed to be 2x faster than normal fp32 or fp16 instructions https://gist.github.com/masahi/2de1a7dc87e2068ffb50ba6135273f95#file-conv2d_nhwc_float16x2-s-L495-L496. I'm going to experiment further.
--
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] echuraev commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-894651590
@jcf94 also, I suppose that we should update statistic in tophub and rename `conv2d_nhwc.cuda` to `conv2d_nhwc.gpu` in all files. Am I right?
--
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] echuraev edited a comment on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev edited a comment on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-912897229
> So it appears @mbrookhart has encountered a regression and that regression wasn't reproduced by @echuraev on a slightly different GPU. The possible solution in order to be extra safe would be to introduce a new set of TOPI schedules for vectorized GPU operator implementations. WDYT @jwfromm and @mbrookhart
Sorry, my mistake. I didn't write an update. @mbrookhart has remeasured autotvm tuning on the latest code from my branch and he got the same numbers as he got on the main branch. Now the table can be updated:
| | Results on NVidia 3090 (1024 trials per kernel) | Results on NVidia 3070 Mobile (512 trials per kernel) |
|----------------------------|-------------------------------------------------|-------------------------------------------------------|
| Main (AutoTVM) | 0.17 ms | 0.4504 ms |
| Main (Ansor) | 0.18 ms | 0.4284 ms |
| Code from the PR (AutoTVM) | 0.173 ms | 0.4200 ms |
| Code from the PR (Ansor) | 0.19 ms | 0.4200 ms |
So, we don't have regression on the latest code.
--
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] jwfromm commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
jwfromm commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-921037478
@jcf94 can you please take another look at this? It's currently being blocked by your request for changes.
--
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] jwfromm commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
jwfromm commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-931458419
Thanks @echuraev, this is now merged.
--
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] echuraev commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-922881575
@jcf94 Could you please take a look one again?
--
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] tmoreau89 commented on a change in pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on a change in pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#discussion_r695005091
##########
File path: python/tvm/topi/gpu/conv2d_nhwc.py
##########
@@ -85,15 +87,17 @@ def schedule_conv2d_nhwc_direct(cfg, s, Conv):
thread_yz = te.thread_axis((0, vthread_n), "vthread", name="vy")
# Schedule for output
- ni, hi, wi, fi = s[output].op.axis
- bz = s[output].fuse(hi, wi)
+ ni, _, wi, fi = s[output].op.axis
+ bz = wi
+ fi, vec = s[output].split(fi, factor=vec_factor)
+ s[output].vectorize(vec)
Review comment:
Thanks @echuraev for providing details on the performance gains here, quite a drastic improvement 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.
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 #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-895541169
Yes, I don't believe optimization like this applies to all GPU backends. I think it's time to have more fine-grained schedules for each GPU backend, rather than one schedule shared by all. Also cc @csullivan I think he has a schedule for Adreno specifically.
> would this also benefit Vulkan FP16 codegen?
I'm also interested in this question. I'll take a look at what kind of asm it generates.
--
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] echuraev commented on a change in pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev commented on a change in pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#discussion_r684124511
##########
File path: python/tvm/topi/cuda/conv2d_nhwc.py
##########
@@ -85,8 +87,10 @@ def schedule_conv2d_nhwc_direct(cfg, s, Conv):
thread_yz = te.thread_axis((0, vthread_n), "vthread", name="vy")
# Schedule for output
- ni, hi, wi, fi = s[output].op.axis
- bz = s[output].fuse(hi, wi)
+ ni, _, wi, fi = s[output].op.axis
+ bz = wi
+ fi, vec = s[output].split(fi, factor=vec_factor)
+ s[output].vectorize(vec)
Review comment:
Yes, it works. I added vectorize to the back of reorder list. Sorry, I'm new in writing schedules, so, maybe I did something wrong :) I checked the performance and accuracy on an OpenCL device, and it works fine and fast.
--
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] echuraev commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-912897229
> So it appears @mbrookhart has encountered a regression and that regression wasn't reproduced by @echuraev on a slightly different GPU. The possible solution in order to be extra safe would be to introduce a new set of TOPI schedules for vectorized GPU operator implementations. WDYT @jwfromm and @mbrookhart
Sorry, my mistake. I didn't write an update. @mbrookhart has remeasured autotvm tuning on the latest code from my branch and he got the same numbers as he got on the main branch. Now the table can be updated:
| | Results on NVidia 3090 (1024 trials per kernel) | Results on NVidia 3070 Mobile (512 trials per kernel) |
|----------------------------|-------------------------------------------------|-------------------------------------------------------|
| Main (AutoTVM) | 0.17 ms | 0.4504 ms |
| Main (Ansor) | 0.18 ms | 0.4284 ms |
| Code from the PR (AutoTVM) | 0.173 ms | 0.4200 ms |
| Code from the PR (Ansor) | 0.19 ms | 0.4200 ms |
--
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] echuraev commented on a change in pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev commented on a change in pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#discussion_r693947325
##########
File path: python/tvm/topi/gpu/conv2d_nhwc.py
##########
@@ -85,15 +87,17 @@ def schedule_conv2d_nhwc_direct(cfg, s, Conv):
thread_yz = te.thread_axis((0, vthread_n), "vthread", name="vy")
# Schedule for output
- ni, hi, wi, fi = s[output].op.axis
- bz = s[output].fuse(hi, wi)
+ ni, _, wi, fi = s[output].op.axis
+ bz = wi
+ fi, vec = s[output].split(fi, factor=vec_factor)
+ s[output].vectorize(vec)
Review comment:
Sorry for the late reply, I was on vacation. Thank you for your question @masahi!
Today in the latest commit I added vectorization to the conv2d inner loop. The diff of generated OpenCL kernel you can see below:
```diff
__kernel void my_conv_kernel0(__global float* restrict inp, __global float* restrict w, __global float* restrict Conv2dOutput) {
float Conv2dOutput_local[4];
__local float PaddedInput_shared[24];
__local float w_shared[256];
float PaddedInput_shared_local[1];
float w_shared_local[4];
for (int yy = 0; yy < 298; ++yy) {
- for (int ff_c_init = 0; ff_c_init < 4; ++ff_c_init) {
- Conv2dOutput_local[(ff_c_init)] = 0.000000e+00f;
- }
+ vstore4(((float4)(0.000000e+00f, 0.000000e+00f, 0.000000e+00f, 0.000000e+00f)), 0, Conv2dOutput_local + 0);
for (int rc_outer = 0; rc_outer < 2; ++rc_outer) {
for (int ry = 0; ry < 3; ++ry) {
for (int rx = 0; rx < 3; ++rx) {
barrier(CLK_LOCAL_MEM_FENCE);
PaddedInput_shared[(((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))))] = inp[((((((((yy * 9600) + (ry * 9600)) + (((int)get_group_id(2)) * 32)) + (rx * 32)) + (rc_outer * 16)) + (((int)get_local_id(1)) * 4)) + ((int)get_local_id(0))))];
for (int ax2_ax3_fused_outer_outer_outer = 0; ax2_ax3_fused_outer_outer_outer < 8; ++ax2_ax3_fused_outer_outer_outer) {
vstore2(vload2(0, w + (((((((ry * 3072) + (rx * 1024)) + (rc_outer * 512)) + (ax2_ax3_fused_outer_outer_outer * 64)) + ((((((int)get_local_id(1)) * 8) + (((int)get_local_id(0)) * 2)) >> 4) * 32)) + (((int)get_group_id(0)) * 16)) + (((((int)get_local_id(1)) * 8) + (((int)get_local_id(0)) * 2)) & 15))), 0, w_shared + (((ax2_ax3_fused_outer_outer_outer * 32) + (((int)get_local_id(1)) * 8)) + (((int)get_local_id(0)) * 2)));
}
barrier(CLK_LOCAL_MEM_FENCE);
for (int rc_inner = 0; rc_inner < 16; ++rc_inner) {
if (((int)get_local_id(1)) < 1) {
PaddedInput_shared_local[(0)] = PaddedInput_shared[(((((int)get_local_id(1)) * 24) + rc_inner))];
}
for (int ax3 = 0; ax3 < 4; ++ax3) {
w_shared_local[(ax3)] = w_shared[((((rc_inner * 16) + (((int)get_local_id(0)) * 4)) + ax3))];
}
- for (int ff_c = 0; ff_c < 4; ++ff_c) {
- if (((int)get_local_id(1)) < 1) {
- Conv2dOutput_local[(ff_c)] = (Conv2dOutput_local[(ff_c)] + (PaddedInput_shared_local[(0)] * w_shared_local[(ff_c)]));
- }
+ if (((int)get_local_id(1)) < 1) {
+ vstore4((vload4(0, Conv2dOutput_local + 0) + (((float4)(PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)])) * vload4(0, w_shared_local + 0))), 0, Conv2dOutput_local + 0);
}
}
}
}
}
for (int ff_outer_inner = 0; ff_outer_inner < 2; ++ff_outer_inner) {
if (((int)get_local_id(1)) < 1) {
vstore2(vload2(0, Conv2dOutput_local + (ff_outer_inner * 2)), 0, Conv2dOutput + ((((((((int)get_local_id(1)) * 2841728) + (yy * 9536)) + (((int)get_group_id(2)) * 32)) + (((int)get_group_id(0)) * 16)) + (((int)get_local_id(0)) * 4)) + (ff_outer_inner * 2)));
}
}
}
}
```
With this vectorization, the execution time didn't change in comparison with previous code generation.
What about performance boost. First, let me share my performance numbers (the numbers are average execution time in 10 runs) which I got today on the Samsung Galaxy A71:
- Original code (w/o any changes): 1225.95 ms
- Code with bug fix (read about bug fix below): 1383.18 ms
- With vectorization: 189.36 ms
**Bug fix.** In the first commit, I also fixed one accuracy problem in OpenCL. [Here](https://github.com/apache/tvm/pull/8636/commits/0a18676916dd8656110f8a27ab0fe6be154136db#diff-194f66fb68e225fb63696eabb5aac5f3e30357c8e8bb25d33e2713eb5de322aaL88-L89) after fusing `hi` and `wi` it was possible that the value of OpenCL `global_work_size` for `clEnqueueNDRangeKernel` was too high and some values in output tensor were not calculated.
> I wonder where 6-7x perf improvement comes from?
To answer on this question, let's compare the generated OpenCL code for version with bug fix and with the latest code:
```diff
__kernel void my_conv_kernel0(__global float* restrict inp, __global float* restrict w, __global float* restrict Conv2dOutput) {
- float Conv2dOutput_local[2];
+ float Conv2dOutput_local[4];
__local float PaddedInput_shared[24];
__local float w_shared[256];
float PaddedInput_shared_local[1];
- float w_shared_local[2];
+ float w_shared_local[4];
for (int yy = 0; yy < 298; ++yy) {
- for (int ff_c_init = 0; ff_c_init < 2; ++ff_c_init) {
- Conv2dOutput_local[(ff_c_init)] = 0.000000e+00f;
- }
+ vstore4(((float4)(0.000000e+00f, 0.000000e+00f, 0.000000e+00f, 0.000000e+00f)), 0, Conv2dOutput_local + 0);
for (int rc_outer = 0; rc_outer < 2; ++rc_outer) {
for (int ry = 0; ry < 3; ++ry) {
for (int rx = 0; rx < 3; ++rx) {
barrier(CLK_LOCAL_MEM_FENCE);
PaddedInput_shared[(((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))))] = inp[((((((((yy * 9600) + (ry * 9600)) + (((int)get_group_id(2)) * 32)) + (rx * 32)) + (rc_outer * 16)) + (((int)get_local_id(1)) * 4)) + ((int)get_local_id(0))))];
- for (int ax2_ax3_fused_outer_outer = 0; ax2_ax3_fused_outer_outer < 8; ++ax2_ax3_fused_outer_outer) {
- w_shared[((((ax2_ax3_fused_outer_outer * 32) + ((((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))) >> 3) * 16)) + (((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))) & 7)))] = w[((((((((ry * 3072) + (rx * 1024)) + (rc_outer * 512)) + (ax2_ax3_fused_outer_outer * 64)) + ((((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))) >> 3) * 32)) + (((int)get_group_id(0)) * 8)) + (((((int)get_local_id(1)) * 4) + ((int)get_local_id(0))) & 7)))];
+ for (int ax2_ax3_fused_outer_outer_outer = 0; ax2_ax3_fused_outer_outer_outer < 8; ++ax2_ax3_fused_outer_outer_outer) {
+ vstore2(vload2(0, w + (((((((ry * 3072) + (rx * 1024)) + (rc_outer * 512)) + (ax2_ax3_fused_outer_outer_outer * 64)) + ((((((int)get_local_id(1)) * 8) + (((int)get_local_id(0)) * 2)) >> 4) * 32)) + (((int)get_group_id(0)) * 16)) + (((((int)get_local_id(1)) * 8) + (((int)get_local_id(0)) * 2)) & 15))), 0, w_shared + (((ax2_ax3_fused_outer_outer_outer * 32) + (((int)get_local_id(1)) * 8)) + (((int)get_local_id(0)) * 2)));
}
barrier(CLK_LOCAL_MEM_FENCE);
for (int rc_inner = 0; rc_inner < 16; ++rc_inner) {
if (((int)get_local_id(1)) < 1) {
PaddedInput_shared_local[(0)] = PaddedInput_shared[(((((int)get_local_id(1)) * 24) + rc_inner))];
}
- for (int ax3 = 0; ax3 < 2; ++ax3) {
- w_shared_local[(ax3)] = w_shared[((((rc_inner * 16) + (((int)get_local_id(0)) * 2)) + ax3))];
+ for (int ax3 = 0; ax3 < 4; ++ax3) {
+ w_shared_local[(ax3)] = w_shared[((((rc_inner * 16) + (((int)get_local_id(0)) * 4)) + ax3))];
}
- for (int ff_c = 0; ff_c < 2; ++ff_c) {
- if (((int)get_local_id(1)) < 1) {
- Conv2dOutput_local[(ff_c)] = (Conv2dOutput_local[(ff_c)] + (PaddedInput_shared_local[(0)] * w_shared_local[(ff_c)]));
- }
+ if (((int)get_local_id(1)) < 1) {
+ vstore4((vload4(0, Conv2dOutput_local + 0) + (((float4)(PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)], PaddedInput_shared_local[(0)])) * vload4(0, w_shared_local + 0))), 0, Conv2dOutput_local + 0);
}
}
}
}
}
- for (int ff_inner = 0; ff_inner < 2; ++ff_inner) {
+ for (int ff_outer_inner = 0; ff_outer_inner < 2; ++ff_outer_inner) {
if (((int)get_local_id(1)) < 1) {
- Conv2dOutput[(((((((((int)get_local_id(1)) * 2841728) + (yy * 9536)) + (((int)get_group_id(2)) * 32)) + (((int)get_group_id(0)) * 8)) + (((int)get_local_id(0)) * 2)) + ff_inner))] = Conv2dOutput_local[(ff_inner)];
+ vstore2(vload2(0, Conv2dOutput_local + (ff_outer_inner * 2)), 0, Conv2dOutput + ((((((((int)get_local_id(1)) * 2841728) + (yy * 9536)) + (((int)get_group_id(2)) * 32)) + (((int)get_group_id(0)) * 16)) + (((int)get_local_id(0)) * 4)) + (ff_outer_inner * 2)));
}
}
}
}
```
I suppose that the performance boost is connected with decreasing memory latency. We read more data in one execution unit and store them in vector data types.
--
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 #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-915938295
I think it is good to go as long as there would be no concern for regression. Does the new schedule subsume the old schedule in its search space?
--
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] tmoreau89 commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-912834612
So it appears @mbrookhart has encountered a regression and that regression wasn't reproduced by @echuraev on a slightly different GPU. The possible solution in order to be extra safe would be to introduce a new set of TOPI schedules for vectorized GPU operator implementations. WDYT @jwfromm and @mbrookhart
--
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 a change in pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#discussion_r688790802
##########
File path: python/tvm/topi/gpu/conv2d_nhwc.py
##########
@@ -85,15 +87,17 @@ def schedule_conv2d_nhwc_direct(cfg, s, Conv):
thread_yz = te.thread_axis((0, vthread_n), "vthread", name="vy")
# Schedule for output
- ni, hi, wi, fi = s[output].op.axis
- bz = s[output].fuse(hi, wi)
+ ni, _, wi, fi = s[output].op.axis
+ bz = wi
+ fi, vec = s[output].split(fi, factor=vec_factor)
+ s[output].vectorize(vec)
Review comment:
Is this supposed to vectorize the conv2d inner loop? Based on generated code, I think it only vectorize the last stage, which can be copying local to global me or fused activation computation. I wonder where 6-7x perf improvement comes from?
Here is an example of generated code where `vec_factor` is fixed to 2.
```extern "C" __global__ void __launch_bounds__(16) default_function_kernel0(float* __restrict__ A, float* __restrict__ W, float* __restrict__ Conv2dOutput) {
float Conv2dOutput_local[8];
__shared__ float PaddedInput_shared[96];
__shared__ float W_shared[256];
float PaddedInput_shared_local[2];
float W_shared_local[4];
for (int yy = 0; yy < 8; ++yy) {
for (int nn_c_init = 0; nn_c_init < 2; ++nn_c_init) {
for (int ff_c_init = 0; ff_c_init < 4; ++ff_c_init) {
Conv2dOutput_local[(((nn_c_init * 4) + ff_c_init))] = 0.000000e+00f;
}
}
for (int rc_outer = 0; rc_outer < 8; ++rc_outer) {
for (int ry = 0; ry < 5; ++ry) {
for (int rx = 0; rx < 5; ++rx) {
__syncthreads();
for (int ax0_ax3_fused_outer_outer = 0; ax0_ax3_fused_outer_outer < 4; ++ax0_ax3_fused_outer_outer) {
PaddedInput_shared[((((ax0_ax3_fused_outer_outer * 24) + (((int)threadIdx.y) * 4)) + ((int)threadIdx.x)))] = (((((2 <= ((yy * 2) + ry)) && (((yy * 2) + ry) < 18)) && (2 <= ((((int)blockIdx.z) * 2) + rx))) && (((((int)blockIdx.z) * 2) + rx) < 18)) ? A[((((((((((ax0_ax3_fused_outer_outer * 32768) + (yy * 4096)) + (ry * 2048)) + (((int)blockIdx.z) * 256)) + (rx * 128)) + (rc_outer * 16)) + (((int)threadIdx.y) * 4)) + ((int)threadIdx.x)) - 4352))] : 0.000000e+00f);
}
for (int ax2_ax3_fused_outer_outer_outer = 0; ax2_ax3_fused_outer_outer_outer < 8; ++ax2_ax3_fused_outer_outer_outer) {
((float2*)(W_shared + ((((ax2_ax3_fused_outer_outer_outer * 32) + (((int)threadIdx.y) * 8)) + (((int)threadIdx.x) * 2)))))[0] = ((float2*)(W + ((((((((ry * 81920) + (rx * 16384)) + (rc_outer * 2048)) + (ax2_ax3_fused_outer_outer_outer * 256)) + ((((((int)threadIdx.y) * 8) + (((int)threadIdx.x) * 2)) >> 4) * 128)) + (((int)blockIdx.x) * 16)) + (((((int)threadIdx.y) * 8) + (((int)threadIdx.x) * 2)) & 15)))))[0];
}
__syncthreads();
for (int rc_inner = 0; rc_inner < 16; ++rc_inner) {
for (int ax0 = 0; ax0 < 2; ++ax0) {
if (((((int)threadIdx.y) * 2) + ax0) < 4) {
PaddedInput_shared_local[(ax0)] = PaddedInput_shared[((((((int)threadIdx.y) * 48) + (ax0 * 24)) + rc_inner))];
}
}
for (int ax3 = 0; ax3 < 4; ++ax3) {
W_shared_local[(ax3)] = W_shared[((((rc_inner * 16) + (((int)threadIdx.x) * 4)) + ax3))];
}
for (int nn_c = 0; nn_c < 2; ++nn_c) {
for (int ff_c = 0; ff_c < 4; ++ff_c) {
if (((((int)threadIdx.y) * 2) + nn_c) < 4) {
Conv2dOutput_local[(((nn_c * 4) + ff_c))] = (Conv2dOutput_local[(((nn_c * 4) + ff_c))] + (PaddedInput_shared_local[(nn_c)] * W_shared_local[(ff_c)]));
}
}
}
}
}
}
}
for (int nn_inner = 0; nn_inner < 2; ++nn_inner) {
for (int ff_outer_inner = 0; ff_outer_inner < 2; ++ff_outer_inner) {
if (((((int)threadIdx.y) * 2) + nn_inner) < 4) {
if (((int)threadIdx.y) < 2) {
((float2*)(Conv2dOutput + ((((((((((int)threadIdx.y) * 16384) + (nn_inner * 8192)) + (yy * 1024)) + (((int)blockIdx.z) * 128)) + (((int)blockIdx.x) * 16)) + (((int)threadIdx.x) * 4)) + (ff_outer_inner * 2)))))[0] = ((float2*)(Conv2dOutput_local + (((nn_inner * 4) + (ff_outer_inner * 2)))))[0];
}
}
}
}
}
}
```
--
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] tmoreau89 commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-895535130
CC-ing @masahi - would this also benefit Vulkan FP16 codegen? And in general do we want to maintain TOPI definitions for both CUDA and GPU, or rename it all to GPU, to show that it encompasses more than just Nvidia backends? That would I believe be the topic of a larger RFC.
@echuraev could you confirm that making those changes to the schedule templates doesn't lead to regressions on the CUDA codegen on Nvidia GPUs? Thank you!
--
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] echuraev commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-915951069
> I think it is good to go as long as there would be no concern for regression. Does the new schedule subsume the old schedule in its search space?
Yes, it does. When vectorization factor is equal to 1, then generated code should be the same as it was with old schedule.
--
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] tmoreau89 commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-904785071
CC @mei-ye on the above comment - investigating if using vector types can lead to improved FP16 codegen on AMD GPUs
--
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 a change in pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#discussion_r688790802
##########
File path: python/tvm/topi/gpu/conv2d_nhwc.py
##########
@@ -85,15 +87,17 @@ def schedule_conv2d_nhwc_direct(cfg, s, Conv):
thread_yz = te.thread_axis((0, vthread_n), "vthread", name="vy")
# Schedule for output
- ni, hi, wi, fi = s[output].op.axis
- bz = s[output].fuse(hi, wi)
+ ni, _, wi, fi = s[output].op.axis
+ bz = wi
+ fi, vec = s[output].split(fi, factor=vec_factor)
+ s[output].vectorize(vec)
Review comment:
Is this supposed to vectorize the conv2d inner loop? Based on generated code, I think it only vectorize the last stage, which can be copying local to global mem or fused activation computation. I wonder where 6-7x perf improvement comes from?
Here is an example of generated code where `vec_factor` is fixed to 2.
```extern "C" __global__ void __launch_bounds__(16) default_function_kernel0(float* __restrict__ A, float* __restrict__ W, float* __restrict__ Conv2dOutput) {
float Conv2dOutput_local[8];
__shared__ float PaddedInput_shared[96];
__shared__ float W_shared[256];
float PaddedInput_shared_local[2];
float W_shared_local[4];
for (int yy = 0; yy < 8; ++yy) {
for (int nn_c_init = 0; nn_c_init < 2; ++nn_c_init) {
for (int ff_c_init = 0; ff_c_init < 4; ++ff_c_init) {
Conv2dOutput_local[(((nn_c_init * 4) + ff_c_init))] = 0.000000e+00f;
}
}
for (int rc_outer = 0; rc_outer < 8; ++rc_outer) {
for (int ry = 0; ry < 5; ++ry) {
for (int rx = 0; rx < 5; ++rx) {
__syncthreads();
for (int ax0_ax3_fused_outer_outer = 0; ax0_ax3_fused_outer_outer < 4; ++ax0_ax3_fused_outer_outer) {
PaddedInput_shared[((((ax0_ax3_fused_outer_outer * 24) + (((int)threadIdx.y) * 4)) + ((int)threadIdx.x)))] = (((((2 <= ((yy * 2) + ry)) && (((yy * 2) + ry) < 18)) && (2 <= ((((int)blockIdx.z) * 2) + rx))) && (((((int)blockIdx.z) * 2) + rx) < 18)) ? A[((((((((((ax0_ax3_fused_outer_outer * 32768) + (yy * 4096)) + (ry * 2048)) + (((int)blockIdx.z) * 256)) + (rx * 128)) + (rc_outer * 16)) + (((int)threadIdx.y) * 4)) + ((int)threadIdx.x)) - 4352))] : 0.000000e+00f);
}
for (int ax2_ax3_fused_outer_outer_outer = 0; ax2_ax3_fused_outer_outer_outer < 8; ++ax2_ax3_fused_outer_outer_outer) {
((float2*)(W_shared + ((((ax2_ax3_fused_outer_outer_outer * 32) + (((int)threadIdx.y) * 8)) + (((int)threadIdx.x) * 2)))))[0] = ((float2*)(W + ((((((((ry * 81920) + (rx * 16384)) + (rc_outer * 2048)) + (ax2_ax3_fused_outer_outer_outer * 256)) + ((((((int)threadIdx.y) * 8) + (((int)threadIdx.x) * 2)) >> 4) * 128)) + (((int)blockIdx.x) * 16)) + (((((int)threadIdx.y) * 8) + (((int)threadIdx.x) * 2)) & 15)))))[0];
}
__syncthreads();
for (int rc_inner = 0; rc_inner < 16; ++rc_inner) {
for (int ax0 = 0; ax0 < 2; ++ax0) {
if (((((int)threadIdx.y) * 2) + ax0) < 4) {
PaddedInput_shared_local[(ax0)] = PaddedInput_shared[((((((int)threadIdx.y) * 48) + (ax0 * 24)) + rc_inner))];
}
}
for (int ax3 = 0; ax3 < 4; ++ax3) {
W_shared_local[(ax3)] = W_shared[((((rc_inner * 16) + (((int)threadIdx.x) * 4)) + ax3))];
}
for (int nn_c = 0; nn_c < 2; ++nn_c) {
for (int ff_c = 0; ff_c < 4; ++ff_c) {
if (((((int)threadIdx.y) * 2) + nn_c) < 4) {
Conv2dOutput_local[(((nn_c * 4) + ff_c))] = (Conv2dOutput_local[(((nn_c * 4) + ff_c))] + (PaddedInput_shared_local[(nn_c)] * W_shared_local[(ff_c)]));
}
}
}
}
}
}
}
for (int nn_inner = 0; nn_inner < 2; ++nn_inner) {
for (int ff_outer_inner = 0; ff_outer_inner < 2; ++ff_outer_inner) {
if (((((int)threadIdx.y) * 2) + nn_inner) < 4) {
if (((int)threadIdx.y) < 2) {
((float2*)(Conv2dOutput + ((((((((((int)threadIdx.y) * 16384) + (nn_inner * 8192)) + (yy * 1024)) + (((int)blockIdx.z) * 128)) + (((int)blockIdx.x) * 16)) + (((int)threadIdx.x) * 4)) + (ff_outer_inner * 2)))))[0] = ((float2*)(Conv2dOutput_local + (((nn_inner * 4) + (ff_outer_inner * 2)))))[0];
}
}
}
}
}
}
```
--
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] jwfromm commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
jwfromm commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-925968890
@tqchen what's the protocol for a case like this where a PR is blocked by requested changes but the requester is away for a long period of time? Is there a way to override that or should @echuraev close and reopen the 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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] jcf94 commented on a change in pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
jcf94 commented on a change in pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#discussion_r683083725
##########
File path: python/tvm/topi/cuda/conv2d_nhwc.py
##########
@@ -85,8 +87,10 @@ def schedule_conv2d_nhwc_direct(cfg, s, Conv):
thread_yz = te.thread_axis((0, vthread_n), "vthread", name="vy")
# Schedule for output
- ni, hi, wi, fi = s[output].op.axis
- bz = s[output].fuse(hi, wi)
+ ni, _, wi, fi = s[output].op.axis
+ bz = wi
+ fi, vec = s[output].split(fi, factor=vec_factor)
+ s[output].vectorize(vec)
Review comment:
Does this actually work? ...... In my experience it's better to put this vectorize at the back of reorder.
And in L96, you didn't add the new vector axis to the reorder list?
##########
File path: python/tvm/topi/cuda/conv2d_nhwc.py
##########
@@ -54,6 +54,7 @@ def schedule_conv2d_nhwc_direct(cfg, s, Conv):
cfg.define_knob("vthread_n", [1] if dynamic_batch else [1, 2])
cfg.define_knob("vthread_c", [1, 2])
cfg.define_knob("step", [16, 3, 32, 64])
+ cfg.define_knob("vectorize", [4, 2, 8, 16])
Review comment:
Better to add a 1 for including no vectorize case.
```suggestion
cfg.define_knob("vectorize", [1, 2, 4, 8, 16])
```
--
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] echuraev commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
echuraev commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-894162828
@jcf94 thank you for your comments. I applied your comments and moved conv2d_nhwc to `topi/gpu` in separate commits to make it easier to review. Please take a look one again.
> Another thing is that you're doing test on a OpenCL device?
Yes, I run topology from #8146 on my Samsung Galaxy A71 and got about 6-7x performance boost.
--
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] tmoreau89 commented on pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#issuecomment-927375110
I've dismissed the review that was blocking progress. In my view we've provided sufficient ping to notify the reviewer who requested changes. @jwfromm I'll leave it to you to merge the PR when convenient.
--
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] jwfromm merged pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
jwfromm merged pull request #8636:
URL: https://github.com/apache/tvm/pull/8636
--
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] jcf94 commented on a change in pull request #8636: [OpenCL] Add vectorization to cuda conv2d_nhwc schedule
Posted by GitBox <gi...@apache.org>.
jcf94 commented on a change in pull request #8636:
URL: https://github.com/apache/tvm/pull/8636#discussion_r683083725
##########
File path: python/tvm/topi/cuda/conv2d_nhwc.py
##########
@@ -85,8 +87,10 @@ def schedule_conv2d_nhwc_direct(cfg, s, Conv):
thread_yz = te.thread_axis((0, vthread_n), "vthread", name="vy")
# Schedule for output
- ni, hi, wi, fi = s[output].op.axis
- bz = s[output].fuse(hi, wi)
+ ni, _, wi, fi = s[output].op.axis
+ bz = wi
+ fi, vec = s[output].split(fi, factor=vec_factor)
+ s[output].vectorize(vec)
Review comment:
Does this actually work? ...... In my experience it's better to put this vectorize at the back of reorder.
And in L96, you didn't add the new vector axis to the reorder list?
##########
File path: python/tvm/topi/cuda/conv2d_nhwc.py
##########
@@ -54,6 +54,7 @@ def schedule_conv2d_nhwc_direct(cfg, s, Conv):
cfg.define_knob("vthread_n", [1] if dynamic_batch else [1, 2])
cfg.define_knob("vthread_c", [1, 2])
cfg.define_knob("step", [16, 3, 32, 64])
+ cfg.define_knob("vectorize", [4, 2, 8, 16])
Review comment:
Better to add a 1 for including no vectorize case.
```suggestion
cfg.define_knob("vectorize", [1, 2, 4, 8, 16])
```
--
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