Skip to content

[mcc] __uint128_t division segfault & variable shift shl_parts failure (reproducible on 4.3.6 & 5.2.0) #153

Open
Open
[mcc] __uint128_t division segfault & variable shift shl_parts failure (reproducible on 4.3.6 & 5.2.0)#153

Description

@SISCHOI
SISCHOI
opened on Aug 25, 2026

mcc 编译器缺陷报告 / mcc Compiler Defects Report

Issue: mcc compiler defects blocking cudf 25.04 MUSA port

**Reporter context **: I'm an Mthreads internal engineer
working on the cudf/cuSpatial MUSA port (auto-musify). These defects were
found and verified during active development — full reproducers below,
happy to coordinate directly with the compiler team (mtcc) if useful.

我是摩尔线程工程师，正在做 cudf/cuSpatial 的 MUSA 移植
（auto-musify 项目）。以下缺陷均在开发过程中发现并实测复现——
完整复现脚本见下文，与编译器团队（mcc）对接沟通。

Where to file this issue / 该提交到哪里（请维护者协助路由）

I'm not sure which repo is the right place for compiler (mcc/mtcc) defects,
since mcc is not open-sourced. Filing here because this is the most active
official Mthreads repo I could find. If this belongs elsewhere, please
point me to the right tracker or route it to the compiler team (mtcc) —
much appreciated.

我不确定 mcc/mtcc 编译器缺陷应该提交到哪个仓库（mcc 未开源）。
选择在此提交是因为这是我能找到的最活跃的摩尔线程官方仓库。
如果应该提交到其他地方，请告知正确的入口，或帮忙转给编译器
团队（mtcc）——非常感谢。

Environment

  • MUSA SDK: 5.2.0 (mcc 5.2.0, clang 14.0.0, target commit 8a8cb297)
  • Hardware: MTT S5000 (mp_31)
  • Container: protenix-v2-musa:2.0.0-musa5.2
  • Project: NVIDIA cudf 25.04 → MUSA port (libcudf, 461 translation units)

Version regression test (Bug 1 & 2 verified on both 4.3.6 and 5.2.0)

We tested the same minimal reproducers with mcc 4.3.6 (commit c64ca08d)
and mcc 5.2.0 (commit 8a8cb297). Results are identical — these bugs
have existed since at least 4.3.6 and are still unfixed in 5.2.0:

Test mcc 4.3.6 mcc 5.2.0
__uint128_t / 10 (division) FAIL — Segfault FAIL — Segfault
__uint128_t << s (variable shift) FAIL — shl_parts Cannot select FAIL — shl_parts Cannot select
__uint128_t << 3 (constant shift) PASS PASS
__uint128_t * 8 (multiplication) PASS PASS

This narrows the defect to: (a) missing shl_parts lowering for variable
shifts, and (b) a crash in the division lowering path. Both are long-standing
(not 5.x regressions).

One-click reproduction (self-contained)

Save the following as repro_mcc_uint128.sh and run it on any machine with
mcc in PATH. It has zero dependencies and zero side effects (writes
only to ./repro_mcc_out/), needs no GPU (compile-only, -c), and
includes a 120s per-case timeout guard against the known compiler-hang
variant of these bugs.

bash repro_mcc_uint128.sh                 # uses default mcc from PATH
MCC=/path/to/other/mcc bash repro_mcc_uint128.sh   # test another version
repro_mcc_uint128.sh — full script (click to expand) 
#!/usr/bin/env bash
set -u
MCC="${MCC:-mcc}"
OUT="$(pwd)/repro_mcc_out"
mkdir -p "$OUT"

echo "==================================================================="
echo " mcc __uint128_t bug reproduction"
echo "==================================================================="
"$MCC" --version 2>&1 | sed 's/^/  /' | head -3
echo ""

cat > "$OUT/div.mu" << 'EOF'
// Bug 1: __uint128_t division -> Segmentation fault
__global__ void k(__uint128_t const* a, __uint128_t* out, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] / 10; }
}
int main() { return 0; }
EOF

cat > "$OUT/shift_var.mu" << 'EOF'
// Bug 2: __uint128_t variable shift -> backend 'shl_parts' Cannot select
__global__ void k(__uint128_t const* a, __uint128_t* out, int s, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] << s; }
}
int main() { return 0; }
EOF

