Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 13 additions & 0 deletions docs/ubenchx-huawei-bandwidth.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,13 @@
# Huawei copy-bandwidth evidence

The displayed series is **Huawei 950PR**, measured on Ascend950PR_9579 (28 Cube, 56 Vector, 128 GiB). The measured SKU remains in API metadata. This series does not change the global GPU specification catalog or inference-data normalization.

`huawei-bandwidth.json` preserves all 42 full-precision plotted points, source input hashes, per-process medians, observed p10/p90 ranges, pool geometry and sample counts. Each size has three processes × six blocks × twenty observations. Aggregation is the pooled median amortized time, converted by `2 * message bytes / time`; bounds invert p90/p10 latency and are not confidence intervals. No observations are smoothed or supplied for missing sizes. The source is the validated October 7 rotating-pool 950PR experiment; the initial scratch-fill curves and overlap-rejected runs are excluded.

The experiment uses FP16 `copy_` operations in complete rotating-pool NPU graph replays. Each direction has at least 512 MiB and two slots. Every block is checked against source bits before and after timing. The time is amortized over a complete rotation, not an isolated-copy latency. Unlike the GPU series, it is not `triton.testing.do_bench` on one FP32 buffer pair. The UI states this above the chart and labels Huawei tooltip times as amortized. No physical HBM-traffic counters prove that every access misses cache.

The endpoint is about 1,447 GB/s at 16 GiB. MBU uses an assumed nominal reference of 1,600 GB/s; the exact board ceiling was not independently measured.

UI/API coverage: the existing bandwidth/MBU/latency selector is presentation-only. `Huawei 950PR` is available through `test=mem-bw`, `gpu=<key>` and `gpu=all`; `test=all` returns empty sections for tests unavailable on an explicitly selected device. API rows retain optional uncertainty/pool fields, `metadata.protocol`, `dtype`, CANN/torch_npu versions and `missingSizesBytes`. `peakBandwidthSource` states the per-run reference qualification. The dashboard and API share `transformUbenchxRun`.

Hover rulers use the currently selected metric and the same scaled coordinates as the displayed point, including after zoom.
90 changes: 90 additions & 0 deletions packages/app/cypress/e2e/ubenchx.cy.ts
Original file line number Diff line number Diff line change
Expand Up @@ -125,3 +125,93 @@ describe('ubenchX hub and test routes', () => {
cy.get('meta[name="robots"]').should('have.attr', 'content').and('contain', 'noindex');
});
});

const dismiss = {
onBeforeLoad(win: Cypress.AUTWindow) {
win.localStorage.setItem('inferencex-star-modal-dismissed', String(Date.now()));
},
};

describe('Huawei copy-bandwidth evidence', () => {
it('renders Huawei 950PR and qualifies the graph timing in English and Chinese', () => {
cy.visit('/ubenchx/mem-bw', dismiss);
cy.contains('Huawei 950PR').should('be.visible');
cy.contains('910B3').should('not.exist');
cy.contains('Graph time is amortized per copy').should('be.visible');
cy.get('[data-testid="ubenchx-metric-select"]').click();
cy.contains('[role="option"]', 'Bandwidth (TB/s)').click();
cy.contains('h2', 'Bandwidth vs Message Size').should('be.visible');
cy.screenshot('huawei-bandwidth-en', { capture: 'fullPage' });
cy.visit('/zh/ubenchx/mem-bw', dismiss);
cy.contains('图执行时间按拷贝次数摊销').should('be.visible');
cy.contains('910B3').should('not.exist');
cy.screenshot('huawei-bandwidth-zh', { capture: 'fullPage' });
});
it('serves measured endpoints and empty unavailable tests through the public API', () => {
cy.request('/api/v1/views/ubenchx?test=all&gpu=Huawei%20950PR').then(({ body }) => {
expect(body.memBw.gpus[0].rows.at(-1).bytes).to.equal(17179869184);
expect(body.memBw.gpus[0].rows.at(-1).mbuPercent).to.be.closeTo(90.4545974, 0.0001);
expect(body.smL2Distance.gpus).to.deep.equal([]);
expect(body.tpcSkyline.gpus).to.deep.equal([]);
});
});
});

