Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[BUG]: Validate tuning refactoring on all affected SM versions #3235

Closed
1 task done
gevtushenko opened this issue Jan 2, 2025 · 3 comments
Closed
1 task done

[BUG]: Validate tuning refactoring on all affected SM versions #3235

gevtushenko opened this issue Jan 2, 2025 · 3 comments
Assignees
Labels
bug Something isn't working right.

Comments

@gevtushenko
Copy link
Collaborator

Is this a duplicate?

Type of Bug

Performance

Component

CUB

Describe the bug

#3138 introduced 120% performance regression in cub::DeviceScan::Sum on Hopper.
Looking at PR description:

No changes in SASS for cub.test.device_scan.lid_0.types_0 except kernel symbol names

I assume that we checked SASS only on default set of architectures from our preset "60;70;80".

diff old.sass new.sass | wc -l 
88

Mentioned refactoring also touches SM90 policies, which results in the following diff:

diff old.90.sass new.90.sass | wc -l 
83618

We should verify SM90 SASS differences on other algorithms we refactored lately.

How to Reproduce

git checkout 52b4b671c111fc289b2f2247a16eff21d5a5a1a8
./ci/build_cub.sh -arch 90
cuobjdump --dump-sass build/cuda12.6-gcc13/cub-cpp17/bin/cub.bench.scan.exclusive.sum.base > new.90.sass
git checkout HEAD~
./ci/build_cub.sh -arch 90
cuobjdump --dump-sass build/cuda12.6-gcc13/cub-cpp17/bin/cub.bench.scan.exclusive.sum.base > old.90.sass
diff old.90.sass new.90.sass | wc -l 
83618

Expected behavior

No SASS difference on any architecture after tuning refactoring

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

@gevtushenko gevtushenko added the bug Something isn't working right. label Jan 2, 2025
@github-project-automation github-project-automation bot moved this to Todo in CCCL Jan 2, 2025
@gevtushenko gevtushenko self-assigned this Jan 2, 2025
@gevtushenko gevtushenko moved this from Todo to In Progress in CCCL Jan 2, 2025
@gevtushenko
Copy link
Collaborator Author

gevtushenko commented Jan 3, 2025

Refactoring Regression on SM90
#3138 Fixed by #3236
#3124 No
#3127 Fixed by #3239
#3125 Fixed by #3174
#3128 No
#3137 Fixed by #3240
#3139 No
#3140 No
#3145 No

@github-project-automation github-project-automation bot moved this from In Progress to Done in CCCL Jan 3, 2025
@bernhardmgruber
Copy link
Contributor

I assume that we checked SASS only on default set of architectures from our preset "60;70;80".

It may be even worse. I cannot find what exact cmake invocation I used, but it could have been just 86. I will revisit your table with all architectures touched by each PR.

@bernhardmgruber bernhardmgruber changed the title [BUG]: Validate tuning refactoring on SM90 [BUG]: Validate tuning refactoring on all affected architectures Jan 7, 2025
@bernhardmgruber bernhardmgruber changed the title [BUG]: Validate tuning refactoring on all affected architectures [BUG]: Validate tuning refactoring on all affected SM versions Jan 7, 2025
@bernhardmgruber bernhardmgruber self-assigned this Jan 7, 2025
@gevtushenko
Copy link
Collaborator Author

I will revisit your table with all architectures touched by each PR.

That's what I did to prepare the table above

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working right.
Projects
Status: Done
Development

No branches or pull requests

2 participants