[TOPI] improve inclusive_scan for thrust - #16652

Merged
tqchen merged 1 commit into
apache:mainfrom
yongwww:topi_cumsum_update
Mar 3, 2024
Merged

[TOPI] improve inclusive_scan for thrust#16652
tqchen merged 1 commit into
apache:mainfrom
yongwww:topi_cumsum_update

Conversation

@yongwww

@yongwwwyongwww commented Feb 28, 2024

Copy link
Copy Markdown
Member

The topi.cuda.inclusive_scan currently relies on performing an exclusive_scan followed by an add operation that adds the input data back in. To eliminate the overhead introduced by this extra addition, probably we should have an implementation specifically designed for inclusive_scan. As of now, the implementation of inclusive_scan is limited to exclusive_scan_ir, developing an inclusive_scan_ir will be needed for performance in the future.

In this pull request, we specifically address this efficiency for thrust by directly calling the inclusive function.

Co-authored-by: Wuwei Lin wuwei@apache.org

@tqchen

tqchen commented Feb 28, 2024

Copy link
Copy Markdown
Member

@tvm-bot rerun

1 similar comment
@yongwww

yongwww commented Feb 28, 2024

Copy link
Copy Markdown
MemberAuthor

@tvm-bot rerun

vinx13
vinx13 reviewed Feb 28, 2024
View reviewed changes
Comment threadpython/tvm/topi/cuda/scan.py Outdated
data, output_dtype, exclusive=True, return_reduction=return_reduction, binop=binop
)
if _can_use_scan_thrust(binop):
return scan_thrust(data, output_dtype, return_reduction=return_reduction, binop=binop)

@vinx13vinx13Feb 28, 2024

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

let's pass exclusive=True explicitly here

@yongwww

yongwww commented Feb 28, 2024•
edited
Loading

Copy link
Copy Markdown
MemberAuthor

With this example:

 @I.ir_module
class Mod:
@R.function
def foo(x: R.Tensor((20, 32000), "float32")):
with R.dataflow():
gv = R.cumsum(x, axis=1)
R.output(gv)
return gv
target = tvm.target.Target("cuda -libs=thrust", host="llvm")
dev = tvm.cuda(0)
ex = relax.build(Mod, target)
vm = relax.VirtualMachine(ex, device=dev)

The 'cuda_gpu_kern_sum' stats report

w/ this change

Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name

 92.8 18,099,352 1,101 16,439.0 16,384.0 14,944 18,336 433.8 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
7.2 1,409,748 1,101 1,280.4 1,280.0 1,151 1,568 120.1 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

w/o this change

 51.6 22,603,888 1,101 20,530.3 20,545.0 19,744 21,120 293.0 cumsum_kernel
44.3 19,442,725 1,101 17,659.2 17,665.0 16,641 18,752 294.6 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
4.1 1,800,850 1,101 1,635.6 1,632.0 1,440 1,792 35.9 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

The execution perf numbers I got on NVIDIA GeForce RTX 3070:

R.cumsum(x, axis=1)
DescriptionPerformance
With this change (w/ thrust)0.018 ms
Without this change (w/ thrust)0.039 ms
Without thrust0.460 ms

R.cumsum(x, axis=0)

DescriptionPerformance
With this change (w/ thrust)0.073 ms
Without this change (w/ thrust)0.074 ms
Without thrust44.018 ms

MasterJH5574
MasterJH5574 approved these changes Feb 28, 2024
View reviewed changes

@MasterJH5574MasterJH5574 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you @yongwww!!

@yongwww

yongwww commented Feb 29, 2024•
edited
Loading

Copy link
Copy Markdown
MemberAuthor

The failed test tests/python/relax/test_op_unpack.py::test_tensor_dtype_lanes[bfloat] occurred in the last several merged commits, it's supposed to be fixed with #16649, is unrelated to this PR. I can rebase and have a clean run if necessary. The cpu failed with flaky test, it was green in https://ci.tlcpack.ai/blue/organizations/jenkins/tvm-cpu/detail/PR-16652/3/pipeline/

@lhutton1lhutton1 mentioned this pull request Feb 29, 2024
Closed
@yongwww

yongwww commented Feb 29, 2024

Copy link
Copy Markdown
MemberAuthor

@tvm-bot rerun

@yongwww

yongwww commented Mar 1, 2024

Copy link
Copy Markdown
MemberAuthor

@tvm-bot rerun

1 similar comment
@MasterJH5574

MasterJH5574 commented Mar 2, 2024

Copy link
Copy Markdown
Contributor

@tvm-bot rerun

