add backward of conv2d #365
Open
Conversation
This file contains 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
Galaxy1458
reviewed
Dec 26, 2024
Galaxy1458
reviewed
Dec 26, 2024
src/flag_gems/ops/conv2d.py
Outdated
| @@ -230,7 +228,7 @@ def conv2d_forward_kernel( | |||
| @triton.autotune( | |||
| configs=[ | |||
| triton.Config( | |||
There was a problem hiding this comment.
Please config these at FlagGems/src/flag_gems/runtime/backend/_nvidia/tune_configs.yaml
StrongSpoon
reviewed
Dec 26, 2024
| @pytest.mark.parametrize("dilation", [1, 2]) | ||
| @pytest.mark.parametrize("bias", [True, False]) | ||
| def test_accuracy_conv2d(shape, kernel, stride, padding, groups, dtype, dilation, bias): | ||
| torch.manual_seed(0) |
There was a problem hiding this comment.
is manual_seed necessary here?
| revert_weight = revert_weight.transpose(1, 2).contiguous() | ||
| revert_weight = revert_weight.reshape( | ||
| groups * weight_c, out_c, weight_height, weight_width | ||
| ).contiguous() |
There was a problem hiding this comment.
redundant contiguous might waste resources.
| if stride_height > 1 or stride_width > 1: | ||
| for i in range(out_grad.shape[2]): | ||
| for j in range(out_grad.shape[3]): | ||
| new_out[:, :, i * (stride_height), j * (stride_width)] = out_grad[ |
There was a problem hiding this comment.
this assignment will cost a lot of time. is there a better way?
There was a problem hiding this comment.
maybe you can reference to the implementation in flip and use a copy_func to fill the elements in new_out.
| device=device, | ||
| ) | ||
|
|
||
| grid_weight = lambda meta: ( |
There was a problem hiding this comment.
since size of weight is generally not large, I suggest not tiling them by BLOCK_CI_HK_WK.
added 2 commits
January 13, 2025 11:39
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.
add backwards of input weight bias for conv.