function assertBandwidthRuler(chart: string) {
cy.get(`${chart} circle.point`)
.filter((_, element) => {
const datum = (
element as unknown as SVGCircleElement & { __data__: { gpuKey: string; bytes: number } }
).__data__;
return datum.gpuKey === 'H200 SXM' && datum.bytes === 33554432;
})
.should('have.length', 1)
.then(($point) => {
const x = Number($point.attr('cx')),
y = Number($point.attr('cy'));
cy.wrap($point).trigger('mouseenter', { force: true });
cy.get(`${chart} .horizontal-ruler`).should(($line) => {
expect(Number($line.attr('y1'))).to.be.closeTo(y, 0.01);
expect(Number($line.attr('y2'))).to.be.closeTo(y, 0.01);
});
cy.get(`${chart} .vertical-ruler`).should(($line) => {
expect(Number($line.attr('x1'))).to.be.closeTo(x, 0.01);
expect(Number($line.attr('x2'))).to.be.closeTo(x, 0.01);
});
});
}

it('keeps bandwidth hover guides on the rendered point in every metric and after zoom', () => {
cy.visit('/ubenchx/mem-bw', dismiss);
assertBandwidthRuler('#ubenchx-mbu');
for (const [label, chart] of [
['Bandwidth (TB/s)', '#ubenchx-bandwidth'],
['Latency (ms)', '#ubenchx-latency'],
]) {
cy.get('[data-testid="ubenchx-metric-select"]').click();
cy.contains('[role="option"]', label).click();
assertBandwidthRuler(chart);
}
cy.get('#ubenchx-latency circle.point')
.first()
.then(($point) => {
const previousX = $point.attr('cx');
cy.get('#ubenchx-latency svg').then(($svg) => {
const rect = $svg[0].getBoundingClientRect();
cy.wrap($svg).trigger('wheel', {
eventConstructor: 'WheelEvent',
deltaY: -200,
shiftKey: true,
clientX: rect.left + 150,
clientY: rect.top + 100,
force: true,
});
});
cy.get('#ubenchx-latency circle.point')
.first()
.should(($current) => {
expect($current.attr('cx')).not.to.equal(previousX);
});
});
assertBandwidthRuler('#ubenchx-latency');
});
44 changes: 44 additions & 0 deletions packages/app/src/app/api/v1/views/ubenchx/route.test.ts
Original file line number Diff line number Diff line change
Expand Up @@ -158,3 +158,47 @@ describe('ubenchx view', () => {
expect(h200Last.mbuPercent).toBeCloseTo((4305.39 / 4800) * 100, 4);
});
});

it('exposes Huawei rotating-copy evidence without inventing missing tests or sizes', async () => {
const response = await GET(request({ test: 'all', gpu: 'huawei 950pr' }));
expect(response.status).toBe(200);
const body = await response.json();
const run = body.memBw.gpus[0];
expect(run.gpu).toBe('Huawei 950PR');
expect(run.peakBandwidthGbps).toBe(1600);
expect(run.peakBandwidthSource).toContain('not independently verified');
expect(run.metadata).toMatchObject({
protocol: 'rotating-pool-npu-graph',
dtype: 'float16',
missingSizesBytes: [],
});
expect(run.rows.some((row: { bytes: number }) => row.bytes === 17179869184)).toBe(true);
const endpoint = run.rows.find((row: { bytes: number }) => row.bytes === 17179869184);
expect(endpoint.bandwidthGbps).toBeCloseTo(1447.2735584561165, 6);
expect(endpoint.mbuPercent).toBeCloseTo(90.45459740350728, 6);
expect(endpoint.observations).toBe(360);
expect(body.smL2Distance.gpus).toEqual([]);
expect(body.tpcSkyline.gpus).toEqual([]);
});

it('preserves full-precision Huawei bandwidth arithmetic and API/UI normalization', async () => {
for (const gpu of ['Huawei 950PR']) {
const response = await GET(request({ test: 'mem-bw', gpu }));
const body = await response.json();
const run = body.gpus[0];
expect(run).toEqual(transformUbenchxRun(gpu, UBENCHX_RUNS[gpu]));
for (const row of run.rows) {
expect(row.bandwidthGbps).toBeCloseTo((2 * row.bytes) / row.timeMs / 1e6, 6);
expect(row.bandwidthP10Gbps).toBeLessThanOrEqual(row.bandwidthGbps);
expect(row.bandwidthP90Gbps).toBeGreaterThanOrEqual(row.bandwidthGbps);
}
}
});

