feat: Remove more unnecessary PTX code from Wave #12380
Conversation
facebook-github-bot
added
the
CLA Signed
dreveman
requested review from
Yuhta
and removed request for
majetideepak and
assignUser
✅ Deploy Preview for meta-velox canceled.
|
dreveman
changed the title
Yuhta
added
the
ready-to-merge
dreveman
force-pushed
the
rm-ptx-2
branch
2 times, most recently
from
69ee8b3 to
540a750
Compare
|
@Yuhta any reason for why this is taking so long to be merged? |
|
@mbasmanova has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
|
@dreveman Seeing format error: @kgpai @majetideepak @assignUser Wondering why this doesn't show up in CI. |
|
@mbasmanova Likely a clang-format version difference. Can you share the version being used internally? |
|
@majetideepak clang-format version 18.1.3 |
That's what's in our container as well. I think the reason is that we only run the check on touched files and the file in question wasn't touched by this PR. Maybe the internal linter works on a dir basis and not a file basis or something? |
|
@assignUser velox/experimental/wave/common/Cuda.cu is modified in this PR. Do we run format-check on .cu files? |
|
Yes now it is but only to fix the lint. Or did I have a weird diff view? |
|
Oh sorry I read the linter wrong, without space is correct. Hm, let me check. |
Shouldn't impact performance. If anything, might be faster as threadIdx.x is likely in a register already and %laneid requires a S2R instruction that can have a relatively long latency compared to a regular register.
Add build option that disables Wave branch kernel test as it requires PTX. This test is enabled by default and can be disabled with: cmake -DVELOX_SKIP_WAVE_BRANCH_KERNEL_TEST=ON
|
Sorry, not sure how that format change become part of that commit. Must have run the formatting script locally with the wrong clang-format version. Rebased and removed that format change from latest version of PR. |
|
When I add the space locally and run But two spaces are corrected to a single one. So this must be an issue with our clang-format file. |
|
@mbasmanova has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
|
I looked at the docs and tested something but was unable to pin down the issue. My local lsp (using the same clang format version) enforces one space after the ctor initializer list and uses the same .clang-format file as @mbasmanova maybe you can diff you internal clang-format file with the public one in the repo to find the misaligned setting. But the manual fix works for now too I guess ^^ |
Krishna, is this something you could help with? |
|
@mbasmanova merged this pull request in 0d8c05a. |
No description provided.