Two and One

Back
HPCGames 题解 A B C 题Blur image
34 min read
中文
hpc games /
parallel computing /
high performance computing /
ai infrastructure

HPCGames 题解 A B C 题

最近参加了HPCGames比赛，写下这篇Blog记录一下自己对A B C题的理解与题解
views | comments

A 题#

首先第一题是签到题，题面给出了一段位置语言的代码，让我们尝试编译并猜测其输出结果。题面代码如下：

program quine
    implicit none
    integer :: i
    character(len=1), parameter :: q = char(34), k = char(44)
    character(len=80), dimension(21) :: s
    s = [ character(len=80) :: &
        "program quine", &
        "    implicit none", &
        "    integer :: i", &
        "    character(len=1), parameter :: q = char(34), k = char(44)", &
        "    character(len=80), dimension(21) :: s", &
        "    s = [ character(len=80) :: &", &
        "    ]", &
        "    do i = 1, 6", &
        "        print '(a)', trim(s(i))", &
        "    end do", &
        "    do i = 1, 21", &
        "        if (i < 21) then", &
        "            print '(5a)', '        ', q, trim(s(i)), q, k // ' &'", &
        "        else", &
        "            print '(5a)', '        ', q, trim(s(i)), q // ' &'", &
        "        end if", &
        "    end do", &
        "    do i = 7, 21", &
        "        print '(a)', trim(s(i))", &
        "    end do", &
        "end program quine" &
    ]
    do i = 1, 6
        print '(a)', trim(s(i))
    end do
    do i = 1, 21
        if (i < 21) then
            print '(5a)', '        ', q, trim(s(i)), q, k // ' &'
        else
            print '(5a)', '        ', q, trim(s(i)), q // ' &'
        end if
    end do
    do i = 7, 21
        print '(a)', trim(s(i))
    end do
end program quine
plaintext

B 题#

这里又来到了经典的小北问答环节，结合一些理论知识和具体论文的查阅，我们可以对题目进行详细的分析和解答。

1. Amdahl & Gustafson#

某程序的代码中 10% 必须串行执行，90% 可完美并行。

  • 根据 Amdahl’s Law，无论核心数如何增加，该程序的理论最大加速比极限是 ____ 倍；
  • 若在 10 核系统中通过扩大问题规模来保持每核计算负载不变，根据 Gustafson’s Law，该系统的加速比将达到 ____ 倍。

答案

首先，根据 Amdahl 定律，加速比 S 可以通过以下公式计算：

S=1(1−P)+PNS = \frac{1}{(1 - P) + \frac{P}{N}}

其中 P 是可并行部分的比例，N 是处理器的数量。对于该问题，P = 0.9（90% 可并行），串行部分为 0.1（10% 必须串行）。当 N 趋近于无穷大时，公式简化为：

Smax=1(1−P)=10.1=10S_{max} = \frac{1}{(1 - P)} = \frac{1}{0.1} = 10

因此，该程序的理论最大加速比极限是 10 倍。

接下来，根据 Gustafson 定律，加速比 S 可以通过以下公式计算：

S=N−(1−P)×(N−1)S = N - (1 - P) \times (N - 1)

在 10 核系统中，N = 10，P = 0.9，因此：

S=10−(1−0.9)×(10−1)=10−0.1×9=10−0.9=9.1S = 10 - (1 - 0.9) \times (10 - 1) = 10 - 0.1 \times 9 = 10 - 0.9 = 9.1

因此，在 10 核系统中通过扩大问题规模来保持每核计算负载不变，该系统的加速比将达到 9.1 倍。

2. OpenMP#

以下代码使用 OpenMP 并行执行循环：

int sum = 0;
#pragma omp parallel for
for (int i = 0; i < 100; i++) {
    sum += i;
}
printf("sum = %d\n", sum);
c

关于该代码，请问以下说法中正确的是 ____ 。

选项描述
A代码一定能正确计算出 0 到 99 的和（4950）
B代码存在数据竞争， 结果不确定。
Csum变量默认为private，每个线程有自己的副本。
DOpenMP 会自动为sum变量添加原子操作，保证结果正确。

答案

正确答案是 B。

