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/09/16 20:53:40 UTC
[GitHub] [tvm] MasterJH5574 opened a new pull request, #12825: [MetaSchedule] Relax conditions of rule Cross-Thread Reduction
MasterJH5574 opened a new pull request, #12825:
URL: https://github.com/apache/tvm/pull/12825
This PR relaxes the conditions of Meta-Schedule schedule rule CrossThreadReduction. The rules are previously a bit over-strict, and some workloads with small reduction loop length are unable to be optimized by cross-thread reduction automatically. In this PR, we relax the rules so that such workloads can be optimized.
cc @junrushao
--
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] junrushao commented on a diff in pull request #12825: [MetaSchedule] Relax conditions of rule Cross-Thread Reduction
Posted by GitBox <gi...@apache.org>.
junrushao commented on code in PR #12825:
URL: https://github.com/apache/tvm/pull/12825#discussion_r973385920
##########
tests/python/unittest/test_meta_schedule_schedule_rule_cross_thread_reduction.py:
##########
@@ -663,8 +685,85 @@ def argmax_1(
)
+def test_gpu_argmax_32():
+ @T.prim_func
+ def argmax_0(
+ idx: T.Buffer[(1, 32), "int32"],
+ val: T.Buffer[(1, 32), "float32"],
+ argmax_v0: T.Buffer[(1,), "int32"],
+ argmax_v1: T.Buffer[(1,), "float32"],
+ ) -> None:
+ # body
+ # with T.block("root")
+ for i0, i1 in T.grid(1, 32):
+ with T.block("argmax"):
+ i, k = T.axis.remap("SR", [i0, i1])
+ T.reads(idx[i, k], val[i, k])
+ T.writes(argmax_v0[i], argmax_v1[i])
+ with T.init():
+ argmax_v0[i] = -1
+ argmax_v1[i] = T.float32(-3.4028234663852886e38)
+ v_argmax_v0: T.int32 = T.Select(argmax_v1[i] >= val[i, k], argmax_v0[i], idx[i, k])
+ v_argmax_v1: T.float32 = T.Select(
+ argmax_v1[i] >= val[i, k], argmax_v1[i], val[i, k]
+ )
+ argmax_v0[i] = v_argmax_v0
+ argmax_v1[i] = v_argmax_v1
+
+ @T.prim_func
+ def argmax_1(
+ idx: T.Buffer[(1, 32), "int32"],
+ val: T.Buffer[(1, 32), "float32"],
+ argmax_v0: T.Buffer[(1,), "int32"],
+ argmax_v1: T.Buffer[(1,), "float32"],
+ ) -> None:
+ # body
+ # with T.block("root")
+ for i0, i1_0 in T.grid(1, 1):
+ for i1_1 in T.thread_binding(64, thread="threadIdx.x"):
+ with T.block("argmax"):
+ i = T.axis.spatial(1, i0)
+ k = T.axis.reduce(32, i1_0 * 64 + i1_1)
+ T.where(i1_0 * 64 + i1_1 < 32)
+ T.reads(idx[i, k], val[i, k])
+ T.writes(argmax_v0[i], argmax_v1[i])
+ with T.init():
+ argmax_v0[i] = -1
+ argmax_v1[i] = T.float32(-3.4028234663852886e38)
+ v_argmax_v0: T.int32 = T.Select(
+ argmax_v1[i] >= val[i, k], argmax_v0[i], idx[i, k]
+ )
+ v_argmax_v1: T.float32 = T.Select(
+ argmax_v1[i] >= val[i, k], argmax_v1[i], val[i, k]
+ )
+ argmax_v0[i] = v_argmax_v0
+ argmax_v1[i] = v_argmax_v1
+
+ decision_0 = [] # type: ignore
+ decision_1 = [
+ ("SampleCategorical", 4),
+ ]
+
+ mod = argmax_32
+ actual = ms.TuneContext(
+ mod=mod,
+ target=Target("nvidia/geforce-rtx-3090", host="llvm"),
+ space_generator=ms.space_generator.PostOrderApply(),
+ sch_rules=get_rules("cuda", ms.schedule_rule.CrossThreadReduction),
+ task_name="test",
+ ).generate_design_space()
+ print(actual[1].mod.script())
Review Comment:
let's remove this debug printing
--
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] MasterJH5574 commented on a diff in pull request #12825: [MetaSchedule] Relax conditions of rule Cross-Thread Reduction
Posted by GitBox <gi...@apache.org>.
MasterJH5574 commented on code in PR #12825:
URL: https://github.com/apache/tvm/pull/12825#discussion_r973393496
##########
tests/python/unittest/test_meta_schedule_schedule_rule_cross_thread_reduction.py:
##########
@@ -663,8 +685,85 @@ def argmax_1(
)
+def test_gpu_argmax_32():
+ @T.prim_func
+ def argmax_0(
+ idx: T.Buffer[(1, 32), "int32"],
+ val: T.Buffer[(1, 32), "float32"],
+ argmax_v0: T.Buffer[(1,), "int32"],
+ argmax_v1: T.Buffer[(1,), "float32"],
+ ) -> None:
+ # body
+ # with T.block("root")
+ for i0, i1 in T.grid(1, 32):
+ with T.block("argmax"):
+ i, k = T.axis.remap("SR", [i0, i1])
+ T.reads(idx[i, k], val[i, k])
+ T.writes(argmax_v0[i], argmax_v1[i])
+ with T.init():
+ argmax_v0[i] = -1
+ argmax_v1[i] = T.float32(-3.4028234663852886e38)
+ v_argmax_v0: T.int32 = T.Select(argmax_v1[i] >= val[i, k], argmax_v0[i], idx[i, k])
+ v_argmax_v1: T.float32 = T.Select(
+ argmax_v1[i] >= val[i, k], argmax_v1[i], val[i, k]
+ )
+ argmax_v0[i] = v_argmax_v0
+ argmax_v1[i] = v_argmax_v1
+
+ @T.prim_func
+ def argmax_1(
+ idx: T.Buffer[(1, 32), "int32"],
+ val: T.Buffer[(1, 32), "float32"],
+ argmax_v0: T.Buffer[(1,), "int32"],
+ argmax_v1: T.Buffer[(1,), "float32"],
+ ) -> None:
+ # body
+ # with T.block("root")
+ for i0, i1_0 in T.grid(1, 1):
+ for i1_1 in T.thread_binding(64, thread="threadIdx.x"):
+ with T.block("argmax"):
+ i = T.axis.spatial(1, i0)
+ k = T.axis.reduce(32, i1_0 * 64 + i1_1)
+ T.where(i1_0 * 64 + i1_1 < 32)
+ T.reads(idx[i, k], val[i, k])
+ T.writes(argmax_v0[i], argmax_v1[i])
+ with T.init():
+ argmax_v0[i] = -1
+ argmax_v1[i] = T.float32(-3.4028234663852886e38)
+ v_argmax_v0: T.int32 = T.Select(
+ argmax_v1[i] >= val[i, k], argmax_v0[i], idx[i, k]
+ )
+ v_argmax_v1: T.float32 = T.Select(
+ argmax_v1[i] >= val[i, k], argmax_v1[i], val[i, k]
+ )
+ argmax_v0[i] = v_argmax_v0
+ argmax_v1[i] = v_argmax_v1
+
+ decision_0 = [] # type: ignore
+ decision_1 = [
+ ("SampleCategorical", 4),
+ ]
+
+ mod = argmax_32
+ actual = ms.TuneContext(
+ mod=mod,
+ target=Target("nvidia/geforce-rtx-3090", host="llvm"),
+ space_generator=ms.space_generator.PostOrderApply(),
+ sch_rules=get_rules("cuda", ms.schedule_rule.CrossThreadReduction),
+ task_name="test",
+ ).generate_design_space()
+ print(actual[1].mod.script())
Review Comment:
Done!
--
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] MasterJH5574 commented on a diff in pull request #12825: [MetaSchedule] Relax conditions of rule Cross-Thread Reduction
Posted by GitBox <gi...@apache.org>.
MasterJH5574 commented on code in PR #12825:
URL: https://github.com/apache/tvm/pull/12825#discussion_r973386916
##########
tests/python/unittest/test_meta_schedule_schedule_rule_cross_thread_reduction.py:
##########
@@ -663,8 +685,85 @@ def argmax_1(
)
+def test_gpu_argmax_32():
+ @T.prim_func
+ def argmax_0(
+ idx: T.Buffer[(1, 32), "int32"],
+ val: T.Buffer[(1, 32), "float32"],
+ argmax_v0: T.Buffer[(1,), "int32"],
+ argmax_v1: T.Buffer[(1,), "float32"],
+ ) -> None:
+ # body
+ # with T.block("root")
+ for i0, i1 in T.grid(1, 32):
+ with T.block("argmax"):
+ i, k = T.axis.remap("SR", [i0, i1])
+ T.reads(idx[i, k], val[i, k])
+ T.writes(argmax_v0[i], argmax_v1[i])
+ with T.init():
+ argmax_v0[i] = -1
+ argmax_v1[i] = T.float32(-3.4028234663852886e38)
+ v_argmax_v0: T.int32 = T.Select(argmax_v1[i] >= val[i, k], argmax_v0[i], idx[i, k])
+ v_argmax_v1: T.float32 = T.Select(
+ argmax_v1[i] >= val[i, k], argmax_v1[i], val[i, k]
+ )
+ argmax_v0[i] = v_argmax_v0
+ argmax_v1[i] = v_argmax_v1
+
+ @T.prim_func
+ def argmax_1(
+ idx: T.Buffer[(1, 32), "int32"],
+ val: T.Buffer[(1, 32), "float32"],
+ argmax_v0: T.Buffer[(1,), "int32"],
+ argmax_v1: T.Buffer[(1,), "float32"],
+ ) -> None:
+ # body
+ # with T.block("root")
+ for i0, i1_0 in T.grid(1, 1):
+ for i1_1 in T.thread_binding(64, thread="threadIdx.x"):
+ with T.block("argmax"):
+ i = T.axis.spatial(1, i0)
+ k = T.axis.reduce(32, i1_0 * 64 + i1_1)
+ T.where(i1_0 * 64 + i1_1 < 32)
+ T.reads(idx[i, k], val[i, k])
+ T.writes(argmax_v0[i], argmax_v1[i])
+ with T.init():
+ argmax_v0[i] = -1
+ argmax_v1[i] = T.float32(-3.4028234663852886e38)
+ v_argmax_v0: T.int32 = T.Select(
+ argmax_v1[i] >= val[i, k], argmax_v0[i], idx[i, k]
+ )
+ v_argmax_v1: T.float32 = T.Select(
+ argmax_v1[i] >= val[i, k], argmax_v1[i], val[i, k]
+ )
+ argmax_v0[i] = v_argmax_v0
+ argmax_v1[i] = v_argmax_v1
+
+ decision_0 = [] # type: ignore
+ decision_1 = [
+ ("SampleCategorical", 4),
+ ]
+
+ mod = argmax_32
+ actual = ms.TuneContext(
+ mod=mod,
+ target=Target("nvidia/geforce-rtx-3090", host="llvm"),
+ space_generator=ms.space_generator.PostOrderApply(),
+ sch_rules=get_rules("cuda", ms.schedule_rule.CrossThreadReduction),
+ task_name="test",
+ ).generate_design_space()
+ print(actual[1].mod.script())
Review Comment:
Sorry my negligence 🤦♂️
--
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] MasterJH5574 commented on pull request #12825: [MetaSchedule] Relax conditions of rule Cross-Thread Reduction
Posted by GitBox <gi...@apache.org>.
MasterJH5574 commented on PR #12825:
URL: https://github.com/apache/tvm/pull/12825#issuecomment-1250170948
> The only question from me is: is it going to make the search space smaller if rfactor or cross-thread reduction is improperly applied? If not, we can get it in
@junrushao This PR is going to make the search space larger for workloads which have small reduction loop length. Previously, the transformations (rfactor or cross-thread reduction) are not applied to these workloads at all. The transformations are not improperly applied.
--
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] junrushao commented on pull request #12825: [MetaSchedule] Relax conditions of rule Cross-Thread Reduction
Posted by GitBox <gi...@apache.org>.
junrushao commented on PR #12825:
URL: https://github.com/apache/tvm/pull/12825#issuecomment-1250139158
The only question from me is: is it going to make the search space smaller if rfactor or cross-thread reduction is improperly applied? If not, we can get it in
--
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] junrushao merged pull request #12825: [MetaSchedule] Relax conditions of rule Cross-Thread Reduction
Posted by GitBox <gi...@apache.org>.
junrushao merged PR #12825:
URL: https://github.com/apache/tvm/pull/12825
--
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