it('does not expose the removed 910B3 bandwidth series', async () => {
const all = await GET(request({ test: 'mem-bw', gpu: 'all' }));
const body = await all.json();
expect(body.gpus.some((run: { gpu: string }) => run.gpu === 'Ascend 910B3')).toBe(false);
const removed = await GET(request({ test: 'mem-bw', gpu: 'Ascend 910B3' }));
expect(removed.status).toBe(400);
});
19 changes: 9 additions & 10 deletions packages/app/src/app/api/v1/views/ubenchx/route.ts
Original file line number Diff line number Diff line change
Expand Up @@ -64,13 +64,11 @@ export function GET(request: NextRequest) {
}

// test === 'all': return every test
const smL2GpuKeys = Object.keys(SM_L2_RUNS);
const smL2Gpu = parseEnumParam(s.get('gpu'), 'gpu', ['all', ...smL2GpuKeys], 'all');
const smL2Results =
smL2Gpu === 'all'
rawGpu === null || gpu === 'all'
? Object.entries(SM_L2_RUNS).map(([key, run]) => transformSmL2Run(key, run))
: SM_L2_RUNS[smL2Gpu]
? [transformSmL2Run(smL2Gpu, SM_L2_RUNS[smL2Gpu])]
: SM_L2_RUNS[gpu]
? [transformSmL2Run(gpu, SM_L2_RUNS[gpu])]
: [];

return Promise.resolve(
Expand All @@ -85,11 +83,12 @@ export function GET(request: NextRequest) {
// Like smL2Distance, default to every GPU rather than mem-bw's first one, and
// match the GPU key case-insensitively as parseEnumParam does for the others.
tpcSkyline: {
gpus: tpcSkylineResults(
Object.keys(TPC_SKYLINE_RUNS).find(
(key) => key.toLowerCase() === rawGpu?.toLowerCase(),
) ?? 'all',
),
gpus:
rawGpu === null || gpu === 'all'
? tpcSkylineResults('all')
: TPC_SKYLINE_RUNS[gpu]
? tpcSkylineResults(gpu)
: [],
},
}),
);
Expand Down
57 changes: 41 additions & 16 deletions packages/app/src/components/ubenchx/UbenchxContent.tsx
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,12 @@ import { useLocale } from '@/lib/i18n/use-locale';
import type { Locale } from '@/lib/i18n/i18n';

import { UBENCHX_GPU_KEYS, UBENCHX_RUNS } from './ubenchx-data';
import { transformUbenchxRun, formatBytes, type UbenchxDerivedRow } from './ubenchx-transform';
import {
transformUbenchxRun,
formatBytes,
sizeAxisTicks,
type UbenchxDerivedRow,
} from './ubenchx-transform';