解释如下：

  • 选项 A 是错误的，因为代码中存在数据竞争，多个线程同时修改共享变量 sum，导致结果不确定。
  • 选项 B 是正确的，因为在并行执行时，多个线程可能同时读取和写入 sum，导致数据竞争，从而使得最终结果不确定。
  • 选项 C 是错误的，因为 sum 变量在默认情况下是共享的（shared），而不是私有的（private）。因此，所有线程都访问同一个 sum 变量。
  • 选项 D 是错误的，因为 OpenMP 不会自动为共享变量添加原子操作。要确保结果正确，需要显式地使用 #pragma omp atomic 或其他同步机制来保护对 sum 的访问。

3. 低精度#

已知 IEEE 754 标准的 FP32 拥有 8 位指数位。请问：

  • BF16 拥有 ____ 位指数位，____ 位尾数位
  • NVFP4 拥有 ____ 位指数位，____ 位尾数位

提示：可以查阅资料，了解 NVFP4 如何在低精度下保持较高的数值范围和动态范围。

答案

经查阅资料可知：

  • BF16 拥有 8 位指数位，7 位尾数位
  • NVFP4 拥有 2 位指数位，1 位尾数位

NVFP4 比较特殊，他只有4个bit，显然如果直接使用的话，其范围会很小，精度也不理想，但经过查阅资料可知： NVFP4 首先将一组数视为了一个块，一个块中会共享一个高精度的scale factor，确定大致的数量级，然后NVFP4只存储每个数相对于这个scale factor的偏移量，这样就能在保持较大数值范围的同时，使用更少的位数来表示每个数，从而提高了存储效率和计算速度。

4. MPI 通信#

4.1 基本原语#

4 个进程执行以下代码，每个进程有局部值 local_val，操作后每个进程都有所有进程的值。

int rank;
MPI_Comm_rank(MPI_COMM_WORLD, &rank);
int local_val = rank; // rank 为进程编号，0~3
int recv_buf[4];
/* 填这一行代码 */
c

4.2 通信器#

创建一个 2 维笛卡尔拓扑，尺寸为 2×2，行优先排列，允许环绕连接。

MPI_Comm comm_cart;
int dims[2] = {2, 2};
int periods[2] = {1, 1}; // 环绕连接
/* 填这一行代码 */
c

答案

4.1 基本原语#

可以使用 MPI_Allgather 来实现该功能。代码如下：

MPI_Allgather(&local_val, 1, MPI_INT, recv_buf, 1, MPI_INT, MPI_COMM_WORLD);
c

这行代码会将每个进程的 local_val 收集到所有进程的 recv_buf 中。

4.2 通信器#

可以使用 MPI_Cart_create 来创建一个二维笛卡尔拓扑。代码如下：

MPI_Cart_create(MPI_COMM_WORLD, 2, dims, periods, 1, &comm_cart);
c

这行代码会创建一个2x2的笛卡尔拓扑通信器 comm_cart，并允许环绕连接。

5.NCCL 延迟#

在深度学习的并行推理与训练中，进程之间会频繁进行集合通信操作。NVIDIA 的开源集合通信库 NCCL，提供了在 GPU 之间进行集合通信的高性能解决方案。

当在异构的硬件上进行大规模集合通信时，如何选择通信的算法将很影响集合通信操作的效率。为了解决这个问题，NCCL 的解决方案是：基于一套硬编码的调优常数，估算不同集合通信算法下的集合通信完成时间，由此选择最优的算法。

问题：在 NCCL 2.28 的默认调优常量中，用 NVLink 连接的两 GPU、采用 Tree 算法和 LL 协议时，在估算时每跳（单步）的硬件延迟取值为 ______ µs。

答案

根据 NCCL 2.28 的默认调优常量，当使用 NVLink 连接的两 GPU，采用 Tree 算法和 LL 协议时，每跳（单步）的硬件延迟取值为 0.6 µs。

具体参考资料可见 NCCL Github src/graph/tunning.cc ↗中的151行。

6. 高性能网络#

Rail-optimized networking 与 Clos 都是高性能网络设计方案。以下说法正确的有：

