Enable Python EVT Compute UT #650
Open
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
There was a problem hiding this comment.
Pull request overview
This PR enables Epilogue Visitor Tree (EVT) Python unit tests for Intel XPU devices (PVC/Xe12 and BMG/Xe20) by adding device detection, platform-specific implementations, and adapting test infrastructure to work with both CUDA and SYCL backends.
Key changes:
- Add Intel Xe12 (PVC) and Xe20 (BMG) support to EVT backend with new node implementations and emitters
- Update device detection to differentiate between BMG (CC=20) and PVC (CC=12) Intel GPUs
- Modify test infrastructure to support XPU tensors and dynamic device selection (CUDA/XPU/CPU)
Reviewed changes
Copilot reviewed 22 out of 22 changed files in this pull request and generated 12 comments.
Show a summary per file
| File | Description |
|---|---|
test/python/cutlass/evt/utils/evt_testbed.py |
Adds dynamic device selection for PyTorch tensors (CUDA/XPU/CPU) |
test/python/cutlass/evt/evt_store_sm80_90.py |
Extends test CC support to include Intel Xe12 (12) and Xe20 (20) |
test/python/cutlass/evt/evt_mixed_sm80_90.py |
Extends test CC support to include Intel Xe12 (12) and Xe20 (20) |
test/python/cutlass/evt/evt_load_sm80_90.py |
Extends test CC support to include Intel Xe12 (12) and Xe20 (20) |
test/python/cutlass/evt/evt_layout_sm80_90.py |
Extends test CC support to include Intel Xe12 (12) and Xe20 (20) |
test/python/cutlass/evt/evt_compute_sm80_90.py |
Extends test CC support to include Intel Xe12 (12) and Xe20 (20) |
python/cutlass_cppgen/backend/utils/device.py |
Adds BMG device detection logic to return CC=20 for Battlemage GPUs |
python/cutlass_cppgen/backend/operation.py |
Replaces CUDA-specific return value with generic success code (0) |
python/cutlass_cppgen/backend/gemm_operation.py |
Updates architecture checks and error handling for XPU compatibility |
python/cutlass_cppgen/backend/frontend.py |
Fixes XPU tensor detection and device pointer handling logic |
python/cutlass_cppgen/backend/evt/passes/pass_manager.py |
Adds Xe-specific function name resolution based on device availability |
python/cutlass_cppgen/backend/evt/passes/pass_get_impl.py |
Adds Xe-prefix support for node implementation lookup |
python/cutlass_cppgen/backend/evt/passes/pass_argument_type.py |
Adds xe12/xe20 argument type handling methods |
python/cutlass_cppgen/backend/evt/frontend/frontend_base.py |
Updates CC checks to include Intel Xe12/Xe20 architectures |
python/cutlass_cppgen/backend/evt/epilogue.py |
Adds Xe emitter selection and updates CC checks |
python/cutlass_cppgen/backend/evt/backend/xe20_nodes.py |
New file implementing Xe20-specific EVT node types |
python/cutlass_cppgen/backend/evt/backend/xe20_emitter.py |
New file implementing Xe20 epilogue emitter |
python/cutlass_cppgen/backend/evt/backend/xe12_nodes.py |
New file implementing Xe12-specific EVT node types |
python/cutlass_cppgen/backend/evt/backend/xe12_emitter.py |
New file implementing Xe12 epilogue emitter |
python/cutlass_cppgen/backend/evt/backend/__init__.py |
Registers new Xe12 and Xe20 emitters and nodes |
python/cutlass_cppgen/backend/compiler.py |
Changes default SYCL target from PVC to BMG |
python/cutlass_cppgen/backend/arguments.py |
Adds comprehensive XPU support with device pointer handling |
Co-authored-by: Copilot <[email protected]>
anamikac-intel
changed the title
Antonyvance
approved these changes
Dec 6, 2025
| @@ -0,0 +1,98 @@ | |||
| ################################################################################################### | |||
| # Copyright (c) 2024 - 2024 Codeplay Software Ltd. All rights reserved. | |||
|
|
||
|
|
||
| class Xe12Emitter: | ||
| def __init__(self, operation: GemmOperationUniversal, graph) -> None: |
There was a problem hiding this comment.
Copilot comment make sense Why cc=90?
| self.nvcc() | ||
| try: | ||
| cc = device_cc() | ||
| if cc in [12, 20]: |
There was a problem hiding this comment.
Can you use the constants here ?https://github.com/intel/sycl-tla/blob/main/python/cutlass_library/arch_constants.py
instead hard coding values. It would be easier to track later.
| @unittest.skipIf(device_cc() not in [12, 20, 80, 86, 89, 90], "This unittest is only supported on CC [12, 20, 80, 86, 89, 90]") | ||
| class TestEVTStore(EVTTestCaseBase): | ||
|
|
||
| @unittest.skipIf(device_cc() != 90, "This test is only for CC 90") |
There was a problem hiding this comment.
This is skipping for anything other than cc=90
| return self._type_decl | ||
|
|
||
| self._type_decl = f""" | ||
| using {self.name_camel} = cutlass::epilogue::fusion::Sm90Compute< |
There was a problem hiding this comment.
Could you please explain this line for fusion:sm90compute. Is that overridden for xe for just calling to provide a default support.
| from typing import Any | ||
|
|
||
| import networkx as nx | ||
| import torch |
There was a problem hiding this comment.
Why are we importing torch in cutlass code? Could you please explain?
| def ensures(self) -> None: | ||
| # Some nodes will be lowered to NoOp, eliminate them | ||
| self.no_op_elimination() | ||
| if torch.xpu.is_available(): |
There was a problem hiding this comment.
Is Torch imported just for this function?
Antonyvance
requested changes
Dec 6, 2025
Antonyvance
left a comment
Antonyvance
left a comment
There was a problem hiding this comment.
Mistakenly approved. Need explanation for the comments.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 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.
This PR depends on #647