[hipDNN] Add SDPA backward operation descriptor and backend lowering/lifting#5674
Merged
SamuelReeder merged 10 commits intodevelopfrom Mar 25, 2026
Merged
Conversation
…tests Adds the SDPA backwards (bprop) operation descriptor to the C-API and wires backend lowering from the frontend node to the backend descriptor using the codegen tool. Changes: - New: SdpaBpropOperationDescriptor (.hpp/.cpp) — backend descriptor impl - New: SdpaBpropPacker.hpp — frontend-to-backend attribute packing - New: SdpaBpropUnpacker.hpp — backend-to-frontend attribute unpacking - New: SdpaBpropNode.hpp updated with create_operation/unpack overrides - New: SdpaBpropConstants.hpp — test SDK constants - New: TestSdpaBpropOperationDescriptor.cpp — unit tests (23 cases) - New: TestGraphDescriptorSdpaBprop.cpp — graph round-trip tests - New: IntegrationSdpaBpropDescriptorLowering.cpp — E2E lowering tests - Modified: HipdnnBackendAttributeName.h — SDPA bprop attribute range - Modified: HipdnnBackendDescriptorType.h — SDPA_BPROP_OPERATION descriptor type - Modified: BackendEnumStringUtils.hpp — string utils for new type/attrs - Modified: DescriptorFactory.cpp — factory case for SdpaBprop descriptor - Modified: NodeFactory.cpp/.hpp — node factory case for SdpaBprop - Modified: OperationUnpacker.hpp — unpacker dispatch case - Modified: CMakeLists.txt files — new sources and tests registered Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
SamuelReeder
commented
Mar 20, 2026
Codecov Report❌ Patch coverage is ❌ Your project status has failed because the head coverage (77.21%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #5674 +/- ##
===========================================
+ Coverage 67.02% 67.06% +0.03%
===========================================
Files 1853 1857 +4
Lines 286550 287515 +965
Branches 40271 40392 +121
===========================================
+ Hits 192053 192795 +742
- Misses 78033 78226 +193
- Partials 16464 16494 +30
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
Fix the SdpaBpropPacker to pack the operation name (was silently dropped). Fix UID collisions in SdpaBpropConstants (50-69 -> 90-109 to avoid overlap with BlockScaleDequantize). Comprehensive test coverage improvements: - Integration tests: verify all tensor properties (UID, dims, strides, data type, name), all scalar/boolean/enum fields, optional tensor presence/absence. Add AutoAssignedUids and OperationNameRoundTrip tests. - Backend descriptor tests: add fromNode() coverage (required-only, all fields, name preservation, missing required, absent optional in map), getAttribute() for all attributes (set/unset), setAttribute() negative cases (wrong type/count/null/unfinalized), graph serialization with exact tensor counts and full field verification. Test count: 5307 -> 5378 (71 new tests), all passing. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
SamuelReeder
commented
Mar 20, 2026
SamuelReeder
commented
Mar 20, 2026
SamuelReeder
commented
Mar 20, 2026
SamuelReeder
changed the title
- Add HIPDNN_ATTR_OPERATION_TYPE_EXT to SdpaBpropOperationDescriptor::getAttribute() so the lifting path can determine operation type during graph reconstruction - Add IntegrationSdpaBpropLifting integration tests covering C-API round-trip, tensor sharing, serialization without finalization, attn scale, attn mask, and graph name preservation - Add CreatesSdpaBpropNode test to TestOperationUnpacker - Improve GetTensorDescriptors and GetTensorDescriptorsRequiredOnly to verify individual tensor dims and strides, not just UID presence - Fix test constant UID collisions (shifted to 90-109 range) - Improve integration tests: verify all optional tensors survive round-trip, use for loop in GetOptionalTensorUnsetReturnsCount0 Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
…iopen-1589/add-sdpa-backward-descriptor
SamuelReeder
commented
Mar 23, 2026
SamuelReeder
commented
Mar 23, 2026
SamuelReeder
commented
Mar 23, 2026
BrianHarrisonAMD
approved these changes
Mar 23, 2026
Contributor
BrianHarrisonAMD
left a comment
BrianHarrisonAMD
left a comment
There was a problem hiding this comment.
Approving, but we should fix comments before merging.
…iopen-1589/add-sdpa-backward-descriptor # Conflicts: # projects/hipdnn/backend/tests/descriptors/TestNodeFactory.cpp
Renumber SDPA bprop attribute range to 3100-3199 (CustomOp claims 3000-3099) and descriptor type to 30 (CustomOp claims 29). Keep both in BackendEnumStringUtils and DescriptorFactory. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
SamuelReeder
deleted the
users/sareeder/almiopen-1589/add-sdpa-backward-descriptor
branch
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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
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.
Motivation
SDPA (Scaled Dot-Product Attention) forward already has a C-API descriptor and backend lowering. This PR adds the corresponding backwards (bprop) path, enabling frontend
SdpaBpropNodeto be lowered to a backendSDPA_BPROP_OPERATIONdescriptor via the C-API.Technical Details
Generated and integrated via the hipDNN codegen tool. Adds
SdpaBpropOperationDescriptorimplementingsetAttribute/getAttribute/buildNode/finalizefor all SDPA backward tensor and data fields (Q, K, V, O, dO, dQ, dK, dV, stats, attn_scale, dropout probability/seed/offset, causal mask flag).SdpaBpropPacker.hpppacks frontend node attributes into the backend descriptor, whileSdpaBpropUnpacker.hppreconstructs the frontend node from a backend descriptor. The existingSdpaBpropNodeis wired withcreate_operation()andunpack_from_descriptor()overrides.Infrastructure changes include a new attribute range in
HipdnnBackendAttributeName.h, aHIPDNN_BACKEND_SDPA_BPROP_OPERATION_DESCRIPTORtype inHipdnnBackendDescriptorType.h, and dispatch cases inDescriptorFactory.cpp,NodeFactory.cpp, andOperationUnpacker.hpp. String utility coverage is added inBackendEnumStringUtils.hpp.Test Plan
TestSdpaBpropOperationDescriptor.cpp— 23 unit tests covering all attributes (set/get, finalize validation)TestGraphDescriptorSdpaBprop.cpp— graph-level round-trip tests (serialize → deserialize → verify)IntegrationSdpaBpropDescriptorLowering.cpp— 2 E2E frontend-to-backend lowering tests (explicit UIDs + auto-assigned UIDs)TestBackendEnumStringUtils.cpp— string coverage for new descriptor type and attributesninja unit-check— 5307 tests passSubmission Checklist