https://www.scnet.cn/ui/console/index.html#/notebook/add?resourceGroupName=hx1hgbwnormal&clusterId=20091  拿到这台海光 BW 实例后，我没有马上安装框架或下载模型，而是先确认三件事：设备是否真的被计算框架识别、容器实际能用多少 CPU/内存，以及现有软件栈能不能完成前向和反向计算。 这一步很重要。平台规格、容器里 `lscpu`/`free` 看到的宿主机资源，以及作业受 cgroup 限制后真正可用的资源，并不一定是同一回事。如果一开始把资源边界判断错了，后面出现 DataLoader 卡顿、编译进程被杀或训练 OOM，很容易误判成加速卡或框架问题。 本文记录一次实际排查和验证过程。所有结果来自同一台 SCNet Notebook 实例，截图做过脱敏；文中的吞吐数据只用于验证计算链路是否正常，不代表产品峰值性能。 ## 一、先说结论 这套环境可以胜任以下工作： - PyTorch 基础教学和深度学习课程作业； - Transformer、RoPE、RMSNorm、SwiGLU、Attention 等算子的实现与调试； - BF16/FP16 混合精度实验； - Triton 自定义算子练习； - FlashAttention 前向、反向验证； - 中小模型推理，以及在显存允许范围内进行 LoRA/QLoRA 微调； - 使用 vLLM 做单卡推理服务实验。 本次实例的核心环境如下。 | 项目 | 实测结果 | |---|---| | 操作系统 | Ubuntu 22.04.5 LTS | | 加速卡 | 海光 BW，1 卡 | | 显存 | 63.98 GiB | | DTK | `/opt/dtk-26.04-DCC2602-0317` | | Python | 3.10.12 | | PyTorch | 2.10.0 | | HIP 运行时 | 6.3.26093 | | Triton | 3.5.0 | | Triton 目标 | `hip / gfx936 / warp_size=64` | | FlashAttention | 2.8.3 定制版本 | | vLLM | 0.18.1 | | BF16 | 支持 |  ## 二、连接实例：浏览器适合试用，VS Code 更适合持续开发 SCNet 页面提供 JupyterLab、工具面板和 SSH 登录信息。临时运行几条命令时，浏览器终端已经够用；如果要反复修改训练脚本、查看 Git diff 或调试 Python，VS Code Remote - SSH 会顺手很多。 本地 `~/.ssh/config` 可以写成下面这样，主机和端口用控制台显示的值替换。密码不要写入配置文件，也不要提交到仓库。 ```sshconfig Host scnet-hygon-bw HostName User root Port ServerAliveInterval 30 ServerAliveCountMax 3 ``` 如果远端不能直接访问 VS Code 的下载站点，可以在本地 VS Code 设置中加入： ```json { "remote.SSH.localServerDownload": "always", "remote.SSH.useExecServer": false } ``` 这样会先在本机下载 VS Code Server，再通过 SSH 传到远端，能绕开远端网络访问受限的问题。连接成功后，我习惯先打开一个单独的实验目录，而不是直接在 `/root` 下堆脚本： ```bash mkdir -p /root/hygon-bw-guide-demo cd /root/hygon-bw-guide-demo ``` ## 三、不要只看 `lscpu` 和 `free` 先检查系统和 cgroup 限额： ```bash cat /etc/os-release lscpu free -h # cgroup v2 cat /sys/fs/cgroup/cpu.max cat /sys/fs/cgroup/memory.max ``` 这台容器里能看到宿主机约 128 个 CPU 线程和 503 GiB 内存，但作业实际受到约 15 个 CPU 核、59 GiB 内存的限制。因此，`num_workers` 不应照着 128 个线程设置，编译并行度也不宜直接拉满。 我会再用下面的方式把限制换算成人能直接读懂的数字： ```bash python - <<'PY' from pathlib import Path quota, period = Path("/sys/fs/cgroup/cpu.max").read_text().split() memory = int(Path("/sys/fs/cgroup/memory.max").read_text()) cpu_limit = "unlimited" if quota == "max" else int(quota) / int(period) print(f"CPU limit: {cpu_limit}") print(f"Memory limit: {memory / 1024**3:.2f} GiB") PY ``` 这个差异会直接影响三类参数： - DataLoader 的 `num_workers`； - 源码编译时的并行任务数； - CPU offload、数据缓存和内存映射文件的规模。 ## 四、DTK 环境必须在当前 Shell 中激活 这台实例已经预装 DTK，但新开的 Shell 不一定自动带上所需环境变量。先执行： ```bash source /opt/dtk/env.sh ``` 然后检查 PyTorch 是否能看到设备： ```bash python - <<'PY' import torch print("torch:", torch.__version__) print("hip:", torch.version.hip) print("available:", torch.cuda.is_available()) print("device_count:", torch.cuda.device_count()) print("device:", torch.cuda.get_device_name(0)) print("bf16:", torch.cuda.is_bf16_supported()) PY ``` 这里有一个容易困惑的地方：在 DTK/HIP 软件栈中，PyTorch 仍然沿用 `torch.cuda` 这套接口名。代码里出现 `device="cuda"` 并不表示程序偷偷切到了 NVIDIA CUDA，它只是 PyTorch 兼容接口的历史命名；实际后端可以从 `torch.version.hip` 和设备名称确认。 设备状态可以用 `hy-smi` 查看： ```bash hy-smi ```  如果输出开头出现 `rocm_smi_lib` 回退警告，但后续仍能正常列出 HCU、温度、功耗和显存占用，我会先把它当作管理库路径提示，而不是立刻认定设备不可用。最终还要以 PyTorch 分配张量和实际计算是否成功为准。 ## 五、一份可复现的四段式验证脚本 只打印 `torch.cuda.is_available() == True` 还不够。我更关心四条链路： 1. FP16/BF16 矩阵乘是否能执行，并且结果有限； 2. PyTorch 原生 SDPA 能否完成 BF16 因果注意力的前向和反向； 3. Triton 能否针对 `gfx936` 编译并运行自定义 kernel； 4. FlashAttention 能否完成 BF16 因果注意力的前向和反向。 下面的脚本会把完整结果保存成 JSON，方便后续比较镜像或依赖版本。 ```python #!/usr/bin/env python3 from __future__ import annotations import json import platform import time from pathlib import Path import torch import torch.nn.functional as F RESULTS_PATH = Path(__file__).with_name("validation_results.json") def gib(num_bytes: int) -> float: return num_bytes / 1024**3 def timed_matmul( dtype: torch.dtype, size: int = 2048, repeats: int = 8, ) -> dict: torch.manual_seed(336) lhs = torch.randn((size, size), device="cuda", dtype=dtype) rhs = torch.randn((size, size), device="cuda", dtype=dtype) for _ in range(2): _ = lhs @ rhs torch.cuda.synchronize() started = time.perf_counter() for _ in range(repeats): result = lhs @ rhs torch.cuda.synchronize() elapsed = (time.perf_counter() - started) / repeats tflops = 2 * size**3 / elapsed / 1e12 assert torch.isfinite(result).all() return { "dtype": str(dtype).removeprefix("torch."), "shape": f"{size}x{size}", "average_ms": round(elapsed * 1000, 3), "tflops": round(tflops, 3), "finite": True, } def check_sdpa() -> dict: torch.manual_seed(336) q = torch.randn( (2, 8, 128, 64), device="cuda", dtype=torch.bfloat16, requires_grad=True, ) k = torch.randn_like(q, requires_grad=True) v = torch.randn_like(q, requires_grad=True) output = F.scaled_dot_product_attention( q, k, v, is_causal=True, ) output.float().square().mean().backward() tensors = (output, q.grad, k.grad, v.grad) assert all(bool(torch.isfinite(item).all()) for item in tensors) return { "shape": list(output.shape), "dtype": "bfloat16", "forward_finite": True, "backward_finite": True, } def check_triton() -> dict: import triton import triton.language as tl @triton.jit def add_kernel( x_ptr, y_ptr, out_ptr, n_elements: tl.constexpr, BLOCK: tl.constexpr, ): offsets = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK) mask = offsets < n_elements x = tl.load(x_ptr + offsets, mask=mask) y = tl.load(y_ptr + offsets, mask=mask) tl.store(out_ptr + offsets, x + y, mask=mask) torch.manual_seed(336) n_elements = 65_537 x = torch.randn(n_elements, device="cuda", dtype=torch.float32) y = torch.randn_like(x) output = torch.empty_like(x) grid = (triton.cdiv(n_elements, 256),) add_kernel[grid](x, y, output, n_elements, BLOCK=256) torch.cuda.synchronize() max_error = float((output - (x + y)).abs().max()) assert torch.allclose(output, x + y) target = triton.runtime.driver.active.get_current_target() return { "version": triton.__version__, "target": str(target), "elements": n_elements, "max_abs_error": max_error, "passed": True, } def check_flash_attention() -> dict: from flash_attn import flash_attn_func torch.manual_seed(336) shape = (2, 128, 8, 64) q = torch.randn( shape, device="cuda", dtype=torch.bfloat16, requires_grad=True, ) k = torch.randn_like(q, requires_grad=True) v = torch.randn_like(q, requires_grad=True) output = flash_attn_func(q, k, v, causal=True) output.float().square().mean().backward() assert torch.isfinite(output).all() assert torch.isfinite(q.grad).all() return { "shape": list(output.shape), "dtype": "bfloat16", "forward_finite": True, "backward_finite": True, } def main() -> None: assert torch.cuda.is_available(), "DTK accelerator is not visible" device = torch.cuda.get_device_properties(0) report = { "system": { "python": platform.python_version(), "torch": torch.__version__, "hip": torch.version.hip, "device": torch.cuda.get_device_name(0), "device_count": torch.cuda.device_count(), "memory_gib": round(gib(device.total_memory), 2), "bf16_supported": torch.cuda.is_bf16_supported(), }, "matmul": [ timed_matmul(torch.float16), timed_matmul(torch.bfloat16), ], "sdpa": check_sdpa(), "triton": check_triton(), "flash_attention": check_flash_attention(), } RESULTS_PATH.write_text( json.dumps(report, ensure_ascii=False, indent=2), encoding="utf-8", ) print(json.dumps(report, ensure_ascii=False, indent=2)) if __name__ == "__main__": main() ``` 运行方式： ```bash source /opt/dtk/env.sh python hygon_bw_validation.py ``` ## 六、实测结果与正确解读方式 本次运行的关键结果如下： ```text Device: BW | VRAM: 63.98 GiB | BF16: True Python: 3.10.12 | PyTorch: 2.10.0 | HIP: 6.3.26093 MatMul float16: 2048x2048 | 0.074 ms | 231.486 TFLOPS | PASS MatMul bfloat16: 2048x2048 | 0.076 ms | 225.437 TFLOPS | PASS SDPA BF16: output=[2, 8, 128, 64] | forward/backward finite | PASS Triton: GPUTarget(backend='hip', arch='gfx936', warp_size=64) vector add | max error=0.00e+00 | PASS FlashAttention BF16: causal forward/backward finite | PASS ```   这里最值得强调的是：`231.486 TFLOPS` 和 `225.437 TFLOPS` 来自规模为 `2048 × 2048`、预热 2 次、重复 8 次的短时烟测。它没有固定频率和功耗，也没有覆盖更大矩阵、长时间稳态、并发任务和不同转置布局，因此不能当成官方峰值，更不能直接拿来与其他平台做横向排名。 这组数据真正能说明的是： - FP16 和 BF16 计算链路已经打通； - 输出没有 NaN 或 Inf； - PyTorch 原生 Attention 的反向传播正常； - Triton 能识别 `gfx936`，自定义向量加法误差为 0； - 当前定制版 FlashAttention 可以完成 BF16 causal 前向和反向。 如果要做严肃的性能报告，还需要固定软件版本、时钟/功耗状态、矩阵规模、预热次数和统计口径，并至少报告中位数、P95、显存峰值及稳态功耗。 ## 七、这台 64GB 单卡实例适合练什么 ### 1. 课程作业和算子实验 实现线性层、Embedding、RMSNorm、RoPE、SwiGLU、Multi-Head Attention、交叉熵和优化器都没有问题。建议每个模块先写一个小规模 CPU 参考实现，再与设备版本比较： ```python torch.testing.assert_close( actual.float().cpu(), expected.float().cpu(), rtol=1e-2, atol=1e-2, ) ``` 这种做法比一上来跑完整训练更容易定位错误，尤其适合检查 mask 方向、张量布局、广播和 dtype 转换。 ### 2. 小模型预训练 可以用它完成小型语言模型从分词、数据加载、训练到采样的完整闭环。为了让实验稳定，我会优先采用： - BF16 参数或自动混合精度； - 梯度累积换取更大的有效 batch； - gradient checkpointing 降低激活占用； - 训练前先用几十个 step 验证 loss 是否下降； - 每个阶段记录峰值显存和 tokens/s。 ### 3. LoRA/QLoRA 与单卡推理 64GB 显存给单卡实验留出了不错的空间，但“能加载”不等于“能全参数训练”。以 7B 模型为例，Adam 类优化器的参数、梯度和状态通常远超仅保存一份权重的空间，全参数训练还要给激活和临时缓冲区留余量。LoRA/QLoRA 更符合这类单卡实例的使用方式，具体上限仍取决于序列长度、batch size、量化算子和内核兼容性。 当前环境已经有 vLLM，可以用于验证模型加载、生成和服务接口。正式跑模型前，建议先用本地已有的小模型或离线上传的权重做 smoke test，不要一开始就搬几十 GB 的权重。 ### 4. Triton 内核练习 本次测试显示 Triton 的后端目标为： ```text GPUTarget(backend='hip', arch='gfx936', warp_size=64) ``` 这意味着可以练习向量加、Softmax、LayerNorm、矩阵乘等内核。不过从其他平台迁移 kernel 时，不要默认 warp size 一定是 32，也不要照搬只针对某一架构调出的 `num_warps` 和 block size。先保证正确，再结合实际目标做调优。 ## 八、离线环境中最容易踩的几个坑 ### 1. 外网域名能解析，不代表 HTTPS 能访问 本次环境中 DNS 可以工作，但访问 PyPI、GitHub 和 Hugging Face 的 HTTPS 请求受限。因此，直接执行 `pip install` 或 `git clone` 可能长时间等待后失败。 更稳妥的流程是： 1. 在可联网机器下载 wheel、源码包和模型文件； 2. 记录文件哈希； 3. 通过平台文件管理、JupyterLab 或 `scp` 上传； 4. 使用本地路径安装； 5. 把依赖版本和验证结果一起归档。 例如： ```bash # 联网机器 pip download -d wheels -r requirements.txt # 上传后，在远端离线安装 pip install --no-index --find-links ./wheels -r requirements.txt ``` ### 2. 每个新终端都要确认 DTK 环境 如果 `torch.cuda.is_available()` 突然从 `True` 变成 `False`，先检查是否忘记执行： ```bash source /opt/dtk/env.sh ``` 不要急着重装 PyTorch。重装很可能破坏平台已经匹配好的定制依赖。 ### 3. 自定义算子包必须做前向和反向验证 能 `import flash_attn` 只能证明 Python 找到了包，不能证明内核与当前设备完全兼容。至少应覆盖： - 目标 dtype； - causal/non-causal 模式； - 真实训练会用到的 head dimension； - 前向输出有限； - 反向梯度有限。 ### 4. 先保存 JSON，再截图 截图适合写文档，但不适合自动比较。我把系统版本、矩阵乘耗时、目标架构和误差同时写入 `validation_results.json`。以后更换镜像或升级 DTK 时，重新运行同一脚本即可做差异比较。 ## 九、我的建议：先建立最小闭环，再扩大实验 在异构加速环境中，最省时间的顺序通常是： 1. 确认 cgroup 资源边界； 2. 激活厂商软件栈； 3. 用小张量完成一次设备计算； 4. 验证核心算子的前向和反向； 5. 保存结构化结果； 6. 再接入数据集、模型和分布式组件。 这次验证之后，我对这台实例的定位比较清楚：它很适合深度学习教学、单卡工程验证、Triton 算子练习和中小规模模型实验；如果目标是长时间大模型预训练或严格的跨平台性能对比，则还需要更完整的监控、数据集、集群和基准测试方案。 对我来说，这次排查最有价值的并不是跑出一个好看的 TFLOPS 数字，而是确认了从 DTK、PyTorch、Triton 到 FlashAttention 的每一层都能给出可复现、可解释的结果。环境是否真正可用，最终还是要由一段能重复运行的代码来回答。