const GPU_COLORS: Record<string, string> = {
'B300 SXM': 'oklch(0.62 0.15 128)',
Expand All @@ -27,13 +32,16 @@ const GPU_COLORS: Record<string, string> = {
MI355X: 'oklch(0.62 0.20 18)',
MI325X: 'oklch(0.55 0.20 28)',
MI300X: 'oklch(0.48 0.20 38)',
'Huawei 950PR': '#C7000B',
};

const STRINGS = {
en: {
pageTitle: 'HBM Bandwidth',
pageSubtitle:
'Microbenchmark measuring device-memory copy bandwidth across message sizes on NVIDIA and AMD GPUs.',
'Microbenchmark measuring device-memory copy bandwidth across message sizes on NVIDIA, AMD and Huawei accelerators.',
protocolNote:
'Huawei uses amortized rotating-buffer graph timing; NVIDIA/AMD use Triton copy timing. Small-message latency is not directly comparable.',
latencyTitle: 'Latency vs Message Size',
latencyY: 'Latency (ms)',
mbuTitle: 'Memory Bandwidth Utilization (MBU) vs Message Size',
Expand All @@ -42,10 +50,16 @@ const STRINGS = {
bwY: 'Bandwidth (TB/s)',
xAxis: 'Message Size',
yMetric: 'Y-axis',
peakNote: "MBU is relative to each GPU's own peak HBM bandwidth from GPU_SPECS.",
peakNote:
'MBU uses each device’s bandwidth reference. Huawei’s 1.6 TB/s references are qualified below.',
methodology: 'Methodology',
methodologyText:
'Times b.copy_(a) on float32 tensors using triton.testing.do_bench; bandwidth = 2 * bytes / time (read + write). Power-of-two sizes from 8 B to 16 GiB.',
'NVIDIA/AMD: times b.copy_(a) on float32 tensors using triton.testing.do_bench; bandwidth = 2 * bytes / time (read + write). Power-of-two sizes from 8 B to 16 GiB.',
huaweiMethodology:
'Huawei: FP16 rotating source/destination pools of at least 512 MiB each, timed as full rotating-buffer NPU graph replays. Three processes, 360 observations per size; full copied-bit checks before/after each block. Graph time is amortized per copy, not isolated-copy latency. Protocols differ from the GPU curves. No physical HBM-traffic counters were collected.',
huaweiReference:
'Huawei MBU uses a 1,600 GB/s reference; the exact board bandwidth ceiling was not independently measured.',
amortized: 'Amortized graph copy time',
source: 'Source',
instructions:
'Shift+Scroll to zoom · Drag to pan · Double-click to reset · Click a point to pin tooltip',
Expand All @@ -54,12 +68,14 @@ const STRINGS = {
tooltipLatency: 'Latency',
tooltipBandwidth: 'Bandwidth',
tooltipMbu: 'MBU',
tooltipPeak: 'Peak BW',
tooltipPeak: 'Bandwidth reference',
aria: 'ubenchX device-memory copy bandwidth charts',
},
zh: {
pageTitle: 'HBM 带宽',
pageSubtitle: '在 NVIDIA 和 AMD GPU 上测量不同消息大小的显存拷贝带宽微基准测试。',
pageSubtitle: '测量 NVIDIA、AMD 和 Huawei 加速器在不同消息大小下的显存拷贝带宽。',
protocolNote:
'Huawei 使用轮转缓冲区的图执行摊销计时;NVIDIA/AMD 使用 Triton 拷贝计时。小消息延迟不能直接比较。',
latencyTitle: '延迟 vs 消息大小',
latencyY: '延迟(ms)',
mbuTitle: '显存带宽利用率(MBU)vs 消息大小',
Expand All @@ -68,18 +84,22 @@ const STRINGS = {
bwY: '带宽(TB/s)',
xAxis: '消息大小',
yMetric: 'Y 轴',
peakNote: 'MBU 基于各 GPU 自身在 GPU_SPECS 中的峰值 HBM 带宽计算。',
peakNote: 'MBU 按各设备的带宽参考值计算。Huawei 的 1.6 TB/s 参考值及其限制见下文。',
methodology: '测试方法',
methodologyText:
'使用 triton.testing.do_bench 计时 b.copy_(a)(float32 张量);带宽 = 2 * bytes / time(读 + 写)。消息大小为 2 的幂次,从 8 B 到 16 GiB。',
'NVIDIA/AMD:使用 triton.testing.do_bench 计时 b.copy_(a)(float32 张量);带宽 = 2 * bytes / time(读 + 写)。消息大小为 2 的幂次,从 8 B 到 16 GiB。',
huaweiMethodology:
'Huawei:源和目标分别使用至少 512 MiB 的 FP16 轮转缓冲池,每次计时覆盖完整缓冲池轮转的一次 NPU 图重放。每个大小在三个进程中共测量 360 次,每个测量块前后逐位校验全部拷贝结果。图执行时间按拷贝次数摊销,不代表单次独立拷贝延迟;与 GPU 曲线的测试协议不同,且未采集物理 HBM 流量计数器。',
huaweiReference: 'Huawei 的 MBU 使用 1,600 GB/s 参考值,该板卡的确切带宽上限未经独立实测。',
amortized: '摊销后的图拷贝时间',
source: '数据来源',
instructions: 'Shift+滚轮缩放 · 拖动平移 · 双击重置 · 点击数据点固定提示框',
dismiss: '点击其他区域关闭',
tooltipSize: '大小',
tooltipLatency: '延迟',
tooltipBandwidth: '带宽',
tooltipMbu: 'MBU',
tooltipPeak: '峰值带宽',
tooltipPeak: '带宽参考值',
aria: 'ubenchX 显存拷贝带宽图表',
},
} as const;
Expand Down Expand Up @@ -143,7 +163,7 @@ function formatTooltip(point: ChartPoint, isPinned: boolean, t: (typeof STRINGS)
<span style="display: inline-block; width: 8px; height: 8px; border-radius: 9999px; background: ${color};"></span>${point.gpuKey}
</div>
${tooltipLine(t.tooltipSize, point.sizeLabel)}
${tooltipLine(t.tooltipLatency, `${point.timeMs.toFixed(4)} ms`)}
${tooltipLine(UBENCHX_RUNS[point.gpuKey].metadata.protocol ? t.amortized : t.tooltipLatency, `${point.timeMs.toFixed(4)} ms`)}
${tooltipLine(t.tooltipBandwidth, `${(point.bandwidthGbps / 1000).toFixed(3)} TB/s`)}
${tooltipLine(t.tooltipMbu, `${point.mbuPercent.toFixed(1)}%`)}
${tooltipLine(t.tooltipPeak, `${point.peakBandwidthGbps / 1000} TB/s`)}
Expand Down Expand Up @@ -215,10 +235,7 @@ function UbenchxChart({
return ticks;
}, [yScaleType, yDomain]);

const xTickValues = useMemo(() => {
const all = [...new Set(allPoints.map((p) => p.x))].sort((a, b) => a - b);
return all.filter((_, i) => i % 3 === 0 || i === all.length - 1);
}, [allPoints]);
const xTickValues = useMemo(() => sizeAxisTicks(...xDomain), [xDomain]);

return (
<D3Chart<ChartPoint>
Expand Down Expand Up @@ -284,6 +301,8 @@ function UbenchxChart({
}}
tooltip={{
rulerType: 'crosshair',
getRulerX: (point, scale) => (scale as d3.ScaleLogarithmic<number, number>)(point.x),
getRulerY: (point, scale) => scale(point.y),
attachToLayer: 1,
content: (point, isPinned) => formatTooltip(point, isPinned, STRINGS[locale]),
}}
Expand Down Expand Up @@ -324,6 +343,7 @@ export function UbenchxContent() {
<div>
<h2 className="text-xl font-semibold tracking-tight">{t.pageTitle}</h2>
<p className="text-sm text-muted-foreground mt-1">{t.pageSubtitle}</p>
<p className="text-sm text-muted-foreground mt-2">{t.protocolNote}</p>
</div>

<div className="flex flex-col gap-1.5 sm:flex-row sm:items-center sm:gap-3">
Expand Down Expand Up @@ -376,11 +396,16 @@ export function UbenchxContent() {
<section className="text-xs text-muted-foreground space-y-1">
<h3 className="text-sm font-semibold">{t.methodology}</h3>
<p>{t.methodologyText}</p>
<p>{t.huaweiMethodology}</p>
<p>{t.huaweiReference}</p>
{viewResults.map((r) => (
<p key={r.gpu}>
<strong>{r.gpu}:</strong> {r.metadata.gpu} | Driver: {r.metadata.driver} | PyTorch:{' '}
{r.metadata.torch} | Triton: {r.metadata.triton} | Container: {r.metadata.container} |
Peak: {r.peakBandwidthGbps / 1000} TB/s (GPU_SPECS)
{r.metadata.torch} |{' '}
{r.metadata.cann
? `CANN: ${r.metadata.cann} | torch_npu: ${r.metadata.torchNpu}`
: `Triton: ${r.metadata.triton} | Container: ${r.metadata.container}`}{' '}
|{r.peakBandwidthGbps / 1000} TB/s
</p>
))}
{sourceUrls.map((url) => (
Expand Down
Loading
Loading