选项描述
A在 Rail-optimized 网络中，来自不同 HB 域（High-Bandwidth Domain）但具有相同 local rank 的 GPU 会被连接到同一个 rail switch 上，以减少跨域通信的延迟
B常见部署模式下，Rail-optimized 网络相比传统 Clos 网络的主要优势是完全不需要 Spine 层交换机，因此可以大大节省网络设备成本
CClos 网络因其使用 Spanning Tree Protocol (STP) 而在大规模部署时存在扩展性问题，这是 Rail-optimized 网络要解决的核心问题之一
DRail-optimized 网络保证了任何情况下集群内任意两个 GPU 之间都能以网络线速（如 400 Gbps InfiniBand）进行通信，无论它们是否在同一个 rail 中
ENCCL 2.12 引入的 PXN 特性可以结合 NVLink 和 PCI 通信来优化网络流量，这个优化对于 Rail-optimized 网络尤为重要
F对于 LLM 训练工作负载，最优的通信策略会将大部分网络流量集中在相同 local rank 的 NIC 之间，并且会多用 NVLink 等高速互联进行跨 rail 交换，这使得 Rail-optimized 架构特别适合此类场景

答案

正确答案是 A、E、F。

解释如下：

  • 选项 A 是正确的，这是 Rail-optimized 网络的核心定义。 在这种架构中，网络拓扑是根据 GPU 的 rank 进行物理隔离的。例如，所有服务器上的 0 号 GPU 都连接到同一组交换机（Rail 0），1 号 GPU 连接到另一组（Rail 1）。这使得在进行数据并行（Data Parallelism）训练时，AllReduce 等操作只需在同一个 Rail 内进行，无需跨越复杂的交换层级，大大降低了拥塞和延迟。

  • 选项 B 是错误的，Rail-optimized 仍然基于 Clos 架构，通常需要 Spine 交换机。 Rail-optimized 描述的是 Leaf 层交换机与 GPU 的连接方式以及流量的导向方式，而不是一种去除了 Spine 的新型拓扑。对于大规模集群（超过一个 Leaf 交换机的容量），Rail 0 的 Leaf 交换机之间仍然需要通过 Spine 交换机互联，以构成一个完整的 Rail 0 网络平面。

  • 选项 C 是错误的，Clos 网络并不使用 STP，且 STP 是传统以太网的痛点。 传统二层以太网使用 STP (Spanning Tree Protocol) 防止环路，这会导致大量链路被阻塞，带宽利用率低。而现代 Data Center Clos 网络（无论是基于 IP 路由的 ECMP 还是 InfiniBand）的设计初衷就是利用所有链路进行负载均衡，完全摒弃了 STP。因此，C 选项描述的前提本身就是错误的。

  • 选项 D 是错误的，Rail-optimized 并不保证“跨 Rail”通信的效率等同于“同 Rail”。 Rail-optimized 的设计哲学是“专路专用”。虽然物理上可以通过 Spine 进行跨 Rail 通信（例如 Node A 的 GPU 0 发给 Node B 的 GPU 1），但这通常不是最优路径，且可能面临 oversubscription（收敛比）的问题。实际上，这种架构倾向于利用 E 选项和 F 选项提到的技术来避免在网络层面上进行跨 Rail 数据传输。

  • 选项 E 是正确的，PXN (PCIe/NVLink Cross-NIC) 是解决 Rail 架构灵活性的关键。 在 Rail-optimized 网络中，如果 GPU 0 需要向网络中的 Rail 1 发送数据，传统的路径非常低效（走 PCIe -> CPU -> NIC -> Switch -> …）。NCCL 的 PXN 特性允许 GPU 0 通过 NVLink 直接把数据传给同机的 GPU 1，然后由 GPU 1 的 NIC（连接着 Rail 1）发送出去。这相当于在节点内部利用 NVLink 完成了“变轨”，从而充分利用 Rail 网络的优势。

  • 选项 F 是正确的，这准确描述了 LLM 训练中的混合通信模式。 在 LLM 训练中，通常结合了数据并行（DP）和模型并行（TP/PP）。

    • TP (Tensor Parallelism) 流量极大，通常限制在单机内部，完全走 NVLink。

    • DP (Data Parallelism) 需要跨机同步梯度，流量发生在相同 rank 的 GPU 之间，这完美契合 Rail-optimized 的网络路径。

    • 如果需要跨 rank 的操作（如 Pipeline Parallelism 的某些阶段或特定的 All-to-All），结合 NVLink（节点内）+ Rail（节点间）是目前最优的策略。

