开发 9 月社区任务「hyperloglog容器开发(950)」（对标 NVIDIA cuCollections hyperloglog，使用仓库 BloomFilter 的 SIMT 工程范式）时，需要复刻 cuCollections 的 Add 路径：所有线程按网格 stride 处理 key，先在每个 AIV 核各自的局部 scratch（UB）中做 sketch 聚合（原子字节位 max），块结束再合并回全局。实测在 Atlas 950 + CANN 9.1.0 下，单核（1 block × 1024 线程）完全正确，但多核（8 block × 1024 线程）并发时每个 block 的局部 UB 聚合大量丢失：8192 字节的局部 sketch 仅约 1/8（≈1024 字节）的寄存器被更新，其余仍为 0；核数越多丢失越严重。 这导致 HLL Add 只能退化为逐 key 全局原子更新（1e8 次约 2.8~5.2ms），无法达到任务书性能基线 0.4×（8KB 档预算≈0.79ms）。

Environment / 环境信息（必填）

• 昇腾硬件：Atlas 950（Ascend950PR），单卡，128GB HBM（npu-smi info 见 Ascend950PR）
• CANN：9.1.0（/usr/local/Ascend/cann-9.1.0）
• 编译器：Bisheng（clang 15.0.5，/usr/local/Ascend/cann-9.1.0/tools/ccec_compiler/bin/bisheng），-x asc --npu-arch=dav-3510 -O3
• 头文件：kernel_operator.h + simt_api/device_atomic_functions.h + simt_api/device_sync_functions.h + simt_api/cpp/kernel_simt_atomic_intf.h
• 启动方式：Kernel<<<numBlocks, launchBytes, aclrtStream>>>(...)，内部 AscendC::Simt::VF_CALL<...>(AscendC::Simt::Dim3{1024}, ...)

Steps to reproduce the issue / 重现步骤（必填）

构造最小内核（每 block）：

1. 在局部 UB 申请 8192B sketch（已分别尝试：静态 __ubuf__ uint8_t ls[8192]、动态 get_imm(0)、动态+blockIdx*8192 偏移）；
2. 各线程按 globalThreadIndex stride 清零该 8192B；
3. asc_syncthreads()；
4. 各线程网格步长遍历 1e6 个 uint32 key：HashRank（复用仓库 hash_functions 的 xxhash64）→ 对该 block 自己的 UB sketch 做 32 位字 CAS 的字节位 max 更新（同一寄存器被多线程竞争）；
5. asc_syncthreads()；
6. 每 block 把局部 sketch 排空到不相交的 GM 区（out[blockIdx*8192 + i]），host 读取并与"宿主全量参考 sketch"（逐 key 求寄存器 max）逐字节比对。

已尝试的变体（均无效，结果完全相同：每个 block 约 7168/8192 的寄存器仍为 0）：

• 静态 __ubuf__ 数组；
• 动态 get_imm(0)，launchBytes = 8192；
• 动态 get_imm(0)，launchBytes = 8 × 8192；
• 动态 get_imm(0) + blockIdx × 8192（per-block 偏移寻址）；
• init 与 update 之间、update 与 drain 之间各调用 两次 asc_syncthreads()。

对照验证（同模型下可用的原语，均正常）：

• 8 block 各自对单字（4B）UB 原子 max：每块独立得到 1023，正确；
• 全局内存（GM）原子字节位 max、32B 向量读、单 block 内 asc_syncthreads：均正常；
• 仓库 BloomFilter 现有的高性能 Add 在此 CANN 下刻意采用 GM 工作区 route/apply 两段式，未使用多块 UB 屏障聚合。

Describe the expected behavior / 预期结果（必填）

期望：多核并发时，每个 AIV 核在其各自的局部 UB 上进行的（原子字节位 max）聚合能够正确保留全部更新（与单核结果一致），从而 HLL/类似容器可以把每核局部聚合与最后的 GM 合并路径作为高性能 Add 方案使用（对标 cuCollections 的 per-block 共享内存 sketch）。

Related log / screenshot / 日志 / 截图（必填）

8 block × 1024 线程、n=1e6、precision=13、每块 8192B 局部 sketch，与宿主参考逐字节比对结果：

block 0: same=73   bad=8119  zero=7168
block 1: same=55   bad=8137  zero=7168
block 2: same=76   bad=8116  zero=7168
block 3: same=65   bad=8127  zero=7168
block 4: same=66   bad=8126  zero=7168
block 5: same=69   bad=8123  zero=7168
block 6: same=70   bad=8122  zero=7168
block 7: same=76   bad=8116  zero=7168

单核对照：全部正确（bad=0）。

Special notes for this issue / 备注（选填）

• HyperLogLog 任务相关：实现与测试已完成、官方 1280 功能用例全绿，仅因上述限制使 Add 性能无法达到 0.4× 标杆；其余操作均达标。完整设计与自测报告将随任务交付。
• 想请教社区：多 block 并发时 per-core 局部 UB 的正确使用方式（声明/寻址/同步）是否是"每 block 各自一份"？还是此 CANN 版本对多块 SIMT-VF 的局部 UB 有已知限制/需要特定开关（如 SPLIT_CORE_VEC、动态 UB 预留）？若有修复版本，请告知最低版本号。
• 可提供最小复现代码（编译即跑的探针）与完整实验记录。