cat > "$OUT/shift_const.mu" << 'EOF'
// Control: constant shift works fine
__global__ void k(__uint128_t const* a, __uint128_t* out, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] << 3; }
}
int main() { return 0; }
EOF

cat > "$OUT/mul.mu" << 'EOF'
// Control: multiplication works fine
__global__ void k(__uint128_t const* a, __uint128_t* out, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] * 8; }
}
int main() { return 0; }
EOF

FLAGS=(-x musa -std=c++17 -O2 --offload-arch=mp_31)

declare -A RESULT
for t in div shift_var shift_const mul; do
  rm -f "$OUT/${t}.o"
  timeout 120 "$MCC" "${FLAGS[@]}" -c "$OUT/${t}.mu" -o "$OUT/${t}.o" \
    > "$OUT/${t}.log" 2>&1
  rc=$?
  if [ "$rc" -eq 0 ] && [ -f "$OUT/${t}.o" ]; then
    RESULT[$t]="PASS"
  elif [ "$rc" -eq 124 ]; then
    RESULT[$t]="HANG(timeout 120s)"
  else
    RESULT[$t]="FAIL(exit=$rc)"
  fi
done

echo "==================== RESULTS ===================="
printf "%-28s %s\n" "uint128_div (a[i]/10)"         "${RESULT[div]}"
printf "%-28s %s\n" "uint128_shift_var (a[i]<<s)"   "${RESULT[shift_var]}"
printf "%-28s %s\n" "uint128_shift_const (a[i]<<3)" "${RESULT[shift_const]}"
printf "%-28s %s\n" "uint128_mul (a[i]*8)"          "${RESULT[mul]}"
echo ""
for t in div shift_var shift_const mul; do
  if [[ "${RESULT[$t]}" == FAIL* ]]; then
    echo "[$t] first error:"
    grep -E 'error|Cannot select|Segmentation' "$OUT/${t}.log" | head -2 | sed 's/^/  /'
  fi
done
echo "DONE"

Verified output on our machine (2026-08-25)

Environment:

Hardware : MTT S5000 (mp_31)
OS       : Ubuntu 22.04 (container: protenix-v2-musa:2.0.0-musa5.2)
SDK      : MUSA Toolkit 5.2.0  (/usr/local/musa)
mcc      : 5.2.0, clang 14.0.0, mtcc commit 8a8cb2971e3084fc442baeacb7443e3d73263dc8
GPU      : NOT required — the script is compile-only (-c)

Actual output:

==================== RESULTS ====================
uint128_div (a[i]/10)        FAIL(exit=254)
uint128_shift_var (a[i]<<s)  FAIL(exit=70)
uint128_shift_const (a[i]<<3) PASS
uint128_mul (a[i]*8)         PASS

---- first error line of each failure ----
[div]
  clang-14: error: unable to execute command: Segmentation fault (core dumped)
[shift_var]
  fatal error: error in backend: Cannot select: 0x...: i64,i64 = shl_parts 0x..., 0x..., 0x...
  clang-14: error: clang frontend command failed with exit code 70 (use -v to see invocation)

For the 4.3.6 cross-check we extracted mcc 4.3.6 (commit c64ca08d) from the
official musa_toolkits_rc4.3.6.tar.gz into an isolated /tmp directory and
ran the same script via MCC=/tmp/.../bin/mcc bash repro_mcc_uint128.sh
(with -nogpulib --no-musa-version-check --musa-path=... appended since the
extracted tree lacks the device runtime libs). Results identical to 5.2.0.

Bug 1: __uint128_t division — backend shl_parts selection failure

Minimal reproducer

// file: uint128_div.mu
__global__ void test_div(__uint128_t const* a, __uint128_t* out, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] / 10; }
}
int main() { return 0; }

Reproduce command

mcc -x musa -std=c++17 -O2 --offload-arch=mp_31 -c uint128_div.mu -o uint128_div.o

Expected

Compiles successfully (nvcc handles this fine; __uint128_t division is supported in CUDA).

Actual

fatal error: error in backend: Cannot select: 0x...: i64,i64 = shl_parts ...
clang-14: error: clang frontend command failed with exit code 70

Impact on cudf