7. GPU#

NVIDIA 的 Hopper 架构引入了 TMA（Tensor Memory Accelerator） 以提升 GPU 内存访问效率。以下说法正确的有：

选项描述
A相比 cp.async，TMA 可以直接将数据从全局内存加载到共享内存，无需经过寄存器中转，从而能节省寄存器
B在 cutlass 的异步流水线抽象中，Producer 调用 producer_acquire 获取空闲的 buffer stage，完成数据加载后调用 producer_commit 通知 Consumer；Consumer 则通过 consumer_wait 等待数据就绪，使用完毕后调用 consumer_release 释放 buffer
C在使用 TMA 进行数据传输时，所有参与的线程都需要执行相同的 TMA 指令，TMA 硬件会自动处理线程间的协调
DCutlass Pipeline 使用多级缓冲（multi-stage buffering），通过 PipelineState 追踪当前读写的 stage index 和 phase，实现 Producer 和 Consumer 之间的流水线重叠
ETMA 的 multicast 功能允许一次 TMA 操作将同一块数据广播到 Cluster 内的多个 Thread Block 的共享内存中，减少了重复的全局内存访问
FTMA 描述符（TMA Descriptor）需要在 kernel 启动前在 host 端创建，描述符中包含了张量的形状、步长和 swizzle 模式等信息，kernel 执行时通过预取描述符（prefetch_tma_descriptor）来减少首次 TMA 操作的延迟

答案

正确答案是 B D E F.

解释 ：

  • 选项A是错误的。Ampere 架构引入的 cp.async 指令同样也是绕过寄存器（Register File），直接将数据从全局内存（GMEM）搬运到共享内存（SMEM）。
  • 选项B是正确的。这描述了 Cutlass 异步流水线中 Producer 和 Consumer 之间的交互方式，符合 Cutlass 的设计理念。
  • 选项C是错误的。TMA 操作允许线程组内的线程根据需要选择性地执行 TMA 指令，而不是所有线程都必须执行相同的指令。如果所有线程都执行，会导致重复发射多个拷贝操作（除非有特殊的掩码处理）。这一点与 Ampere 的 cp.async（通常每个线程负责一部分）不同。
  • 选项D是正确的。Cutlass Pipeline 确实使用多级缓冲，通过 PipelineState 来追踪当前读写的 stage index 和 phase，从而实现 Producer 和 Consumer 之间的流水线重叠。
  • 选项E是正确的。TMA 的 multicast 功能允许一次 TMA 操作将同一块数据广播到 Cluster 内的多个 Thread Block 的共享内存中，减少了重复的全局内存访问。
  • 选项F是正确的。TMA 描述符需要在 kernel 启动前在 host 端创建，包含张量的形状、步长和 swizzle 模式等信息，kernel 执行时通过预取描述符来减少首次 TMA 操作的延迟。

8. LLM#

对于参数如下的一个标准的 Transformer-Decoder 模型，所有的 all reduce 操作都使用 ring all reduce。假设一共有 4 张卡。

模型参数#

参数值
层数  32 层
隐藏层维度 (h)  4096
FFN 结构两层线性层，中间层维度为 4h
序列长度2048
Batch Size32
优化器Adam + 混合精度训练
精度设置参数和梯度使用 fp16，Adam 优化器状态使用 fp32（包括 momentum、variance 和 master weights）

问题#

请计算在以下三种并行方式下，进行一个 batch 的前向传播和反向传播，每张卡需要的发送量（以 GB 为单位）：

  • 数据并行：每张卡上存放完整的模型，把 batch 均匀拆分到每张卡上，分别计算完成后对梯度进行 All-Reduce 操作

  • 流水并行：按层拆分模型放到不同卡上，只需要前向传播的时候发送 activation，反向传播的时候发送 gradient。（计算通信量时只考虑中间的卡）

  • 张量并行：对于 MHA 操作，按照 head 拆分到不同卡上。对于 FFN，第一个线性层按照输出维度进行拆分，第二个线性层按照输入维度进行拆分

答案

