メニュー
Review pull requests for XPU operator or backend code. Use when reviewing PRs, when asked to review code changes, or when the user mentions "review PR", "code review", or "check this PR".
Codex または Claude でインストール この Prompt をコピーして Codex、Claude、または他のアシスタントに貼り付けると、Skill ページを確認してインストールできます。
SOC 職業分類に基づく
Analyze UT (unit test) results for a torch-xpu-ops PR. Use when asked to check test results, analyze CI failures, or evaluate test coverage for a PR. Produces a structured report of new failures, failure relevance, and new test coverage.
Rules for automatically determining which disable_* labels to apply to a PR based on the file paths changed. Used by the auto-label workflow.
Extract GPU ISA from oneDNN ngen-JIT kernels. This is the ONLY codegen path that bypasses the standard SYCL/SPIR-V/IGC stack. oneDNN uses its own native code generator (ngen) that directly emits GPU ISA bytes — no SPIR-V, no IGC, no zebin ELF, no .debug_line. Use when extracting ASM from oneDNN kernels (gemm_kernel, gen_conv_kernel), matmul, linear, conv, or SDPA-graph ops dispatched via mkldnn.
Extract GPU ISA from AOT-compiled SYCL binaries using NEO driver runtime dump. Use when the kernel is a SYCL symbol (_ZTS...) from an AOT binary such as libtorch_xpu.so, SYCL-TLA FMHA, or any DPC++ binary compiled with -fsycl-targets=spir64_gen -Xs "-device <gpu>".
Extract GPU ISA from JIT-compiled SYCL kernels. The binary only contains SPIR-V; IGC compiles it to native ISA at first launch. The zebin only exists in memory at runtime, so IGC_ShaderDumpEnable=1 is required to capture it. Use when extracting ASM from SYCL JIT kernels, torch-xpu-ops built with TORCH_XPU_ARCH_LIST=none, or standalone DPC++ without AOT flags.
Extract GPU ISA from Triton kernels on XPU. Triton compiles through the same IGC backend as SYCL (Triton IR → SPIR-V → IGC → zebin). Extraction is identical to sycl-jit: IGC_ShaderDumpEnable=1 captures the zebin at runtime. Use when extracting ASM from torch.compile fusions (triton_per_fused, triton_poi, triton_red), standalone @triton.jit kernels, or Inductor-generated XPU kernels.
| name | pr-review |
| description | Review pull requests for XPU operator or backend code. Use when reviewing PRs, when asked to review code changes, or when the user mentions "review PR", "code review", or "check this PR". |
Review torch-xpu-ops pull requests focusing on what CI cannot check: correctness against CPU/CUDA semantics, XPU-specific risks (synchronization, indexing, precision), test adequacy, and backward compatibility.
Detailed reference:
If the user invokes /pr-review with no arguments, do not perform a review. Instead, ask:
What would you like me to review?
- A PR number or URL (e.g.,
/pr-review 12345)- A local branch (e.g.,
/pr-review branch)
The user provides a PR number or URL:
/pr-review 12345
/pr-review https://github.com/intel/torch-xpu-ops/pull/12345
For a detailed review with line-by-line specific comments:
/pr-review 12345 detailed
Use gh CLI to fetch PR data:
# Get PR details
gh pr view <PR_NUMBER> --json title,body,author,baseRefName,headRefName,files,additions,deletions,commits
# Get the diff
gh pr diff <PR_NUMBER>
# Get PR comments
gh pr view <PR_NUMBER> --json comments,reviews
Review changes in the current branch that are not in main:
/pr-review branch
/pr-review branch detailed
Use git commands to get branch changes:
# Get current branch name
git branch --show-current
# Get list of changed files compared to main
git diff --name-only main...HEAD
# Get full diff compared to main
git diff main...HEAD
# Get commit log for the branch
git log main..HEAD --oneline
# Get diff stats (files changed, insertions, deletions)
git diff --stat main...HEAD
For local branch reviews:
When invoked via @copilot /pr-review or @claude /pr-review on a GitHub PR, detect this mode by the presence of <formatted_context>, <pr_or_issue_body>, and <comments> tags in the prompt.
The prompt already contains:
Use git commands to get the diff and commit history. The base branch name is in the
prompt context (look for PR Branch: <head> -> <base> or the baseBranch field).
# Get the full diff against the base branch
git diff origin/<baseBranch>...HEAD
# Get diff stats
git diff --stat origin/<baseBranch>...HEAD
# Get commit history for this PR
git log origin/<baseBranch>..HEAD --oneline
# If the base branch ref is not available, fetch it first
git fetch origin <baseBranch> --depth=1
Do NOT use gh CLI commands in this mode -- only git commands are available.
All PR metadata, comments, and reviews are already in the prompt context;
only the diff and commit log need to be fetched via git.
pytorch/pytorch.The review checklist is large. Spawn sub-agents to investigate whether checklist items apply: read surrounding code, check upstream PyTorch implementation for parity, or verify tests exist. Spawn them in parallel for independent areas.
Before reviewing, build understanding of what the PR touches and why:
For every changed kernel or operator file, fetch and read the corresponding upstream PyTorch implementation BEFORE evaluating correctness:
src/ATen/native/xpu/<Op>.cpp → read aten/src/ATen/native/<Op>.cppsrc/ATen/native/xpu/sycl/<Op>Kernels.cpp → read aten/src/ATen/native/cuda/<Op>.cuMathExtensions.h) → read aten/src/ATen/native/Math.hUse gh api or spawn a sub-agent to fetch the upstream file content. Do NOT proceed to the deep review until upstream code has been read. Quote or summarize relevant upstream patterns in your working notes before continuing.
Go through every changed line in the diff and evaluate against the review checklist in review-checklist.md.
If the diff adds or modifies any agent instruction files (SKILL.md, AGENTS.md, claude.md, copilot-instructions.md, or files under .claude/ / .github/), load the skill-writer skill and evaluate those changes against its guidelines. Do NOT skip this — treat it as a blocking gate for those files.
Pay special attention to XPU-specific risks:
Evaluate BC implications per bc-guidelines.md. For non-trivial BC questions, spawn a sub-agent to search for existing callers of the modified API.
Structure your review with actionable feedback organized by category. Every finding should be traceable to a specific line in the diff.
After drafting the review, spawn a sub-agent per reported issue (in parallel) to independently verify the claim by re-reading the relevant code. Drop invalid issues, reword uncertain ones with a note about confidence level.
Omit sections where you have no problems to report. Every sentence must identify a problem or request a change.
## PR Review: #<number>
<!-- Or for local branch reviews: -->
## Branch Review: <branch-name> (vs main)
### Summary
What the PR does (1 sentence), then the overall verdict.
### Correctness
[Problems only — semantic parity, edge cases, behavioral issues]
### XPU-Specific Risks
[Problems only — synchronization, indexing, precision, kernel issues]
### Dispatch & Registration
[Problems only — yaml wiring, fallback, backend path]
### Testing
[Problems only — missing tests, wrong patterns, inadequate coverage]
### Backward Compatibility
[Problems only]
### Performance
[Problems only]
### Recommendation
**Approve** / **Request Changes** / **Needs Discussion**
Missing tests (new functionality without tests, bug fixes without regression tests) always means **Request Changes**.
[Brief justification — focus on what blocks approval. IMPORTANT: Do NOT use `#N` (e.g., #1, #2, #3) to reference findings — GitHub auto-links these to real issues/PRs. Instead use descriptive references like "the step numbering issue", "the stale path in auto-labeling", or inline the file path.]
Only include this section if the user requests a "detailed" or "in depth" review.
When performing a detailed review, group findings by severity:
For each finding, quote the offending line and provide a concrete fix.
### 🔴 Must Fix (N issues)
**[Category] file.cpp:42** — <description>
<quoted code>
→ <suggested fix>
### 🟡 Should Fix (N issues)
...
### 🟢 Suggestions (N issues)
...
### ✅ What looks good
<briefly note well-written parts — good reviews are balanced>
If there are zero issues in a severity level, omit that section. Always include the "What looks good" section in detailed reviews.
Principle: For SYCL programming, always use SYCL programming model terms. Hardware architecture terms should only appear in comments explaining the motivation for a particular optimization.
This is a 🔴 Must Fix category.
Use these terms in all code, variable names, and function names:
| ❌ Deprecated Term | ✅ Current Term | Notes |
|---|---|---|
| SIMD width / SIMD length / simd_width | subgroup size | SYCL/oneAPI standard term |
| SIMD lane | work-item (within subgroup) | Aligns with SYCL spec |
| SIMD-8 / SIMD-16 / SIMD-32 | subgroup size 8 / 16 / 32 | Use numeric subgroup size |
| thread block / threadblock | work-group | SYCL/oneAPI standard term |
Use these only in comments to explain why a particular optimization choice was made (e.g., occupancy, register pressure, memory alignment). They must NOT appear in variable names, function names, or general code.
| ❌ Deprecated Term | ✅ Current Intel Term | Generic Term | Abbreviation |
|---|---|---|---|
| Execution Unit (EU) | Xe Vector Engine | Vector Engine | XVE |
| Systolic / "DPAS part of EU" | Xe Matrix eXtension | Matrix Engine | XMX |
| Subslice (SS) / Dual Subslice (DSS) | Xe-core | — | XC |
| HW thread | XVE thread | — | Each XVE thread executes a subgroup (subgroup size 16 or 32) |
| SIMD-16 / SIMD-32 (hardware context) | XVE thread width | — | The number of data elements processed per thread; maps to subgroup size |
| GRF file / GRF count | register file / register count | — | Use architecture-neutral where possible |
| SLM (ambiguous) | SLM (Shared Local Memory) | — | Spell out on first use |
subslice_count → xc_count or xecore_countgetSimdWidth() → getSubgroupSize()// 🔴 Bad
int num_subslices = device.get_info<ext::info::device::gpu_subslices_per_slice>();
int simd_len = 16;
// Each EU has 8 HW threads
// ✅ Good
int num_xecores = device.get_info<ext::info::device::gpu_subslices_per_slice>();
// Note: API name still uses "subslices" — wrap in a helper if possible
int subgroup_size = 16;
// Each XVE supports 8 concurrent threads
Edge case — API boundaries: When the underlying API (Level Zero, OpenCL, SYCL extensions) still uses old terms in function/enum names, it's acceptable to use them at the call site only. Wrap in a helper with modern naming, and add a comment:
// ✅ OK — API uses legacy name, but our abstraction uses modern term
// Level Zero API still exposes "subslice" in its info query
int xecore_count = zeDeviceGetSubsliceCount(device); // legacy API name
Principle: Never use sycl::shift_group_left combined with a reduction operation to perform subgroup reductions. Always use sycl::reduce_over_group instead.
This is a 🔴 Must Fix category.
Flag any code where sycl::shift_group_left is combined with a reduction combiner (any binary operation that accumulates values: add, min, max, mean, product, bitwise-or/and, etc.). The true signal is the shift + combine combination, not the loop itself — the loop may be at a different call site or abstracted behind a helper.
sycl::shift_group_left produces excessive integer ALU instructions for index manipulation, which stalls the ALU-INT pipeline and degrades performance. The SYCL runtime's built-in reduce_over_group can use hardware-optimized reduction paths.
// 🔴 Bad — shift_group_left generates massive int-related instructions,
// causing ALU-INT pipe stall
for (int offset = 1; offset < sg_size; offset <<= 1) {
arg_t other = sycl::shift_group_left(sg, value[i], offset);
value[i] = combine(value[i], other); // combine = any reduction op
}
Replace with sycl::reduce_over_group using the appropriate SYCL binary operation:
// ✅ Good — uses hardware-optimized reduction
value[i] = sycl::reduce_over_group(sg, value[i], sycl::plus<arg_t>()); // for sum/mean
value[i] = sycl::reduce_over_group(sg, value[i], sycl::minimum<arg_t>()); // for min
value[i] = sycl::reduce_over_group(sg, value[i], sycl::maximum<arg_t>()); // for max
All of these are the same anti-pattern:
// 🔴 Bad — inline loop (any direction)
for (int offset = 1; offset < sg_size; offset <<= 1) { ... shift_group_left ... }
for (int offset = (sg_size >> 1); offset > 0; offset >>= 1) { ... shift_group_left ... }
// 🔴 Bad — shift+combine split into a functor (loop lives elsewhere)
struct ShiftAndCombine {
arg_t operator()(sycl::sub_group sg, arg_t value, int offset) const {
arg_t other = sycl::shift_group_left(sg, value, offset);
return combine_(value, other);
}
BinaryOp combine_;
};
// 🔴 Bad — with vectorized inner loop
for (int offset = 1; offset < sg_size; offset <<= 1) {
for (int i = 0; i < out_vec_sz; ++i) {
arg_t other = sycl::shift_group_left(sg, value[i], offset);
value[i] = combine(value[i], other);
}
}
src/ATen/native/xpu/sycl/ — All SYCL kernel filessycl::shift_group_left in combination with a reduction combinerWhen reviewing, consult these for context:
src/ATen/native/xpu/ — XPU operator implementationssrc/ATen/native/xpu/sycl/ — SYCL kernel implementationssrc/ATen/native/xpu/XPUFallback.template — Fallback logictest/xpu/ — XPU-specific teststest/test_ops_xpu.py — OpInfo-based XPU operator testsIf a calling workflow explicitly requires a skill marker, append this exact literal final line: Custom skills applied: pr-review.
Otherwise, keep the reply in the requested review format and do not force an extra trailing sentence.