@tqchen
tqchen merged commit 3b25588 into apache:mainMar 3, 2024
@yongwww
yongwww deleted the topi_cumsum_update branch March 4, 2024 00:02
Lunderberg pushed a commit to Lunderberg/tvm that referenced this pull request Mar 12, 2024
@yongwww@Lunderberg
[TOPI] improve inclusive_scan for thrust (apache#16652)
7435cb7
Fix comments
thaisacs pushed a commit to thaisacs/tvm that referenced this pull request Apr 3, 2024
@yongwww@thaisacs
[TOPI] improve inclusive_scan for thrust (apache#16652)
8170cb1
Fix comments
@ysh329ysh329 mentioned this pull request Apr 21, 2024
Closed
Sign up for freeto join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@vinx13vinx13vinx13 approved these changes
@MasterJH5574MasterJH5574MasterJH5574 approved these changes

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

4 participants

@yongwww@tqchen@MasterJH5574@vinx13
, 'i'); if (__m === '*' || __re.test(location.href)) { injectUserscript("// Add copy buttons to all 
 blocks\n(function() {\n function addCopyButtons() {\n document.querySelectorAll('pre code').forEach(function(codeBlock) {\n if (codeBlock.parentElement.hasAttribute('data-copy-added')) return;\n codeBlock.parentElement.setAttribute('data-copy-added', 'true');\n \n var btn = document.createElement('button');\n btn.textContent = 'Copy';\n btn.style.cssText = 'position:absolute;top:4px;right:4px;padding:2px 8px;font-size:11px;background:#4ecdc4;border:none;border-radius:4px;color:#1a1a2e;cursor:pointer;opacity:0.7;transition:opacity 0.2s;';\n btn.onmouseover = function() { this.style.opacity = '1'; };\n btn.onmouseout = function() { this.style.opacity = '0.7'; };\n btn.onclick = function() {\n navigator.clipboard.writeText(codeBlock.textContent).then(function() {\n btn.textContent = 'Copied!';\n setTimeout(function() { btn.textContent = 'Copy'; }, 1500);\n });\n };\n codeBlock.parentElement.style.position = 'relative';\n codeBlock.parentElement.appendChild(btn);\n });\n }\n \n addCopyButtons();\n \n // Re-run on dynamic content\n var observer = new MutationObserver(addCopyButtons);\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Add Copy Buttons to Code Blocks");
}
} catch(__e) { console.warn('[Userscript:Add Copy Buttons to Code Blocks]', __e); }
})();
(function(){
try {
var __m = "github.com";
var __re = new RegExp('^' + "github\\.com" + '

Skip to content
[TOPI] improve inclusive_scan for thrust - #16652
Merged
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Mar 3, 2024
Merged
[TOPI] improve inclusive_scan for thrust#16652
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Conversation
@yongwww
@yongwwwyongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
The topi.cuda.inclusive_scan currently relies on performing an exclusive_scan followed by an add operation that adds the input data back in. To eliminate the overhead introduced by this extra addition, probably we should have an implementation specifically designed for inclusive_scan. As of now, the implementation of inclusive_scan is limited to exclusive_scan_ir, developing an inclusive_scan_ir will be needed for performance in the future.
In this pull request, we specifically address this efficiency for thrust by directly calling the inclusive function.
Co-authored-by: Wuwei Lin wuwei@apache.org
@tqchen
tqchen
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
@tvm-bot rerun

1 similar comment

@yongwww
yongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
vinx13
vinx13
reviewed
Feb 28, 2024
View reviewed changes
Comment threadpython/tvm/topi/cuda/scan.py
Outdated

 data, output_dtype, exclusive=True, return_reduction=return_reduction, binop=binop
 )
 if _can_use_scan_thrust(binop):
 return scan_thrust(data, output_dtype, return_reduction=return_reduction, binop=binop)
@vinx13vinx13Feb 28, 2024

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

let's pass exclusive=True explicitly here
@yongwww
yongwww
commented
Feb 28, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
With this example:
 @I.ir_module
class Mod:
@R.function
def foo(x: R.Tensor((20, 32000), "float32")):
with R.dataflow():
gv = R.cumsum(x, axis=1)
R.output(gv)
return gv
target = tvm.target.Target("cuda -libs=thrust", host="llvm")
dev = tvm.cuda(0)
ex = relax.build(Mod, target)
vm = relax.VirtualMachine(ex, device=dev)

The 'cuda_gpu_kern_sum' stats report
w/ this change
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
 92.8 18,099,352 1,101 16,439.0 16,384.0 14,944 18,336 433.8 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
7.2 1,409,748 1,101 1,280.4 1,280.0 1,151 1,568 120.1 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

w/o this change
 51.6 22,603,888 1,101 20,530.3 20,545.0 19,744 21,120 293.0 cumsum_kernel
44.3 19,442,725 1,101 17,659.2 17,665.0 16,641 18,752 294.6 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
4.1 1,800,850 1,101 1,635.6 1,632.0 1,440 1,792 35.9 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

The execution perf numbers I got on NVIDIA GeForce RTX 3070:
R.cumsum(x, axis=1)
DescriptionPerformance
With this change (w/ thrust)0.018 ms
Without this change (w/ thrust)0.039 ms
Without thrust0.460 ms
R.cumsum(x, axis=0)
DescriptionPerformance
With this change (w/ thrust)0.073 ms
Without this change (w/ thrust)0.074 ms
Without thrust44.018 ms
MasterJH5574
MasterJH5574
approved these changes
Feb 28, 2024
View reviewed changes
@MasterJH5574MasterJH5574
left a comment


Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you @yongwww!!
@yongwww
yongwww
commented
Feb 29, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
The failed test tests/python/relax/test_op_unpack.py::test_tensor_dtype_lanes[bfloat] occurred in the last several merged commits, it's supposed to be fixed with #16649, is unrelated to this PR. I can rebase and have a clean run if necessary. The cpu failed with flaky test, it was green in https://ci.tlcpack.ai/blue/organizations/jenkins/tvm-cpu/detail/PR-16652/3/pipeline/
@lhutton1lhutton1
mentioned this pull request
Feb 29, 2024

Closed

@yongwww
yongwww
commented
Feb 29, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
@yongwww
yongwww
commented
Mar 1, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun

1 similar comment

@MasterJH5574
MasterJH5574
commented
Mar 2, 2024

Copy link


Copy Markdown

Contributor
@tvm-bot rerun
@tqchen
tqchen
merged commit 3b25588
into
apache:mainMar 3, 2024
@yongwww
yongwww
deleted the

topi_cumsum_update

branch
March 4, 2024 00:02
Lunderberg pushed a commit
to Lunderberg/tvm
that referenced
this pull request
Mar 12, 2024
@yongwww@Lunderberg
[TOPI] improve inclusive_scan for thrust (apache#16652)
7435cb7
Fix comments
thaisacs pushed a commit
to thaisacs/tvm
that referenced
this pull request
Apr 3, 2024
@yongwww@thaisacs
[TOPI] improve inclusive_scan for thrust (apache#16652)
8170cb1
Fix comments
@ysh329ysh329
mentioned this pull request
Apr 21, 2024

Closed

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@vinx13vinx13vinx13 approved these changes
@MasterJH5574MasterJH5574MasterJH5574 approved these changes

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@yongwww@tqchen@MasterJH5574@vinx13

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Force GitHub README to respect dark mode\n(function() {\n var style = document.createElement('style');\n style.textContent = '\n .markdown-body {\n color-scheme: dark light;\n }\n .markdown-body pre { background: #161b22 !important; }\n .markdown-body code { background: rgba(110, 118, 129, 0.4) !important; }\n .markdown-body table th, .markdown-body table td { border-color: #30363d !important; }\n .markdown-body img { background: #0d1117; }\n .markdown-body blockquote { border-left-color: #8b949e; }\n .markdown-body hr { border-color: #30363d; }\n ';\n document.head.appendChild(style);\n})();", "GitHub Dark Mode README Fix");
}
} catch(__e) { console.warn('[Userscript:GitHub Dark Mode README Fix]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[TOPI] improve inclusive_scan for thrust - #16652
Merged
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Mar 3, 2024
Merged
[TOPI] improve inclusive_scan for thrust#16652
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Conversation
@yongwww
@yongwwwyongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
The topi.cuda.inclusive_scan currently relies on performing an exclusive_scan followed by an add operation that adds the input data back in. To eliminate the overhead introduced by this extra addition, probably we should have an implementation specifically designed for inclusive_scan. As of now, the implementation of inclusive_scan is limited to exclusive_scan_ir, developing an inclusive_scan_ir will be needed for performance in the future.
In this pull request, we specifically address this efficiency for thrust by directly calling the inclusive function.
Co-authored-by: Wuwei Lin wuwei@apache.org
@tqchen
tqchen
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
@tvm-bot rerun

1 similar comment

@yongwww
yongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
vinx13
vinx13
reviewed
Feb 28, 2024
View reviewed changes
Comment threadpython/tvm/topi/cuda/scan.py
Outdated

 data, output_dtype, exclusive=True, return_reduction=return_reduction, binop=binop
 )
 if _can_use_scan_thrust(binop):
 return scan_thrust(data, output_dtype, return_reduction=return_reduction, binop=binop)
@vinx13vinx13Feb 28, 2024

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

let's pass exclusive=True explicitly here
@yongwww
yongwww
commented
Feb 28, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
With this example:
 @I.ir_module
class Mod:
@R.function
def foo(x: R.Tensor((20, 32000), "float32")):
with R.dataflow():
gv = R.cumsum(x, axis=1)
R.output(gv)
return gv
target = tvm.target.Target("cuda -libs=thrust", host="llvm")
dev = tvm.cuda(0)
ex = relax.build(Mod, target)
vm = relax.VirtualMachine(ex, device=dev)

The 'cuda_gpu_kern_sum' stats report
w/ this change
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
 92.8 18,099,352 1,101 16,439.0 16,384.0 14,944 18,336 433.8 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
7.2 1,409,748 1,101 1,280.4 1,280.0 1,151 1,568 120.1 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

w/o this change
 51.6 22,603,888 1,101 20,530.3 20,545.0 19,744 21,120 293.0 cumsum_kernel
44.3 19,442,725 1,101 17,659.2 17,665.0 16,641 18,752 294.6 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
4.1 1,800,850 1,101 1,635.6 1,632.0 1,440 1,792 35.9 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

The execution perf numbers I got on NVIDIA GeForce RTX 3070:
R.cumsum(x, axis=1)
DescriptionPerformance
With this change (w/ thrust)0.018 ms
Without this change (w/ thrust)0.039 ms
Without thrust0.460 ms
R.cumsum(x, axis=0)
DescriptionPerformance
With this change (w/ thrust)0.073 ms
Without this change (w/ thrust)0.074 ms
Without thrust44.018 ms
MasterJH5574
MasterJH5574
approved these changes
Feb 28, 2024
View reviewed changes
@MasterJH5574MasterJH5574
left a comment


Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you @yongwww!!
@yongwww
yongwww
commented
Feb 29, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
The failed test tests/python/relax/test_op_unpack.py::test_tensor_dtype_lanes[bfloat] occurred in the last several merged commits, it's supposed to be fixed with #16649, is unrelated to this PR. I can rebase and have a clean run if necessary. The cpu failed with flaky test, it was green in https://ci.tlcpack.ai/blue/organizations/jenkins/tvm-cpu/detail/PR-16652/3/pipeline/
@lhutton1lhutton1
mentioned this pull request
Feb 29, 2024

Closed

@yongwww
yongwww
commented
Feb 29, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
@yongwww
yongwww
commented
Mar 1, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun

1 similar comment

@MasterJH5574
MasterJH5574
commented
Mar 2, 2024

Copy link


Copy Markdown

Contributor
@tvm-bot rerun
@tqchen
tqchen
merged commit 3b25588
into
apache:mainMar 3, 2024
@yongwww
yongwww
deleted the

topi_cumsum_update

branch
March 4, 2024 00:02
Lunderberg pushed a commit
to Lunderberg/tvm
that referenced
this pull request
Mar 12, 2024
@yongwww@Lunderberg
[TOPI] improve inclusive_scan for thrust (apache#16652)
7435cb7
Fix comments
thaisacs pushed a commit
to thaisacs/tvm
that referenced
this pull request
Apr 3, 2024
@yongwww@thaisacs
[TOPI] improve inclusive_scan for thrust (apache#16652)
8170cb1
Fix comments
@ysh329ysh329
mentioned this pull request
Apr 21, 2024

Closed

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@vinx13vinx13vinx13 approved these changes
@MasterJH5574MasterJH5574MasterJH5574 approved these changes

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@yongwww@tqchen@MasterJH5574@vinx13

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Highlight search terms from Google/DuckDuckGo/Bing referrer\n(function() {\n var ref = document.referrer;\n var terms = [];\n \n if (ref.includes('google.com') || ref.includes('duckduckgo.com') || ref.includes('bing.com')) {\n var url = new URL(ref);\n var q = url.searchParams.get('q') || url.searchParams.get('p');\n if (q) {\n terms = q.split(/\\s+/).filter(function(t) { return t.length > 2; });\n }\n }\n \n if (terms.length === 0) return;\n \n var style = document.createElement('style');\n style.textContent = '.userscript-highlight { background: #fbbf24; color: #1a1a2e; padding: 1px 3px; border-radius: 2px; }';\n document.head.appendChild(style);\n \n function highlight(node) {\n if (node.nodeType === 3) { // text node\n var text = node.textContent;\n var found = false;\n terms.forEach(function(term) {\n var regex = new RegExp('(' + term.replace(/[.*+?^${}()|[\\]\\\\]/g, '\\\\') + ')', 'gi');\n if (regex.test(text)) {\n found = true;\n var frag = document.createDocumentFragment();\n var parts = text.split(regex);\n parts.forEach(function(part, i) {\n if (i % 2 === 0) {\n frag.appendChild(document.createTextNode(part));\n } else {\n var span = document.createElement('span');\n span.className = 'userscript-highlight';\n span.textContent = part;\n frag.appendChild(span);\n }\n });\n node.parentNode.replaceChild(frag, node);\n }\n });\n } else if (node.nodeType === 1 && node.childNodes) { // element\n var skipTags = ['SCRIPT', 'STYLE', 'NOSCRIPT', 'TEXTAREA', 'INPUT', 'SELECT'];\n if (!skipTags.includes(node.tagName)) {\n Array.from(node.childNodes).forEach(highlight);\n }\n }\n }\n \n highlight(document.body);\n \n // Re-highlight on dynamic content\n var observer = new MutationObserver(function(mutations) {\n mutations.forEach(function(m) {\n m.addedNodes.forEach(function(node) {\n if (node.nodeType === 1 || node.nodeType === 3) highlight(node);\n });\n });\n });\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Highlight Search Terms");
}
} catch(__e) { console.warn('[Userscript:Highlight Search Terms]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[TOPI] improve inclusive_scan for thrust - #16652
Merged
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Mar 3, 2024
Merged
[TOPI] improve inclusive_scan for thrust#16652
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Conversation
@yongwww
@yongwwwyongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
The topi.cuda.inclusive_scan currently relies on performing an exclusive_scan followed by an add operation that adds the input data back in. To eliminate the overhead introduced by this extra addition, probably we should have an implementation specifically designed for inclusive_scan. As of now, the implementation of inclusive_scan is limited to exclusive_scan_ir, developing an inclusive_scan_ir will be needed for performance in the future.
In this pull request, we specifically address this efficiency for thrust by directly calling the inclusive function.
Co-authored-by: Wuwei Lin wuwei@apache.org
@tqchen
tqchen
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
@tvm-bot rerun

1 similar comment

@yongwww
yongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
vinx13
vinx13
reviewed
Feb 28, 2024
View reviewed changes
Comment threadpython/tvm/topi/cuda/scan.py
Outdated

 data, output_dtype, exclusive=True, return_reduction=return_reduction, binop=binop
 )
 if _can_use_scan_thrust(binop):
 return scan_thrust(data, output_dtype, return_reduction=return_reduction, binop=binop)
@vinx13vinx13Feb 28, 2024

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

let's pass exclusive=True explicitly here
@yongwww
yongwww
commented
Feb 28, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
With this example:
 @I.ir_module
class Mod:
@R.function
def foo(x: R.Tensor((20, 32000), "float32")):
with R.dataflow():
gv = R.cumsum(x, axis=1)
R.output(gv)
return gv
target = tvm.target.Target("cuda -libs=thrust", host="llvm")
dev = tvm.cuda(0)
ex = relax.build(Mod, target)
vm = relax.VirtualMachine(ex, device=dev)

The 'cuda_gpu_kern_sum' stats report
w/ this change
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
 92.8 18,099,352 1,101 16,439.0 16,384.0 14,944 18,336 433.8 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
7.2 1,409,748 1,101 1,280.4 1,280.0 1,151 1,568 120.1 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

w/o this change
 51.6 22,603,888 1,101 20,530.3 20,545.0 19,744 21,120 293.0 cumsum_kernel
44.3 19,442,725 1,101 17,659.2 17,665.0 16,641 18,752 294.6 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
4.1 1,800,850 1,101 1,635.6 1,632.0 1,440 1,792 35.9 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

The execution perf numbers I got on NVIDIA GeForce RTX 3070:
R.cumsum(x, axis=1)
DescriptionPerformance
With this change (w/ thrust)0.018 ms
Without this change (w/ thrust)0.039 ms
Without thrust0.460 ms
R.cumsum(x, axis=0)
DescriptionPerformance
With this change (w/ thrust)0.073 ms
Without this change (w/ thrust)0.074 ms
Without thrust44.018 ms
MasterJH5574
MasterJH5574
approved these changes
Feb 28, 2024
View reviewed changes
@MasterJH5574MasterJH5574
left a comment


Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you @yongwww!!
@yongwww
yongwww
commented
Feb 29, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
The failed test tests/python/relax/test_op_unpack.py::test_tensor_dtype_lanes[bfloat] occurred in the last several merged commits, it's supposed to be fixed with #16649, is unrelated to this PR. I can rebase and have a clean run if necessary. The cpu failed with flaky test, it was green in https://ci.tlcpack.ai/blue/organizations/jenkins/tvm-cpu/detail/PR-16652/3/pipeline/
@lhutton1lhutton1
mentioned this pull request
Feb 29, 2024

Closed

@yongwww
yongwww
commented
Feb 29, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
@yongwww
yongwww
commented
Mar 1, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun

1 similar comment

@MasterJH5574
MasterJH5574
commented
Mar 2, 2024

Copy link


Copy Markdown

Contributor
@tvm-bot rerun
@tqchen
tqchen
merged commit 3b25588
into
apache:mainMar 3, 2024
@yongwww
yongwww
deleted the

topi_cumsum_update

branch
March 4, 2024 00:02
Lunderberg pushed a commit
to Lunderberg/tvm
that referenced
this pull request
Mar 12, 2024
@yongwww@Lunderberg
[TOPI] improve inclusive_scan for thrust (apache#16652)
7435cb7
Fix comments
thaisacs pushed a commit
to thaisacs/tvm
that referenced
this pull request
Apr 3, 2024
@yongwww@thaisacs
[TOPI] improve inclusive_scan for thrust (apache#16652)
8170cb1
Fix comments
@ysh329ysh329
mentioned this pull request
Apr 21, 2024

Closed

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@vinx13vinx13vinx13 approved these changes
@MasterJH5574MasterJH5574MasterJH5574 approved these changes

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@yongwww@tqchen@MasterJH5574@vinx13

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Strip utm_, fbclid, gclid, etc. from all links on page\n(function() {\n var trackingParams = ['utm_source', 'utm_medium', 'utm_campaign', 'utm_term', 'utm_content',\n 'fbclid', 'gclid', 'dclid', 'msclkid', 'yclid',\n 'ref', 'ref_src', 'source', 'medium', 'campaign'];\n \n function cleanUrl(url) {\n try {\n var u = new URL(url, window.location.origin);\n var changed = false;\n trackingParams.forEach(function(p) {\n if (u.searchParams.has(p)) {\n u.searchParams.delete(p);\n changed = true;\n }\n });\n return changed ? u.toString() : url;\n } catch (e) {\n return url;\n }\n }\n \n function cleanLinks() {\n document.querySelectorAll('a[href]').forEach(function(a) {\n var clean = cleanUrl(a.href);\n if (clean !== a.href) a.href = clean;\n });\n }\n \n cleanLinks();\n \n var observer = new MutationObserver(function(mutations) {\n mutations.forEach(function(m) {\n m.addedNodes.forEach(function(node) {\n if (node.nodeType === 1) {\n if (node.tagName === 'A') cleanLinks();\n node.querySelectorAll('a[href]').forEach(function(a) {\n var clean = cleanUrl(a.href);\n if (clean !== a.href) a.href = clean;\n });\n }\n });\n });\n });\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Remove Tracking Parameters from Links");
}
} catch(__e) { console.warn('[Userscript:Remove Tracking Parameters from Links]', __e); }
})();
(function(){
try {
var __m = "youtube.com";
var __re = new RegExp('^' + "youtube\\.com" + '

Skip to content
[TOPI] improve inclusive_scan for thrust - #16652
Merged
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Mar 3, 2024
Merged
[TOPI] improve inclusive_scan for thrust#16652
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Conversation
@yongwww
@yongwwwyongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
The topi.cuda.inclusive_scan currently relies on performing an exclusive_scan followed by an add operation that adds the input data back in. To eliminate the overhead introduced by this extra addition, probably we should have an implementation specifically designed for inclusive_scan. As of now, the implementation of inclusive_scan is limited to exclusive_scan_ir, developing an inclusive_scan_ir will be needed for performance in the future.
In this pull request, we specifically address this efficiency for thrust by directly calling the inclusive function.
Co-authored-by: Wuwei Lin wuwei@apache.org
@tqchen
tqchen
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
@tvm-bot rerun

1 similar comment

@yongwww
yongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
vinx13
vinx13
reviewed
Feb 28, 2024
View reviewed changes
Comment threadpython/tvm/topi/cuda/scan.py
Outdated

 data, output_dtype, exclusive=True, return_reduction=return_reduction, binop=binop
 )
 if _can_use_scan_thrust(binop):
 return scan_thrust(data, output_dtype, return_reduction=return_reduction, binop=binop)
@vinx13vinx13Feb 28, 2024

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

let's pass exclusive=True explicitly here
@yongwww
yongwww
commented
Feb 28, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
With this example:
 @I.ir_module
class Mod:
@R.function
def foo(x: R.Tensor((20, 32000), "float32")):
with R.dataflow():
gv = R.cumsum(x, axis=1)
R.output(gv)
return gv
target = tvm.target.Target("cuda -libs=thrust", host="llvm")
dev = tvm.cuda(0)
ex = relax.build(Mod, target)
vm = relax.VirtualMachine(ex, device=dev)

The 'cuda_gpu_kern_sum' stats report
w/ this change
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
 92.8 18,099,352 1,101 16,439.0 16,384.0 14,944 18,336 433.8 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
7.2 1,409,748 1,101 1,280.4 1,280.0 1,151 1,568 120.1 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

w/o this change
 51.6 22,603,888 1,101 20,530.3 20,545.0 19,744 21,120 293.0 cumsum_kernel
44.3 19,442,725 1,101 17,659.2 17,665.0 16,641 18,752 294.6 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
4.1 1,800,850 1,101 1,635.6 1,632.0 1,440 1,792 35.9 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

The execution perf numbers I got on NVIDIA GeForce RTX 3070:
R.cumsum(x, axis=1)
DescriptionPerformance
With this change (w/ thrust)0.018 ms
Without this change (w/ thrust)0.039 ms
Without thrust0.460 ms
R.cumsum(x, axis=0)
DescriptionPerformance
With this change (w/ thrust)0.073 ms
Without this change (w/ thrust)0.074 ms
Without thrust44.018 ms
MasterJH5574
MasterJH5574
approved these changes
Feb 28, 2024
View reviewed changes
@MasterJH5574MasterJH5574
left a comment


Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you @yongwww!!
@yongwww
yongwww
commented
Feb 29, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
The failed test tests/python/relax/test_op_unpack.py::test_tensor_dtype_lanes[bfloat] occurred in the last several merged commits, it's supposed to be fixed with #16649, is unrelated to this PR. I can rebase and have a clean run if necessary. The cpu failed with flaky test, it was green in https://ci.tlcpack.ai/blue/organizations/jenkins/tvm-cpu/detail/PR-16652/3/pipeline/
@lhutton1lhutton1
mentioned this pull request
Feb 29, 2024

Closed

@yongwww
yongwww
commented
Feb 29, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
@yongwww
yongwww
commented
Mar 1, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun

1 similar comment

@MasterJH5574
MasterJH5574
commented
Mar 2, 2024

Copy link


Copy Markdown

Contributor
@tvm-bot rerun
@tqchen
tqchen
merged commit 3b25588
into
apache:mainMar 3, 2024
@yongwww
yongwww
deleted the

topi_cumsum_update

branch
March 4, 2024 00:02
Lunderberg pushed a commit
to Lunderberg/tvm
that referenced
this pull request
Mar 12, 2024
@yongwww@Lunderberg
[TOPI] improve inclusive_scan for thrust (apache#16652)
7435cb7
Fix comments
thaisacs pushed a commit
to thaisacs/tvm
that referenced
this pull request
Apr 3, 2024
@yongwww@thaisacs
[TOPI] improve inclusive_scan for thrust (apache#16652)
8170cb1
Fix comments
@ysh329ysh329
mentioned this pull request
Apr 21, 2024

Closed

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@vinx13vinx13vinx13 approved these changes
@MasterJH5574MasterJH5574MasterJH5574 approved these changes

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@yongwww@tqchen@MasterJH5574@vinx13

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Auto-enable theater mode on YouTube\n(function() {\n function tryTheater() {\n var btn = document.querySelector('button[aria-label=\"Theater mode\"], ytd-player #player button[title=\"Theater mode\"]');\n if (btn && !btn.classList.contains('activated')) {\n btn.click();\n }\n }\n \n // Try immediately\n tryTheater();\n \n // Try after navigation (SPA)\n var lastUrl = location.href;\n setInterval(function() {\n if (location.href !== lastUrl) {\n lastUrl = location.href;\n setTimeout(tryTheater, 500);\n }\n }, 1000);\n \n // Also try on player load\n var observer = new MutationObserver(tryTheater);\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "YouTube Theater Mode Default");
}
} catch(__e) { console.warn('[Userscript:YouTube Theater Mode Default]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[TOPI] improve inclusive_scan for thrust - #16652
Merged
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Mar 3, 2024
Merged
[TOPI] improve inclusive_scan for thrust#16652
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Conversation
@yongwww
@yongwwwyongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
The topi.cuda.inclusive_scan currently relies on performing an exclusive_scan followed by an add operation that adds the input data back in. To eliminate the overhead introduced by this extra addition, probably we should have an implementation specifically designed for inclusive_scan. As of now, the implementation of inclusive_scan is limited to exclusive_scan_ir, developing an inclusive_scan_ir will be needed for performance in the future.
In this pull request, we specifically address this efficiency for thrust by directly calling the inclusive function.
Co-authored-by: Wuwei Lin wuwei@apache.org
@tqchen
tqchen
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
@tvm-bot rerun

1 similar comment

@yongwww
yongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
vinx13
vinx13
reviewed
Feb 28, 2024
View reviewed changes
Comment threadpython/tvm/topi/cuda/scan.py
Outdated

 data, output_dtype, exclusive=True, return_reduction=return_reduction, binop=binop
 )
 if _can_use_scan_thrust(binop):
 return scan_thrust(data, output_dtype, return_reduction=return_reduction, binop=binop)
@vinx13vinx13Feb 28, 2024

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

let's pass exclusive=True explicitly here
@yongwww
yongwww
commented
Feb 28, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
With this example:
 @I.ir_module
class Mod:
@R.function
def foo(x: R.Tensor((20, 32000), "float32")):
with R.dataflow():
gv = R.cumsum(x, axis=1)
R.output(gv)
return gv
target = tvm.target.Target("cuda -libs=thrust", host="llvm")
dev = tvm.cuda(0)
ex = relax.build(Mod, target)
vm = relax.VirtualMachine(ex, device=dev)

The 'cuda_gpu_kern_sum' stats report
w/ this change
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
 92.8 18,099,352 1,101 16,439.0 16,384.0 14,944 18,336 433.8 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
7.2 1,409,748 1,101 1,280.4 1,280.0 1,151 1,568 120.1 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

w/o this change
 51.6 22,603,888 1,101 20,530.3 20,545.0 19,744 21,120 293.0 cumsum_kernel
44.3 19,442,725 1,101 17,659.2 17,665.0 16,641 18,752 294.6 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
4.1 1,800,850 1,101 1,635.6 1,632.0 1,440 1,792 35.9 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

The execution perf numbers I got on NVIDIA GeForce RTX 3070:
R.cumsum(x, axis=1)
DescriptionPerformance
With this change (w/ thrust)0.018 ms
Without this change (w/ thrust)0.039 ms
Without thrust0.460 ms
R.cumsum(x, axis=0)
DescriptionPerformance
With this change (w/ thrust)0.073 ms
Without this change (w/ thrust)0.074 ms
Without thrust44.018 ms
MasterJH5574
MasterJH5574
approved these changes
Feb 28, 2024
View reviewed changes
@MasterJH5574MasterJH5574
left a comment


Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you @yongwww!!
@yongwww
yongwww
commented
Feb 29, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
The failed test tests/python/relax/test_op_unpack.py::test_tensor_dtype_lanes[bfloat] occurred in the last several merged commits, it's supposed to be fixed with #16649, is unrelated to this PR. I can rebase and have a clean run if necessary. The cpu failed with flaky test, it was green in https://ci.tlcpack.ai/blue/organizations/jenkins/tvm-cpu/detail/PR-16652/3/pipeline/
@lhutton1lhutton1
mentioned this pull request
Feb 29, 2024

Closed

@yongwww
yongwww
commented
Feb 29, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
@yongwww
yongwww
commented
Mar 1, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun

1 similar comment

@MasterJH5574
MasterJH5574
commented
Mar 2, 2024

Copy link


Copy Markdown

Contributor
@tvm-bot rerun
@tqchen
tqchen
merged commit 3b25588
into
apache:mainMar 3, 2024
@yongwww
yongwww
deleted the

topi_cumsum_update

branch
March 4, 2024 00:02
Lunderberg pushed a commit
to Lunderberg/tvm
that referenced
this pull request
Mar 12, 2024
@yongwww@Lunderberg
[TOPI] improve inclusive_scan for thrust (apache#16652)
7435cb7
Fix comments
thaisacs pushed a commit
to thaisacs/tvm
that referenced
this pull request
Apr 3, 2024
@yongwww@thaisacs
[TOPI] improve inclusive_scan for thrust (apache#16652)
8170cb1
Fix comments
@ysh329ysh329
mentioned this pull request
Apr 21, 2024

Closed

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@vinx13vinx13vinx13 approved these changes
@MasterJH5574MasterJH5574MasterJH5574 approved these changes

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@yongwww@tqchen@MasterJH5574@vinx13

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Remove or un-stick sticky/fixed headers that block content\n(function() {\n function unstick() {\n document.querySelectorAll('header, nav, [role=\"banner\"], .header, .navbar, .sticky, .fixed-top, [style*=\"position: fixed\"], [style*=\"position:sticky\"]').forEach(function(el) {\n if (el.style.position === 'fixed' || el.style.position === 'sticky' || \n getComputedStyle(el).position === 'fixed' || getComputedStyle(el).position === 'sticky') {\n el.style.position = 'static';\n el.style.top = 'auto';\n el.style.zIndex = 'auto';\n }\n });\n }\n \n unstick();\n \n var observer = new MutationObserver(unstick);\n observer.observe(document.body, { childList: true, subtree: true, attributes: true, attributeFilter: ['style', 'class'] });\n})();", "Kill Sticky Headers");
}
} catch(__e) { console.warn('[Userscript:Kill Sticky Headers]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[TOPI] improve inclusive_scan for thrust - #16652
Merged
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Mar 3, 2024
Merged
[TOPI] improve inclusive_scan for thrust#16652
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Conversation
@yongwww
@yongwwwyongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
The topi.cuda.inclusive_scan currently relies on performing an exclusive_scan followed by an add operation that adds the input data back in. To eliminate the overhead introduced by this extra addition, probably we should have an implementation specifically designed for inclusive_scan. As of now, the implementation of inclusive_scan is limited to exclusive_scan_ir, developing an inclusive_scan_ir will be needed for performance in the future.
In this pull request, we specifically address this efficiency for thrust by directly calling the inclusive function.
Co-authored-by: Wuwei Lin wuwei@apache.org
@tqchen
tqchen
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
@tvm-bot rerun

1 similar comment

@yongwww
yongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
vinx13
vinx13
reviewed
Feb 28, 2024
View reviewed changes
Comment threadpython/tvm/topi/cuda/scan.py
Outdated

 data, output_dtype, exclusive=True, return_reduction=return_reduction, binop=binop
 )
 if _can_use_scan_thrust(binop):
 return scan_thrust(data, output_dtype, return_reduction=return_reduction, binop=binop)
@vinx13vinx13Feb 28, 2024

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

let's pass exclusive=True explicitly here
@yongwww
yongwww
commented
Feb 28, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
With this example:
 @I.ir_module
class Mod:
@R.function
def foo(x: R.Tensor((20, 32000), "float32")):
with R.dataflow():
gv = R.cumsum(x, axis=1)
R.output(gv)
return gv
target = tvm.target.Target("cuda -libs=thrust", host="llvm")
dev = tvm.cuda(0)
ex = relax.build(Mod, target)
vm = relax.VirtualMachine(ex, device=dev)

The 'cuda_gpu_kern_sum' stats report
w/ this change
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
 92.8 18,099,352 1,101 16,439.0 16,384.0 14,944 18,336 433.8 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
7.2 1,409,748 1,101 1,280.4 1,280.0 1,151 1,568 120.1 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

w/o this change
 51.6 22,603,888 1,101 20,530.3 20,545.0 19,744 21,120 293.0 cumsum_kernel
44.3 19,442,725 1,101 17,659.2 17,665.0 16,641 18,752 294.6 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
4.1 1,800,850 1,101 1,635.6 1,632.0 1,440 1,792 35.9 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

The execution perf numbers I got on NVIDIA GeForce RTX 3070:
R.cumsum(x, axis=1)
DescriptionPerformance
With this change (w/ thrust)0.018 ms
Without this change (w/ thrust)0.039 ms
Without thrust0.460 ms
R.cumsum(x, axis=0)
DescriptionPerformance
With this change (w/ thrust)0.073 ms
Without this change (w/ thrust)0.074 ms
Without thrust44.018 ms
MasterJH5574
MasterJH5574
approved these changes
Feb 28, 2024
View reviewed changes
@MasterJH5574MasterJH5574
left a comment


Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you @yongwww!!
@yongwww
yongwww
commented
Feb 29, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
The failed test tests/python/relax/test_op_unpack.py::test_tensor_dtype_lanes[bfloat] occurred in the last several merged commits, it's supposed to be fixed with #16649, is unrelated to this PR. I can rebase and have a clean run if necessary. The cpu failed with flaky test, it was green in https://ci.tlcpack.ai/blue/organizations/jenkins/tvm-cpu/detail/PR-16652/3/pipeline/
@lhutton1lhutton1
mentioned this pull request
Feb 29, 2024

Closed

@yongwww
yongwww
commented
Feb 29, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
@yongwww
yongwww
commented
Mar 1, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun

1 similar comment

@MasterJH5574
MasterJH5574
commented
Mar 2, 2024

Copy link


Copy Markdown

Contributor
@tvm-bot rerun
@tqchen
tqchen
merged commit 3b25588
into
apache:mainMar 3, 2024
@yongwww
yongwww
deleted the

topi_cumsum_update

branch
March 4, 2024 00:02
Lunderberg pushed a commit
to Lunderberg/tvm
that referenced
this pull request
Mar 12, 2024
@yongwww@Lunderberg
[TOPI] improve inclusive_scan for thrust (apache#16652)
7435cb7
Fix comments
thaisacs pushed a commit
to thaisacs/tvm
that referenced
this pull request
Apr 3, 2024
@yongwww@thaisacs
[TOPI] improve inclusive_scan for thrust (apache#16652)
8170cb1
Fix comments
@ysh329ysh329
mentioned this pull request
Apr 21, 2024

Closed

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@vinx13vinx13vinx13 approved these changes
@MasterJH5574MasterJH5574MasterJH5574 approved these changes

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@yongwww@tqchen@MasterJH5574@vinx13

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Universal Dark Mode - works on any site\n(function() {\n var enabled = true;\n \n function applyDarkMode() {\n if (!enabled) return;\n \n // Create style element if it doesn't exist\n var style = document.getElementById('universal-dark-mode-style');\n if (!style) {\n style = document.createElement('style');\n style.id = 'universal-dark-mode-style';\n document.head.appendChild(style);\n }\n \n // Dark mode CSS - inverts colors but preserves images/video\n style.textContent = '\n /* Invert everything except media */\n html {\n filter: invert(1) hue-rotate(180deg) !important;\n background: #1a1a2e !important;\n }\n \n /* Restore images, videos, iframes, canvas */\n img, video, iframe, canvas, svg, picture, [style*=\"background-image\"] {\n filter: invert(1) hue-rotate(180deg) !important;\n }\n \n /* Preserve specific elements that should not be inverted */\n .no-dark-mode, .no-dark-mode *,\n [data-theme=\"light\"], [data-theme=\"light\"],\n .ace_editor, .ace_editor *,\n .CodeMirror, .CodeMirror *,\n .monaco-editor, .monaco-editor *,\n .markdown-body pre, .markdown-body pre *,\n .highlight, .highlight *,\n pre code, pre code * {\n filter: none !important;\n }\n \n /* Fix common UI elements */\n .modal, .popup, .dropdown-menu, .tooltip, .popover {\n filter: invert(1) hue-rotate(180deg) !important;\n background: #2d2d44 !important;\n border-color: #444 !important;\n }\n \n /* Scrollbars */\n ::-webkit-scrollbar { background: #1a1a2e !important; }\n ::-webkit-scrollbar-thumb { background: #444 !important; }\n ::-webkit-scrollbar-thumb:hover { background: #555 !important; }\n \n /* Selection */\n ::selection { background: #4ecdc4 !important; color: #1a1a2e !important; }\n ::-moz-selection { background: #4ecdc4 !important; color: #1a1a2e !important; }\n ';\n }\n \n function removeDarkMode() {\n var style = document.getElementById('universal-dark-mode-style');\n if (style) style.remove();\n }\n \n // Toggle with Alt+Shift+D\n document.addEventListener('keydown', function(e) {\n if (e.altKey && e.shiftKey && e.key === 'D') {\n e.preventDefault();\n enabled = !enabled;\n if (enabled) {\n applyDarkMode();\n console.log('[Universal Dark Mode] Enabled');\n } else {\n removeDarkMode();\n console.log('[Universal Dark Mode] Disabled');\n }\n }\n });\n \n // Apply on load\n applyDarkMode();\n \n // Re-apply on dynamic content\n var observer = new MutationObserver(function(mutations) {\n if (enabled && !document.getElementById('universal-dark-mode-style')) {\n applyDarkMode();\n }\n });\n observer.observe(document.head, { childList: true });\n \n console.log('[Universal Dark Mode] Loaded - Press Alt+Shift+D to toggle');\n})();", "Universal Dark Mode");
}
} catch(__e) { console.warn('[Userscript:Universal Dark Mode]', __e); }
})();
})();

Skip to content
[TOPI] improve inclusive_scan for thrust - #16652
Merged
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Mar 3, 2024
Merged
[TOPI] improve inclusive_scan for thrust#16652
tqchen merged 1 commit into
apache:mainfrom 
yongwww:topi_cumsum_update
Conversation
@yongwww
@yongwwwyongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
The topi.cuda.inclusive_scan currently relies on performing an exclusive_scan followed by an add operation that adds the input data back in. To eliminate the overhead introduced by this extra addition, probably we should have an implementation specifically designed for inclusive_scan. As of now, the implementation of inclusive_scan is limited to exclusive_scan_ir, developing an inclusive_scan_ir will be needed for performance in the future.
In this pull request, we specifically address this efficiency for thrust by directly calling the inclusive function.
Co-authored-by: Wuwei Lin wuwei@apache.org
@tqchen
tqchen
commented
Feb 28, 2024

Copy link


Copy Markdown

Member
@tvm-bot rerun

1 similar comment

@yongwww
yongwww
commented
Feb 28, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
vinx13
vinx13
reviewed
Feb 28, 2024
View reviewed changes
Comment threadpython/tvm/topi/cuda/scan.py
Outdated

 data, output_dtype, exclusive=True, return_reduction=return_reduction, binop=binop
 )
 if _can_use_scan_thrust(binop):
 return scan_thrust(data, output_dtype, return_reduction=return_reduction, binop=binop)
@vinx13vinx13Feb 28, 2024

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

let's pass exclusive=True explicitly here
@yongwww
yongwww
commented
Feb 28, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
With this example:
 @I.ir_module
class Mod:
@R.function
def foo(x: R.Tensor((20, 32000), "float32")):
with R.dataflow():
gv = R.cumsum(x, axis=1)
R.output(gv)
return gv
target = tvm.target.Target("cuda -libs=thrust", host="llvm")
dev = tvm.cuda(0)
ex = relax.build(Mod, target)
vm = relax.VirtualMachine(ex, device=dev)

The 'cuda_gpu_kern_sum' stats report
w/ this change
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
 92.8 18,099,352 1,101 16,439.0 16,384.0 14,944 18,336 433.8 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
7.2 1,409,748 1,101 1,280.4 1,280.0 1,151 1,568 120.1 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

w/o this change
 51.6 22,603,888 1,101 20,530.3 20,545.0 19,744 21,120 293.0 cumsum_kernel
44.3 19,442,725 1,101 17,659.2 17,665.0 16,641 18,752 294.6 void cub::CUB_200200_750_NS::DeviceScanByKeyKernel<cub::CUB_200200_750_NS::DeviceScanByKeyPolicy<th…
4.1 1,800,850 1,101 1,635.6 1,632.0 1,440 1,792 35.9 void cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel<cub::CUB_200200_750_NS::ReduceByKeyScanTileS…

The execution perf numbers I got on NVIDIA GeForce RTX 3070:
R.cumsum(x, axis=1)
DescriptionPerformance
With this change (w/ thrust)0.018 ms
Without this change (w/ thrust)0.039 ms
Without thrust0.460 ms
R.cumsum(x, axis=0)
DescriptionPerformance
With this change (w/ thrust)0.073 ms
Without this change (w/ thrust)0.074 ms
Without thrust44.018 ms
MasterJH5574
MasterJH5574
approved these changes
Feb 28, 2024
View reviewed changes
@MasterJH5574MasterJH5574
left a comment


Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you @yongwww!!
@yongwww
yongwww
commented
Feb 29, 2024•

edited

Loading

Copy link


Copy Markdown

MemberAuthor
The failed test tests/python/relax/test_op_unpack.py::test_tensor_dtype_lanes[bfloat] occurred in the last several merged commits, it's supposed to be fixed with #16649, is unrelated to this PR. I can rebase and have a clean run if necessary. The cpu failed with flaky test, it was green in https://ci.tlcpack.ai/blue/organizations/jenkins/tvm-cpu/detail/PR-16652/3/pipeline/
@lhutton1lhutton1
mentioned this pull request
Feb 29, 2024

Closed

@yongwww
yongwww
commented
Feb 29, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun
@yongwww
yongwww
commented
Mar 1, 2024

Copy link


Copy Markdown

MemberAuthor
@tvm-bot rerun

1 similar comment

@MasterJH5574
MasterJH5574
commented
Mar 2, 2024

Copy link


Copy Markdown

Contributor
@tvm-bot rerun
@tqchen
tqchen
merged commit 3b25588
into
apache:mainMar 3, 2024
@yongwww
yongwww
deleted the

topi_cumsum_update

branch
March 4, 2024 00:02
Lunderberg pushed a commit
to Lunderberg/tvm
that referenced
this pull request
Mar 12, 2024
@yongwww@Lunderberg
[TOPI] improve inclusive_scan for thrust (apache#16652)
7435cb7
Fix comments
thaisacs pushed a commit
to thaisacs/tvm
that referenced
this pull request
Apr 3, 2024
@yongwww@thaisacs
[TOPI] improve inclusive_scan for thrust (apache#16652)
8170cb1
Fix comments
@ysh329ysh329
mentioned this pull request
Apr 21, 2024

Closed

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@vinx13vinx13vinx13 approved these changes
@MasterJH5574MasterJH5574MasterJH5574 approved these changes

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@yongwww@tqchen@MasterJH5574@vinx13