计算过程如下：

  • 数据并行：

    • 模型参数量：

      P=32×12×(40962)P = 32 \times 12 \times (4096 ^ 2)

    • 梯度大小：

      G=P×2 bytes (fp16)=12 GBG = P \times 2 \text{ bytes (fp16)} = 12 \text{ GB}

    • 通信量：

      通信量=2×N−1N×G=18 GB\text{通信量} = 2 \times \frac{N-1}{N} \times G = 18 \text{ GB}

  • 流水并行：

    • 每层激活大小：

      A=32×2048×4096×2 bytes (fp16)=0.5 GBA = 32 \times 2048 \times 4096 \times 2 \text{ bytes (fp16)} = 0.5 \text{ GB}

    • 通信量：

      通信量=2×A=1 GB\text{通信量} = 2 \times A = 1 \text{ GB}

  • 张量并行：

    • 单次 All-Reduce 大小：

      AR=32×2048×4096×2 bytes (fp16)=0.5 GBAR = 32 \times 2048 \times 4096 \times 2 \text{ bytes (fp16)} = 0.5 \text{ GB}

    • 单次 Ring All-Reduce 通信量：

      单次通信量=2×N−1N×AR=0.75 GB\text{单次通信量} = 2 \times \frac{N-1}{N} \times AR = 0.75 \text{ GB}

    • 总通信量：

      通信量=32×4×0.75 GB=96 GB\text{通信量} = 32 \times 4 \times 0.75 \text{ GB} = 96 \text{ GB}

9. UB 互联#

在高性能计算系统中，集合通信（Collective Communication）的性能主要受带宽（Bandwidth） 与延迟（Latency） 两个因素制约。

NVIDIA 通过 NVLink 与 NVSwitch 构建 GPU 间的高速 Scale-up 互联网络，而华为则提出了 Unified Bus（UB）协议，作为面向 NPU 的统一互联与内存访问机制。UB 协议基于华为自研的 UB Switch 交换芯片，并通过高带宽物理链路 HCCS（High-Capacity Coherent System） 进行连接。

传统 AI 集群通常以 8 卡服务器为基本单元进行 Scale-out 扩展，而华为在 CloudMatrix 384（CM384） 架构中，通过两级 UB Switch 组网，将 384 颗昇腾 910C NPU 构建为一个统一的超节点（SuperPod）。在该超节点范围内，所有 NPU 均处于同一个低延迟的轨道优化网络中，实现全对等 Scale-up 互联。

CM384 进一步将 UB 网络划分为 7 个相互独立的物理平面。每颗 NPU 的 7 个 HCCS 接口分别接入不同的交换平面，从而保证大规模并行通信过程中，数据流在物理路径上完全隔离、无链路冲突。

问题#

在 CloudMatrix 384 的标准满配部署方案中，为了支撑 384 颗昇腾 910C NPU 实现无收敛、全对等的 Scale-up 互联，系统采用两级交换架构。在该超节点的物理拓扑中，分别使用了：

  • ____ 个 Level 1 UB Switch
  • ____ 个 Level 2 UB Switch
  • 最终实现了理论上 ____ GB/s 的系统级聚合带宽

假设 switch chip 提供的单个 Port 可以提供 28GB/s 的通信带宽

答案

该问题直接查阅华为 CloudMatrix 384 的白皮书 ↗即可得到答案。

10. Cache 行为分析#

假设我们需要进行一个矩阵乘法 C=A×BC=A×B。

测试环境

为了简化分析，假设：

参数类型配置
数据类型double (8 Bytes)
L1 Cache 大小4KB (4096 Bytes)
相联度直接映射 (Direct Mapped, E=1)
块大小64 Bytes（1 个 Cache Line 可存 8 个 double）
矩阵规模A,B,C 均为 64×64 的方阵 (N=64)
存储方式数组按行优先存储
内存对齐A,B,C 的起始地址均对齐到 Cache 的起始 Set

代码实现

// 假设变量 sum 已优化到寄存器中，忽略 C 的访存影响
// 仅考虑内层循环中 A 和 B 的读取
for (int j = 0; j < 64; ++j) {       // Loop 1
    for (int i = 0; i < 64; ++i) {   // Loop 2
        double sum = 0.0;
        for (int k = 0; k < 64; ++k) { // Loop 3
            sum += A[i][k] * B[k][j];
        }
        C[i][j] = sum;
    }
}
cpp

