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/06 19:51:56 UTC
[GitHub] [tvm] ArmageddonKnight opened a new pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
ArmageddonKnight opened a new pull request #8678:
URL: https://github.com/apache/tvm/pull/8678
# Short Summary
Sometimes, when executing CUDA kernels, we might encounter the error `CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES` (e.g., [here](https://discuss.tvm.apache.org/t/cuda-got-error-cuda-error-launch-out-of-resources/4173)). This happens because **the nvcc compiler allocates too many registers per thread**. In the case when we launch the CUDA kernel using too many threads, the GPU will notice that the CUDA kernel requests more registers than what are available on the chip and therefore refuse to launch the kernel.
This hence implies that we need a way of telling nvcc what to expect in terms of the number of threads per block. Luckily, the `__launch_bounds__` directive can help us achieve what we want. In this patch, we add `__launch_bounds__` as part of the CUDA code generation procedure. `__launch_bounds__` will be automatically printed if it is detected that the number of threads per block is a constant integer value. Passing this information to nvcc allows it to spill registers if needed, which might hurt performance, but is still better than having a CUDA kernel that is not functional.
# Q & A
**Q: Would this affect the AutoTVM and the auto-scheduler submodule?**
A: No. Although in those cases the number of threads keeps changing at each trial, the number will be set to a constant when it comes to the code generation phase. Furthermore, in the case when the number of threads per block is not a constant, `__launch_bounds__` will simply not be printed.
Any feedback on this patch is appreciated. @comaniac @icemelon @yzhliu @yidawang
--
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] junrushao1994 merged pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
junrushao1994 merged pull request #8678:
URL: https://github.com/apache/tvm/pull/8678
--
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] junrushao1994 commented on a change in pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on a change in pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#discussion_r684504421
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -46,6 +48,39 @@ void CodeGenCUDA::Init(bool output_ssa) {
void CodeGenCUDA::PrintFuncPrefix() { stream << "extern \"C\" __global__ void"; }
+class threadIdxExtractor : public tir::StmtVisitor {
Review comment:
Wait please fix the naming style
```suggestion
class ThreadIdxExtractor : public tir::StmtVisitor {
```
--
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] ArmageddonKnight commented on pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
ArmageddonKnight commented on pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#issuecomment-894610419
@junrushao1994 Thanks for letting me know. I had a look into this issue. The problem is caused by assigning `threadIdx.x` to `iv->thread_tag` rather than `iv->var->name_hint` (and therefore, the extractor is unable to correctly extract the number of threads per block). To address this issue, I extend the extractor to cover `thread_tag`'s as well. At the same time, if the number of threads per block is extracted as 1, then `__launch_bounds__` will NOT be printed.
I ran the test case again and it works locally.
--
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] vinx13 commented on a change in pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
vinx13 commented on a change in pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#discussion_r684474049
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -46,6 +48,42 @@ void CodeGenCUDA::Init(bool output_ssa) {
void CodeGenCUDA::PrintFuncPrefix() { stream << "extern \"C\" __global__ void"; }
+class threadIdxExtractor : public tir::StmtVisitor {
+ private:
+ void VisitStmt_(const AttrStmtNode* op) final {
+ if (op->attr_key == tir::attr::thread_extent) {
+ IterVar iv = Downcast<IterVar>(op->node);
+ LOG(INFO) << iv->var->name_hint;
Review comment:
remove this
--
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] junrushao1994 commented on pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#issuecomment-894506980
I am very much in favor of this approach. Thanks Bojian!
--
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] junrushao1994 commented on pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#issuecomment-894606943
The syntax should be `__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)`.
However, in the failing test, I am seeing:
```
grid=(16,8,1), block=(32,2,2)
...
__launch_bounds__(1)
```
Would you like to double check? @ArmageddonKnight
--
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] comaniac commented on a change in pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
comaniac commented on a change in pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#discussion_r684513225
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -46,6 +48,39 @@ void CodeGenCUDA::Init(bool output_ssa) {
void CodeGenCUDA::PrintFuncPrefix() { stream << "extern \"C\" __global__ void"; }
+class ThreadIdxExtractor : public tir::StmtVisitor {
+ private:
+ void VisitStmt_(const AttrStmtNode* op) final {
+ if (op->attr_key == tir::attr::thread_extent) {
+ IterVar iv = Downcast<IterVar>(op->node);
+ if (iv->var->name_hint == "threadIdx.x") {
+ threadIdx_x_ext = op->value;
+ }
+ if (iv->var->name_hint == "threadIdx.y") {
+ threadIdx_y_ext = op->value;
+ }
+ if (iv->var->name_hint == "threadIdx.z") {
+ threadIdx_z_ext = op->value;
+ }
+ }
+ StmtVisitor::VisitStmt_(op);
+ }
+
+ public:
+ PrimExpr threadIdx_x_ext = Integer(1), threadIdx_y_ext = Integer(1), threadIdx_z_ext = Integer(1);
+};
+
+void CodeGenCUDA::PrintExtraAttrs(const PrimFunc& f) {
+ ThreadIdxExtractor extractor;
+ extractor(f->body);
+ arith::Analyzer analyzer;
+ PrimExpr threadIdx_ext = analyzer.Simplify(extractor.threadIdx_x_ext * extractor.threadIdx_y_ext *
+ extractor.threadIdx_z_ext);
+ if (const IntImmNode* const threadIdx_ext_int = threadIdx_ext.as<IntImmNode>()) {
+ stream << " __launch_bounds__(" << threadIdx_ext_int->value << ")";
Review comment:
Oh I see. Then we don't need a new line 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] junrushao1994 edited a comment on pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
junrushao1994 edited a comment on pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#issuecomment-894606943
The syntax should be `__launch_bounds__(maxThreadsPerBlock) ` or `__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)`.
However, in the failing test, I am seeing:
```
grid=(16,8,1), block=(32,2,2)
...
__launch_bounds__(1)
```
Would you like to double check? @ArmageddonKnight
--
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] ArmageddonKnight commented on a change in pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
ArmageddonKnight commented on a change in pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#discussion_r684504964
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -46,6 +48,39 @@ void CodeGenCUDA::Init(bool output_ssa) {
void CodeGenCUDA::PrintFuncPrefix() { stream << "extern \"C\" __global__ void"; }
+class threadIdxExtractor : public tir::StmtVisitor {
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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ArmageddonKnight commented on a change in pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
ArmageddonKnight commented on a change in pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#discussion_r684474908
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -46,6 +48,42 @@ void CodeGenCUDA::Init(bool output_ssa) {
void CodeGenCUDA::PrintFuncPrefix() { stream << "extern \"C\" __global__ void"; }
+class threadIdxExtractor : public tir::StmtVisitor {
+ private:
+ void VisitStmt_(const AttrStmtNode* op) final {
+ if (op->attr_key == tir::attr::thread_extent) {
+ IterVar iv = Downcast<IterVar>(op->node);
+ LOG(INFO) << iv->var->name_hint;
Review comment:
Removed.
--
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] junrushao1994 commented on pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#issuecomment-894619557
Thanks Bojian! It makes perfect sense to me, which reminds me of a similar bug we encountered in TensorIR lowering :-)
--
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] ArmageddonKnight edited a comment on pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
ArmageddonKnight edited a comment on pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#issuecomment-894610419
@junrushao1994 Thanks for letting me know. I had a look into this issue. The problem is caused by assigning `threadIdx.x` to `iv->thread_tag` rather than `iv->var->name_hint` (and therefore, the extractor is unable to correctly extract the number of threads per block). To address this issue, I extend the extractor to cover `thread_tag`'s as well. At the same time, if the number of threads per block is extracted as 1, then `__launch_bounds__` will NOT be printed.
I ran the test case again and it works locally.
--
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] ArmageddonKnight commented on a change in pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
ArmageddonKnight commented on a change in pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#discussion_r684512382
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -46,6 +48,39 @@ void CodeGenCUDA::Init(bool output_ssa) {
void CodeGenCUDA::PrintFuncPrefix() { stream << "extern \"C\" __global__ void"; }
+class ThreadIdxExtractor : public tir::StmtVisitor {
+ private:
+ void VisitStmt_(const AttrStmtNode* op) final {
+ if (op->attr_key == tir::attr::thread_extent) {
+ IterVar iv = Downcast<IterVar>(op->node);
+ if (iv->var->name_hint == "threadIdx.x") {
+ threadIdx_x_ext = op->value;
+ }
+ if (iv->var->name_hint == "threadIdx.y") {
+ threadIdx_y_ext = op->value;
+ }
+ if (iv->var->name_hint == "threadIdx.z") {
+ threadIdx_z_ext = op->value;
+ }
+ }
+ StmtVisitor::VisitStmt_(op);
+ }
+
+ public:
+ PrimExpr threadIdx_x_ext = Integer(1), threadIdx_y_ext = Integer(1), threadIdx_z_ext = Integer(1);
+};
+
+void CodeGenCUDA::PrintExtraAttrs(const PrimFunc& f) {
+ ThreadIdxExtractor extractor;
+ extractor(f->body);
+ arith::Analyzer analyzer;
+ PrimExpr threadIdx_ext = analyzer.Simplify(extractor.threadIdx_x_ext * extractor.threadIdx_y_ext *
+ extractor.threadIdx_z_ext);
+ if (const IntImmNode* const threadIdx_ext_int = threadIdx_ext.as<IntImmNode>()) {
+ stream << " __launch_bounds__(" << threadIdx_ext_int->value << ")";
Review comment:
I would say currently it looks ok
--
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] ArmageddonKnight commented on a change in pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
ArmageddonKnight commented on a change in pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#discussion_r684513647
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -46,6 +48,39 @@ void CodeGenCUDA::Init(bool output_ssa) {
void CodeGenCUDA::PrintFuncPrefix() { stream << "extern \"C\" __global__ void"; }
+class ThreadIdxExtractor : public tir::StmtVisitor {
+ private:
+ void VisitStmt_(const AttrStmtNode* op) final {
+ if (op->attr_key == tir::attr::thread_extent) {
+ IterVar iv = Downcast<IterVar>(op->node);
+ if (iv->var->name_hint == "threadIdx.x") {
+ threadIdx_x_ext = op->value;
+ }
+ if (iv->var->name_hint == "threadIdx.y") {
+ threadIdx_y_ext = op->value;
+ }
+ if (iv->var->name_hint == "threadIdx.z") {
+ threadIdx_z_ext = op->value;
+ }
+ }
+ StmtVisitor::VisitStmt_(op);
+ }
+
+ public:
+ PrimExpr threadIdx_x_ext = Integer(1), threadIdx_y_ext = Integer(1), threadIdx_z_ext = Integer(1);
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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao1994 commented on pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#issuecomment-894650408
Thanks Bojian! It’s really super useful improvement in cuda codegen
--
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] ArmageddonKnight commented on pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
ArmageddonKnight commented on pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#issuecomment-894650040
@vinx13 @comaniac @junrushao1994 FYI, the patch has passed all the checks. Could you please merge it if possible?
--
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] comaniac commented on a change in pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
comaniac commented on a change in pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#discussion_r684511293
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -46,6 +48,39 @@ void CodeGenCUDA::Init(bool output_ssa) {
void CodeGenCUDA::PrintFuncPrefix() { stream << "extern \"C\" __global__ void"; }
+class ThreadIdxExtractor : public tir::StmtVisitor {
+ private:
+ void VisitStmt_(const AttrStmtNode* op) final {
+ if (op->attr_key == tir::attr::thread_extent) {
+ IterVar iv = Downcast<IterVar>(op->node);
+ if (iv->var->name_hint == "threadIdx.x") {
+ threadIdx_x_ext = op->value;
+ }
+ if (iv->var->name_hint == "threadIdx.y") {
+ threadIdx_y_ext = op->value;
+ }
+ if (iv->var->name_hint == "threadIdx.z") {
+ threadIdx_z_ext = op->value;
+ }
+ }
+ StmtVisitor::VisitStmt_(op);
+ }
+
+ public:
+ PrimExpr threadIdx_x_ext = Integer(1), threadIdx_y_ext = Integer(1), threadIdx_z_ext = Integer(1);
Review comment:
nit: It would be clearer to put them in 3 lines.
##########
File path: src/target/source/codegen_cuda.cc
##########
@@ -46,6 +48,39 @@ void CodeGenCUDA::Init(bool output_ssa) {
void CodeGenCUDA::PrintFuncPrefix() { stream << "extern \"C\" __global__ void"; }
+class ThreadIdxExtractor : public tir::StmtVisitor {
+ private:
+ void VisitStmt_(const AttrStmtNode* op) final {
+ if (op->attr_key == tir::attr::thread_extent) {
+ IterVar iv = Downcast<IterVar>(op->node);
+ if (iv->var->name_hint == "threadIdx.x") {
+ threadIdx_x_ext = op->value;
+ }
+ if (iv->var->name_hint == "threadIdx.y") {
+ threadIdx_y_ext = op->value;
+ }
+ if (iv->var->name_hint == "threadIdx.z") {
+ threadIdx_z_ext = op->value;
+ }
+ }
+ StmtVisitor::VisitStmt_(op);
+ }
+
+ public:
+ PrimExpr threadIdx_x_ext = Integer(1), threadIdx_y_ext = Integer(1), threadIdx_z_ext = Integer(1);
+};
+
+void CodeGenCUDA::PrintExtraAttrs(const PrimFunc& f) {
+ ThreadIdxExtractor extractor;
+ extractor(f->body);
+ arith::Analyzer analyzer;
+ PrimExpr threadIdx_ext = analyzer.Simplify(extractor.threadIdx_x_ext * extractor.threadIdx_y_ext *
+ extractor.threadIdx_z_ext);
+ if (const IntImmNode* const threadIdx_ext_int = threadIdx_ext.as<IntImmNode>()) {
+ stream << " __launch_bounds__(" << threadIdx_ext_int->value << ")";
Review comment:
Do we need a `\n`?
--
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] ArmageddonKnight commented on pull request #8678: [Target] Add __launch_bounds__ directive as part of the CUDA code generation
Posted by GitBox <gi...@apache.org>.
ArmageddonKnight commented on pull request #8678:
URL: https://github.com/apache/tvm/pull/8678#issuecomment-894664430
Thanks @junrushao1994
--
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