Argmax(cpp wrapper) #784
Merged
Argmax(cpp wrapper) #784
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Bowen12992
reviewed
Jul 15, 2025
ctests/test_triton_argmax.cpp
Outdated
| #include "flag_gems/operators.h" | ||
| #include "torch/torch.h" | ||
|
|
||
| TEST(reduction_op_test, argmax) { |
There was a problem hiding this comment.
It's better to test with larger shapes and with different dtypes.
Bowen12992
reviewed
Jul 15, 2025
ctests/test_triton_argmax.cpp
Outdated
|
|
||
| TEST(reduction_op_test, argmax_keepdim_option) { | ||
| const torch::Device device(torch::kCUDA, 0); | ||
| torch::Tensor input = torch::randn({2, 2, 2, 2}, device); |
There was a problem hiding this comment.
ditto, the test shape is too small
Bowen12992
reviewed
Jul 15, 2025
lib/argmax.cpp
Outdated
Comment on lines
29
to
32
| auto shape = self.sizes().vec(); | ||
| for (auto &s : shape) { | ||
| s = 1; | ||
| } |
There was a problem hiding this comment.
Suggested change
| auto shape = self.sizes().vec(); | |
| for (auto &s : shape) { | |
| s = 1; | |
| } | |
| const auto shape = std::vector<int64_t>(self.dim(), 1); |
Bowen12992
reviewed
Jul 15, 2025
lib/argmax.cpp
Outdated
| c10::DeviceGuard guard(self.device()); | ||
| c10::cuda::CUDAStream stream = c10::cuda::getCurrentCUDAStream(); | ||
|
|
||
| f1(stream, mid_size, 1, 1, 4 /*num_warps*/, 2 /*num_stages*/, self, mid_value, mid_index, M, block_size); |
There was a problem hiding this comment.
Suggested change
| f1(stream, mid_size, 1, 1, 4 /*num_warps*/, 2 /*num_stages*/, self, mid_value, mid_index, M, block_size); | |
| f1(stream, mid_size, 1, 1, /* num_warps = */ 4, /* num_stages = */ 2 , self, mid_value, mid_index, M, block_size); |
Bowen12992
reviewed
Jul 15, 2025
lib/argmax.cpp
Outdated
|
|
||
| f1(stream, mid_size, 1, 1, 4 /*num_warps*/, 2 /*num_stages*/, self, mid_value, mid_index, M, block_size); | ||
|
|
||
| f2(stream, 1, 1, 1, 4 /*num_warps*/, 2 /*num_stages*/, mid_value, mid_index, out, mid_size, block_mid); |
Bowen12992
reviewed
Jul 15, 2025
lib/argmax.cpp
Outdated
| int64_t dim_val = dim.value(); | ||
| dim_val = at::maybe_wrap_dim(dim_val, self.dim()); | ||
|
|
||
| auto shape = self.sizes(); |
There was a problem hiding this comment.
Suggested change
| auto shape = self.sizes(); | |
| const auto& shape = self.sizes(); |
Bowen12992
previously approved these changes
Jul 15, 2025
added 2 commits
July 15, 2025 15:27
0x45f
reviewed
Jul 16, 2025
Bowen12992
approved these changes
Jul 16, 2025
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
PR Category
Operator
Type of Change
Refactor
Description
cpp wrapper
Issue
Progress
Performance