问题 10.1

我们试图分析上述代码中最内层循环 Loop 3 对矩阵 BB 的访存行为。

已知 Cache 总共有 4096/64=644096/64=64 个 Set。

在计算 C[i][j]C[i][j] 的过程中（即一次完整的 Loop 3），关于 B[k][j]B[k][j] 的 Cache Miss Rate（不命中率），下列说法正确的是：

选项描述
A12.5% - 这里有良好的空间局部性，每 8 个 double 只有 1 次 Miss
B25% - 虽然是列优先访问，但 Cache 够大，只有冷不命中
C约 50% - A 和 B 互相打架（冲突），导致一半的数据被驱逐
D100% - 发生了严重的 Cache Thrashing（抖动），每次读取都是 Miss

💡 提示：计算一下访问 B[k][j]B[k][j] 和 B[k+1][j]B[k+1][j] 时的内存地址差值（Stride），以及它们映射到的 Set Index 的跨度。

答案

正确答案是 D。

解释如下：

  • 矩阵 B 是按行优先存储的，因此访问 B[k][j] 时，k 的变化会导致访问的内存地址以列为单位跳跃。

  • 计算地址差值（Stride）：

    Stride=sizeof(double)×N=8 Bytes×64=512 Bytes\text{Stride} = \text{sizeof(double)} \times N = 8 \text{ Bytes} \times 64 = 512 \text{ Bytes}

  • 每次访问 B[k][j] 时，地址增加 512 Bytes，而每个 Cache Line 大小为 64 Bytes，因此每次访问都会跨越多个 Cache Line。

  • 计算 Set Index 的跨度：

    Set Index Span=StrideCache Line Size=512 Bytes64 Bytes=8\text{Set Index Span} = \frac{\text{Stride}}{\text{Cache Line Size}} = \frac{512 \text{ Bytes}}{64 \text{ Bytes}} = 8

  • 因为 Cache 有 64 个 Set，跨度为 8 意味着每次访问都会映射到不同的 Set，但由于 k 从 0 到 63，共有 64 次访问，这些访问会循环映射到同一组 Set 上，导致频繁的冲突和驱逐。

  • 最终结果是每次读取 B[k][j] 都会导致 Cache Miss，即 Cache Thrashing。

问题 10.2 为了进一步提升矩阵乘法的效率，我们决定使用分块技术。你将矩阵分成了 8×88×8 的小块（Block Size = 8）。

// 8x8 分块优化演示
for (int jj = 0; jj < 64; jj += 8) {
    for (int ii = 0; ii < 64; ii += 8) {
        for (int kk = 0; kk < 64; kk += 8) {
            // 在这里处理 8x8 的子块乘法
            for (int j = jj; j < jj + 8; ++j) {
                for (int i = ii; i < ii + 8; ++i) {
                     double sum = C[i][j]; // 简化写法
                     for (int k = kk; k < kk + 8; ++k) {
                         sum += A[i][k] * B[k][j];
                     }
                     C[i][j] = sum;
                }
            }
        }
    }
}
cpp

针对一个 8×88×8 的 BB 矩阵子块（假设该子块已预加载），在处理该子块内部的计算时，关于其在 L1 Cache 中的状态，下列分析正确的是：

选项描述
A一个 8×8 的子块大小为 512 Bytes，远小于 Cache 大小，因此完全没有冲突，所有数据都能驻留在 Cache 中
B尽管子块很小，但由于 B 的原始列宽（Stride）很大，导致子块内的 8 行数据全部映射到了同一个 Set 中，依然存在严重的冲突
C子块内的 8 行数据分别映射到了 8 个不同的 Set 中（Set 索引间隔为 8），且在子块计算期间不会发生自我冲突（Self-Conflict）
D分块主要是为了利用 L2/L3 Cache，对这么小的 L1 Cache (4KB) 来说，8×8 的分块没有任何意义

答案

正确答案是 C。

解释如下：

  • 一个 8×8 的子块大小为：

    Block Size=8×8×sizeof(double)=64×8 Bytes=512 Bytes\text{Block Size} = 8 \times 8 \times \text{sizeof(double)} = 64 \times 8 \text{ Bytes} = 512 \text{ Bytes}

  • 该子块的大小（512 Bytes）确实远小于 L1 Cache 大小（4KB），但关键在于访问模式。

  • 在处理该子块时，访问 B[k][j] 时，k 的变化会导致访问的内存地址以列为单位跳跃。

  • 计算 Set Index 的跨度：

    Set Index Span=StrideCache Line Size=512 Bytes64 Bytes=8\text{Set Index Span} = \frac{\text{Stride}}{\text{Cache Line Size}} = \frac{512 \text{ Bytes}}{64 \text{ Bytes}} = 8

  • 因此，子块内的 8 行数据分别映射到了 8 个不同的 Set 中，且在子块计算期间不会发生自我冲突（Self-Conflict）。

  • 分块技术有效地利用了 Cache 的空间局部性，减少了冲突，提高了数据的命中率。

C 题#

Ticker#

背景#

在分秒必争的高频交易领域，系统的响应速度直接决定了盈亏。我们的核心交易引擎在最新的鲲鹏服务器上部署后，发现了一个令人困惑的现象：虽然我们为每个交易对分配了独立的计算核心，理论上应当实现完美的线性加速，但实际吞吐量却远低于预期，甚至不如少用几个核心时的表现。

作为团队的首席性能架构师，你需要深入底层，找出阻碍性能释放的元凶，并对关键数据结构进行重构，让这台多核怪兽展现出它在这个并行任务上应有的爆发力。

任务描述#

程序包含两个文件：

main.cpp：负责启动 16 个 OpenMP 线程，生成随机模拟行情，并统计最终结果。 market.h：定义了核心数据结构 Candle 。 你需要修改 market.h 中的结构体定义，以提升程序在多核环境下的运行效率。

修改限制：允许调整结构体的定义方式、内存布局等；禁止修改成员变量的名称（如 high, vol 等），禁止修改变量类型及其精度（必须保持 double/long long）。

请提交修改后的 market.h 文件。

输入输出#

输入 format 程序从标准输入读取数据：

  1. N (long long): 每个线程模拟的 Tick 数。
  2. Seed (int): 随机种子。 输出 format 程序向标准输出打印结果：

第一行：所有合约的总成交额，即

Total Turnover=∑Price×Volume\text{Total Turnover} = \sum Price \times Volume

第二行：所有合约的 VWAP (成交量加权平均价)，计算方式为 Total Turnover / Total Volume。

样例输入

50000000 123
plaintext

样例输出

769984094225.13
174.9963
plaintext

运行环境、编译与测试#

本题将在 Huawei Kunpeng 920A (ARM64, 64-Core) 的 16 个核心上运行。评测容器镜像：cr.hpc.lcpu.dev/hpcgame/3rd-kunpeng920:latest，基础发行版为Fedora 43，安装了clang 21,flang 21和gcc 15。

可以使用以下命令编译程序：

g++ main.cpp -o main -O3 -fopenmp -std=c++17 -march=native
plaintext

我们会使用如下命令运行程序：

export OMP_NUM_THREADS=16
export OMP_PLACES=cores
export OMP_PROC_BIND=close
./main < input.txt > output.txt
plaintext

评分标准#

  • 正确性：输出结果必须与标准答案的相对误差不超过1e-12。
  • 性能：程序运行时间为 t，满分时间为t_0则得分为 min(20, 20^(t_0/t))。 共5个测试点，迭代次数有所不同。满分标准：
迭代次数满分时间最长运行时间分值
5e70.35s0.7s20
5e83.5s7.0s20
1e97.0s14.0s20
5e935.0s70.0s20
5e10349.0s698.0s20

Hint#

鲲鹏 CPU 有自己的硬件特质，了解这些特质有助于优化性能。
HPCGames 题解 A B C 题
https://www.hjcheng0602.cn/blog/hpcgamesabc/hpcgamesabc
Author Han Jincheng
Published at February 4, 2026
Copyright CC BY-NC-SA 4.0
Buy me a cup of coffee ☕.

HPCGames 题解 D E 题

SF3D 论文阅读记录
Comment seems to stuck. Try to refresh?✨