[HCU][TLE] Support tle.exclusive_cumsum

[alexshuang]

HCU backend support for tle.exclusive_cumsum primitive and TLE topk tutorial.

Lowering path

• Avoid incorrectly rewriting the LDS write operation tt.store(smem_ptr + offsets) into a buffer store op.
• Add the Passthrough="alwaysinline" attribute to functions with noinline=false in the LLVM IR, to avoid generating dynamic LDS instructions.

Unit Test

Reduce num_warps to avoid triton.runtime.errors.OutOfResources: out of resource: threads, Required: 2048, Hardware limit: 1024.

Performance Data

Benchmark source:

• python/tutorials/tle/03-topk.py.

Environment:

• Docker: harbor.baai.ac.cn/flagtree/flagtree-hcu3.6-py310-torch2.4.1-ubuntu22.04:202604-base

Performance:

• TLE vs Triton: 1.25x speedup in large size

[sunnycase]

Hi, thanks for adding HCU support for tle.exclusive_cumsum.

I reproduced the NV CI failure locally: test_tle_cumsum_exclusive_and_total reports the same four floating-point failures while all five integer cases pass. I traced the regression to commit 4e9c373a7, specifically createWarpScanStepI32 in third_party/tle/dialect/lib/Conversion/TleToLLVM/ExclusiveCumsumOpToLLVM.cpp.

Before this change, the NV inline-PTX path was restricted to i32:

auto intTy = dyn_cast<IntegerType>(val.getType());
if (!intTy || intTy.getWidth() != 32)
  return Value();

That guard was removed while adding the HCU branch. As a result, the non-HCU/NV branch now also emits the following integer operation for floating-point values:

shfl.sync.up.b32 ...
@p add.s32 r0, r0, $1;

The generated f32 LLVM IR confirms this:

tail call float asm "... @p add.s32 r0, r0, $1; ...",
                    "=r,r,r,r,r"(float %20, ...)

This performs signed integer addition on IEEE-754 bit patterns and then interprets the result as a float, which explains the very large values and NaNs in NV CI. The integer cases pass because int8/int16 are promoted to i32; bf16 is promoted to f32 and fails for the same reason.

A path-control experiment also isolates the issue: the same f32 kernel is correct with num_warps=2 when it uses the typed fallback (maximum error 4.77e-7), but produces an error of about 1.84e38 with num_warps=4 when it enters this warp-scan fast path.

Could you please restore the IntegerType && width == 32 guard specifically in the non-HCU/PTX branch? Non-i32 values should return an empty Value and use the existing targetInfo.shuffleUp + createAdd fallback, which emits FAdd for floats. Keeping the guard inside the non-HCU branch preserves the new type-aware HCU implementation as well as the optimized NV i32 path.

Please also rerun:

python -m pytest -s -vv \
  python/test/tle/unit/test_tle_cumsum.py::test_tle_cumsum_exclusive_and_total

It would also be helpful to extend the PTX regression coverage with a floating-point case so this type-domain regression cannot recur. Thank you!