cudf/fixed_point/detail/floating_conversion.hpp uses __uint128_t as the
intermediate type for decimal↔floating-point conversion (ShiftingRep for
double). Division by powers of 10 (divide_power10) is used throughout.
This blocks all 35 binaryop/compiled/*.mu translation units — the
entire cudf binary operations JIT module.

Workaround attempted (all failed)

  • -O0 through -O3: all fail
  • -fno-inline, -fno-builtin, -fno-vectorize: all fail
  • __attribute__((noinline)) on the division function: fails
  • Replacing division with loop multiplication by reciprocal: mathematically
    incorrect for exact decimal conversion

Bug 2: __uint128_t variable shift — LLVM APInt::trunc crash

Minimal reproducer

// file: uint128_shift.mu
__global__ void test_shift(__uint128_t const* a, __uint128_t* out, int s, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] << s; }
}
int main() { return 0; }

Reproduce command

mcc -x musa -std=c++17 -O2 --offload-arch=mp_31 -c uint128_shift.mu -o uint128_shift.o

Actual

PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/
Stack dump:
0.  Program arguments: /usr/local/musa-5.2.0/bin/clang-14 -cc1 ...
1.  <eof> parser at end of file
2.  Code generation
3.  Running pass 'Function Pass Manager' on module '...'
4.  Running pass 'Early CSE' on function '...'
 #3  llvm::APInt::trunc(unsigned int) const
 #14 llvm::isKnownNonZero(...)
 #17 llvm::SimplifyInstruction(...)

Note: Constant shifts (x << 3) work fine. Only variable shifts
(x << s where s is a runtime value) crash.

Impact on cudf

Same file as Bug 1: floating_conversion.hpp uses shifting_rep << pow2
where pow2 is a runtime variable. This independently blocks all
binaryop compiled TUs.

Bug 3: Large template TUs segfault without diagnostics (~32 TUs)

Affected files (representative list)

cudf/src_musa/groupby/hash/compute_groupby.mu
cudf/src_musa/groupby/hash/groupby.mu
cudf/src_musa/groupby/hash/compute_aggregations.mu
cudf/src_musa/groupby/hash/compute_aggregations_null.mu
cudf/src_musa/reductions/mean.mu
cudf/src_musa/reductions/min.mu
cudf/src_musa/reductions/max.mu
cudf/src_musa/reductions/std.mu
cudf/src_musa/reductions/var.mu
cudf/src_musa/reductions/sum.mu
cudf/src_musa/reductions/product.mu
cudf/src_musa/quantiles/quantiles.mu
cudf/src_musa/quantiles/quantile.mu
cudf/src_musa/round/round.mu
cudf/src_musa/sort/is_sorted.mu
cudf/src_musa/stream_compaction/unique.mu
cudf/src_musa/stream_compaction/distinct_count.mu
cudf/src_musa/search/contains_scalar.mu
cudf/src_musa/transform/one_hot_encode.mu
... (32 total)

Reproduce command

mcc -x musa -std=c++17 -O2 --offload-arch=mp_31 \
  -DCUDF_MUSA_PORT=1 -DRMM_MUSA_PORT=1 -DCUCO_MUSA_PORT=1 \
  -I<cudf_include_musa> -I<cudf_src_musa> -I<rmm_include> \
  -c cudf/src_musa/reductions/mean.mu -o mean.o

Actual

Segmentation fault (core dumped)

No compiler diagnostics. No stack dump. Just segfault.

Additional notes

  • Some of these TUs also compile when using different optimization levels,
    but the majority segfault at all levels from -O0 to -O3.
  • These are all files with heavy template instantiation (thrust/cub/cuco
    templates expanded for all cudf data types).

Bug 4: Compiler deadlock — hangs forever, CPU=0%, never exits

Affected files

cudf/src_musa/groupby/hash/create_sparse_results_table.mu
(possibly others; at least 1 confirmed)

Symptom

$ mcc -x musa ... create_sparse_results_table.mu &
$ ps -o pid,stat,etime,time,cmd -p <pid>
  PID STAT     ELAPSED TIME COMMAND
70650 R+      80:23  0:00 mcc -x musa ... create_sparse_results_table.mu

The process shows TIME=0:00 (zero CPU time consumed) even after 80+
minutes. It never exits, never produces output, never crashes. It must be
manually killed.

This is distinct from Bug 3 (segfault) — this is a hang, not a crash.

Impact

Any build system that doesn't have per-TU timeout will hang indefinitely.
In parallel builds, this blocks the entire batch.

Bug 5: CUB MaxPolicyT::Invoke incompatible with MUSA

Error signature

error: no matching function for call to 'Invoke'
    error = MaxPolicyT::Invoke(ptx_version, dispatch);

Root cause

CUB's algorithm dispatch mechanism uses PTX version to select
implementation strategy. MUSA does not have PTX. The MUSA SDK's bundled
CUB library (/usr/local/musa/include/cub/) does not properly handle
this, resulting in template resolution failure.

Affected cudf TUs

Approximately 27-40 TUs that use cub::DeviceReduce, cub::DeviceSort,
cub::DeviceScan, cub::DeviceSegmentedReduce, etc.

Related: array-of-references in CUB agent templates

error: 'input_items' declared as array of references of type 'signed char &&'
error: 'input_items' declared as array of references of type 'long &&'

The CUB agent templates (e.g., agent_scan_by_key.cuh) expand to
T(&)[N] (array of references) which is ill-formed C++. nvcc tolerates
this as an extension; mcc correctly rejects it, but this means the
bundled CUB is not compatible with cudf's usage patterns.

Summary of impact on cudf 25.04 MUSA port

Category TU count Percentage
Total non-binaryop TUs 426 100%
Compile successfully (after source-level fixes) 307 72%
Blocked by Bug 1+2 (uint128, binaryop only) 35 8.2%
Blocked by Bug 3 (segfault) ~32 7.5%
Blocked by Bug 5 (CUB dispatch) ~40 9.4%
Blocked by missing third-party deps (nvcomp etc.) ~17 4.0%
Total blocked by compiler/SDK issues ~119 28%

Environment details

$ mcc --version
clang version 14.0.0 (git@sh-code.mthreads.com:sw/mtcc.git 8a8cb2971e3084fc442baeacb7443e3d73263dc8)
mcc version 5.2.0
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /usr/local/musa/bin

$ mthreads-gmi
mthreads-gmi:2.4.2  Driver Version:5.2.0-server
GPU 0: MTT S5000

中文版 / Chinese Version

Where to file this issue / 该提交到哪里（请维护者协助路由）

I'm not sure which repo is the right place for compiler (mcc/mtcc) defects,
since mcc is not open-sourced. Filing here because this is the most active
official Mthreads repo I could find. If this belongs elsewhere, please
point me to the right tracker or route it to the compiler team (mtcc) —
much appreciated.

我不确定 mcc/mtcc 编译器缺陷应该提交到哪个仓库（mcc 未开源）。
选择在此提交是因为这是我能找到的最活跃的摩尔线程官方仓库。
如果应该提交到其他地方，请告知正确的入口，或帮忙转给编译器
团队（mtcc）——非常感谢。

环境信息

  • MUSA SDK: 5.2.0（mcc 5.2.0，clang 14.0.0，commit 8a8cb297）
  • 硬件: MTT S5000（mp_31 架构）
  • 容器: protenix-v2-musa:2.0.0-musa5.2
  • 项目: NVIDIA cudf 25.04 → MUSA 移植（libcudf，共 461 个编译单元）

版本回归测试（Bug 1 & 2 在 4.3.6 和 5.2.0 上均实测复现）

我们用mcc 4.3.6（commit c64ca08d）和 mcc 5.2.0（commit 8a8cb297）
分别编译了相同的最小复现代码，结果完全一致——这些 bug 至少从
4.3.6 时代就存在，5.2.0 仍未修复：

测试 mcc 4.3.6 mcc 5.2.0
__uint128_t / 10（除法） 失败 — 段错误 失败 — 段错误
__uint128_t << s（变量移位） 失败 — shl_parts 不可选 失败 — shl_parts 不可选
__uint128_t << 3（常量移位） 通过 通过
__uint128_t * 8（乘法） 通过 通过

这说明缺陷可精确圈定为：(a) 变量移位的 shl_parts lowering 缺失；
(b) 除法 lowering 路径的崩溃。均为长期存在的老 bug（非 5.x 回归）。

一键复现（自包含脚本）

把下面的脚本保存为 repro_mcc_uint128.sh，在任何装了 mcc 的机器上
直接运行即可。零依赖、零副作用（只在当前目录写 ./repro_mcc_out/），
不需要 GPU（纯编译 -c），并对每个用例加了 120 秒超时保护
（防同类 bug 的"编译器挂起"变体把脚本卡死）：

bash repro_mcc_uint128.sh                          # 用 PATH 里的默认 mcc
MCC=/path/to/other/mcc bash repro_mcc_uint128.sh   # 测其他版本
repro_mcc_uint128.sh 完整脚本（点击展开） 
#!/usr/bin/env bash
set -u
MCC="${MCC:-mcc}"
OUT="$(pwd)/repro_mcc_out"
mkdir -p "$OUT"

echo "==================================================================="
echo " mcc __uint128_t bug reproduction"
echo "==================================================================="
"$MCC" --version 2>&1 | sed 's/^/  /' | head -3
echo ""

cat > "$OUT/div.mu" << 'EOF'
// Bug 1: __uint128_t division -> Segmentation fault
__global__ void k(__uint128_t const* a, __uint128_t* out, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] / 10; }
}
int main() { return 0; }
EOF

cat > "$OUT/shift_var.mu" << 'EOF'
// Bug 2: __uint128_t variable shift -> backend 'shl_parts' Cannot select
__global__ void k(__uint128_t const* a, __uint128_t* out, int s, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] << s; }
}
int main() { return 0; }
EOF

cat > "$OUT/shift_const.mu" << 'EOF'
// 对照组：常量移位正常
__global__ void k(__uint128_t const* a, __uint128_t* out, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] << 3; }
}
int main() { return 0; }
EOF

cat > "$OUT/mul.mu" << 'EOF'
// 对照组：乘法正常
__global__ void k(__uint128_t const* a, __uint128_t* out, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] * 8; }
}
int main() { return 0; }
EOF

FLAGS=(-x musa -std=c++17 -O2 --offload-arch=mp_31)

declare -A RESULT
for t in div shift_var shift_const mul; do
  rm -f "$OUT/${t}.o"
  timeout 120 "$MCC" "${FLAGS[@]}" -c "$OUT/${t}.mu" -o "$OUT/${t}.o" \
    > "$OUT/${t}.log" 2>&1
  rc=$?
  if [ "$rc" -eq 0 ] && [ -f "$OUT/${t}.o" ]; then
    RESULT[$t]="PASS"
  elif [ "$rc" -eq 124 ]; then
    RESULT[$t]="HANG(120秒超时)"
  else
    RESULT[$t]="FAIL(exit=$rc)"
  fi
done

echo "==================== RESULTS ===================="
printf "%-28s %s\n" "uint128_div (a[i]/10)"         "${RESULT[div]}"
printf "%-28s %s\n" "uint128_shift_var (a[i]<<s)"   "${RESULT[shift_var]}"
printf "%-28s %s\n" "uint128_shift_const (a[i]<<3)" "${RESULT[shift_const]}"
printf "%-28s %s\n" "uint128_mul (a[i]*8)"          "${RESULT[mul]}"
echo ""
for t in div shift_var shift_const mul; do
  if [[ "${RESULT[$t]}" == FAIL* ]]; then
    echo "[$t] 首个错误行:"
    grep -E 'error|Cannot select|Segmentation' "$OUT/${t}.log" | head -2 | sed 's/^/  /'
  fi
done
echo "DONE"

我们的实测环境与实测输出（2026-08-25）

复现环境：

硬件     : MTT S5000（mp_31 架构）
系统     : Ubuntu 22.04（容器：protenix-v2-musa:2.0.0-musa5.2）
SDK      : MUSA Toolkit 5.2.0（/usr/local/musa）
mcc      : 5.2.0，clang 14.0.0，mtcc commit 8a8cb2971e3084fc442baeacb7443e3d73263dc8
GPU      : 不需要——脚本纯编译（-c），任何装了 mcc 的 x86 机器可跑

实测输出（完整、未删节的结果部分）：

==================== RESULTS ====================
uint128_div (a[i]/10)        FAIL(exit=254)
uint128_shift_var (a[i]<<s)  FAIL(exit=70)
uint128_shift_const (a[i]<<3) PASS
uint128_mul (a[i]*8)         PASS

---- first error line of each failure ----
[div]
  clang-14: error: unable to execute command: Segmentation fault (core dumped)
[shift_var]
  fatal error: error in backend: Cannot select: 0x...: i64,i64 = shl_parts 0x..., 0x..., 0x...
  clang-14: error: clang frontend command failed with exit code 70 (use -v to see invocation)

4.3.6 交叉验证方法：从官方 musa_toolkits_rc4.3.6.tar.gz 中把 mcc 4.3.6
（commit c64ca08d）解压到独立的 /tmp 目录（不触碰系统环境），用
MCC=/tmp/.../bin/mcc bash repro_mcc_uint128.sh 运行（因解压目录缺设备
运行时库，追加 -nogpulib --no-musa-version-check --musa-path=... 参数）。
结果与 5.2.0 完全一致。

Bug 1：__uint128_t 除法 — 后端 shl_parts 指令选择失败

最小复现代码

// 文件名: uint128_div.mu
__global__ void test_div(__uint128_t const* a, __uint128_t* out, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] / 10; }
}
int main() { return 0; }

复现命令

mcc -x musa -std=c++17 -O2 --offload-arch=mp_31 -c uint128_div.mu -o uint128_div.o

预期结果

编译成功（nvcc 可以正常处理；CUDA 完整支持 __uint128_t 除法）。

实际结果

fatal error: error in backend: Cannot select: 0x...: i64,i64 = shl_parts ...
clang-14: error: clang frontend command failed with exit code 70

对 cudf 的影响

cudf/fixed_point/detail/floating_conversion.hpp 使用 __uint128_t 作为
十进制↔浮点数转换的中间类型（double 的 ShiftingRep）。除以 10 的幂
（divide_power10）在整个文件中大量使用。

此 bug 导致 全部 35 个 binaryop/compiled/*.mu 编译单元无法编译——
即 cudf 的整个二元运算 JIT 模块。

已尝试的绕过方案（全部失败）

  • -O0 到 -O3：全部失败
  • -fno-inline、-fno-builtin、-fno-vectorize：全部失败
  • 在除法函数上加 __attribute__((noinline))：失败
  • 用循环乘以倒数替代除法：对精确十进制转换数学上不正确

Bug 2：__uint128_t 变量移位 — LLVM APInt::trunc 崩溃

最小复现代码

// 文件名: uint128_shift.mu
__global__ void test_shift(__uint128_t const* a, __uint128_t* out, int s, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) { out[i] = a[i] << s; }
}
int main() { return 0; }

复现命令

mcc -x musa -std=c++17 -O2 --offload-arch=mp_31 -c uint128_shift.mu -o uint128_shift.o

实际结果

PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/
Stack dump:
0.  Program arguments: /usr/local/musa-5.2.0/bin/clang-14 -cc1 ...
1.  <eof> parser at end of file
2.  Code generation
3.  Running pass 'Function Pass Manager' on module '...'
4.  Running pass 'Early CSE' on function '...'
 #3  llvm::APInt::trunc(unsigned int) const
 #14 llvm::isKnownNonZero(...)
 #17 llvm::SimplifyInstruction(...)

注意：常量移位（x << 3）可以正常工作。只有变量移位
（x << s，其中 s 是运行时变量）会崩溃。

对 cudf 的影响

与 Bug 1 相同的文件：floating_conversion.hpp 中使用
shifting_rep << pow2，其中 pow2 是运行时变量。
此 bug 独立地导致所有 binaryop 编译单元无法编译。

Bug 3：大型模板编译单元编译时段错误（无任何诊断信息，约 32 个编译单元）

受影响的文件（代表性列表）

cudf/src_musa/groupby/hash/compute_groupby.mu
cudf/src_musa/groupby/hash/groupby.mu
cudf/src_musa/groupby/hash/compute_aggregations.mu
cudf/src_musa/groupby/hash/compute_aggregations_null.mu
cudf/src_musa/reductions/mean.mu
cudf/src_musa/reductions/min.mu
cudf/src_musa/reductions/max.mu
cudf/src_musa/reductions/std.mu
cudf/src_musa/reductions/var.mu
cudf/src_musa/reductions/sum.mu
cudf/src_musa/reductions/product.mu
cudf/src_musa/quantiles/quantiles.mu
cudf/src_musa/quantiles/quantile.mu
cudf/src_musa/round/round.mu
cudf/src_musa/sort/is_sorted.mu
cudf/src_musa/stream_compaction/unique.mu
cudf/src_musa/stream_compaction/distinct_count.mu
cudf/src_musa/search/contains_scalar.mu
cudf/src_musa/transform/one_hot_encode.mu
...（共 32 个）

复现命令

mcc -x musa -std=c++17 -O2 --offload-arch=mp_31 \
  -DCUDF_MUSA_PORT=1 -DRMM_MUSA_PORT=1 -DCUCO_MUSA_PORT=1 \
  -I<cudf_include_musa> -I<cudf_src_musa> -I<rmm_include> \
  -c cudf/src_musa/reductions/mean.mu -o mean.o

实际结果

Segmentation fault (core dumped)

无编译器诊断信息，无堆栈转储，直接段错误。

补充说明

  • 这些文件都包含大量模板实例化（thrust/cub/cuco 模板为所有 cudf
    数据类型展开）
  • 部分文件在降低优化级别后可以编译，但大多数在 -O0 到 -O3 所有
    级别下都段错误

Bug 4：编译器死锁 — 永久挂起，CPU 使用率 0%，永不退出

受影响的文件

cudf/src_musa/groupby/hash/create_sparse_results_table.mu
（可能还有其他文件；至少确认此 1 个）

症状

$ mcc -x musa ... create_sparse_results_table.mu &
$ ps -o pid,stat,etime,time,cmd -p <pid>
  PID STAT     ELAPSED TIME COMMAND
70650 R+      80:23  0:00 mcc -x musa ... create_sparse_results_table.mu

进程显示 TIME=0:00（CPU 时间为零），即使已运行超过 80 分钟。
它永不退出、永不产生输出、永不崩溃。必须手动杀死。

这与 Bug 3（段错误）不同——这是一个挂起，不是崩溃。

影响

任何没有按编译单元设置超时的构建系统都会无限期挂起。
在并行构建中，这会阻塞整个批次。

Bug 5：CUB MaxPolicyT::Invoke 与 MUSA 不兼容

错误特征

error: no matching function for call to 'Invoke'
    error = MaxPolicyT::Invoke(ptx_version, dispatch);

根因分析

CUB 的算法分发机制使用 PTX 版本号来选择实现策略。MUSA 没有 PTX。
MUSA SDK 自带的 CUB 库（/usr/local/musa/include/cub/）没有正确
处理这个问题，导致模板解析失败。

受影响的 cudf 编译单元

约 27-40 个使用 cub::DeviceReduce、cub::DeviceSort、
cub::DeviceScan、cub::DeviceSegmentedReduce 等的编译单元。

相关问题：CUB agent 模板中的引用数组

error: 'input_items' declared as array of references of type 'signed char &&'
error: 'input_items' declared as array of references of type 'long &&'

CUB agent 模板（如 agent_scan_by_key.cuh）展开为
T(&)[N]（引用数组），这是不合法的 C++。nvcc 将其作为扩展容忍了；
mcc 正确地拒绝了它，但这意味着自带的 CUB 与 cudf 的使用模式不兼容。

对 cudf 25.04 MUSA 移植的影响汇总

类别 编译单元数量 百分比
非二元运算总编译单元 426 100%
编译成功（经过源码级修复后） 307 72%
被 Bug 1+2 阻塞（uint128，仅限 binaryop） 35 8.2%
被 Bug 3 阻塞（段错误） ~32 7.5%
被 Bug 5 阻塞（CUB 分发） ~40 9.4%
被缺失的第三方依赖阻塞（nvcomp 等） ~17 4.0%
被编译器/SDK 问题阻塞的总数 ~119 28%

环境详细信息

$ mcc --version
clang version 14.0.0 (git@sh-code.mthreads.com:sw/mtcc.git 8a8cb2971e3084fc442baeacb7443e3d73263dc8)
mcc version 5.2.0
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /usr/local/musa/bin

$ mthreads-gmi
mthreads-gmi:2.4.2  Driver Version:5.2.0-server
GPU 0: MTT S5000

备注

  • 以上所有 bug 均在 MUSA SDK 5.2.0 发布说明的"已知问题与限制"中未被记录
  • 这些问题在 CUDA 12.x / nvcc 环境下均不存在
  • 我们的迁移项目（cuSpatial/cudf → MUSA）已因此被阻塞约 28% 的编译单元
  • 希望能在后续 SDK 版本中修复，或提供已知的绕过方案

Activity

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions