[Bug] trandom 将循环代码展开精度可以通过，用for循环则精度有问题

[yuqiha]

Component

PTO Dialect / ODS (include/PTO/IR)

Description

可跑通代码已提 pr （展开版本）：#328
跑不通代码（for循环版本，其他地方基本上没有动）： https://github.com/yuqiha/PTOAS_fork/blob/error_test/lib/TileOps/trandom_template.py

ST 代码建议使用pr中的，代码基本上就是将展开循环写成了for循环，所以应该模板没有问题，agent定位原因如下：

Reproduction (minimal)

python3 test/tilelang_st/script/run_st.py -r sim -v a5 -t trandom &> ST_tset.log

Expected behavior

pass

Actual behavior / error logs

[ERROR] Mismatch: max diff=4232234439.0 at flat idx=909 (golden=-2121847184.0, output=2110387255.0)
[ERROR] int32_4x256: compare failed

Git commit

e5d45dedaa9fa15de299e8ddbdb6e07b7f50e27d

Host platform

None

Target Ascend arch (if relevant)

None

PTOAS build level (if relevant)

None

[KurrinQu]

看起来确实是一个ccec的编译器问题：

• 右边循环内使用每次迭代vmull指令的计算结果更新tmp_ctr1和tmp_ctr3
• 对应左边汇编中v18和v19的值在vloop内并没有更新，正确的做法应该是在vloop末尾，用v16和v12分别更新v18和v19。
• 右边代码生成的LLVM IR中，数据流是正确的，怀疑是编译器后端代码生成错误：

; ModuleID = 'ptoas.hivm.official'
source_filename = "ptoas.hivm.official"

declare <256 x i1> @llvm.hivm.pset.b32(i32)

declare <64 x i32> @llvm.hivm.vbr.s32.v300(i32)

declare <64 x i32> @llvm.hivm.vbr.u32.v300(i32)

declare <64 x i32> @llvm.hivm.vci.v64s32(i32, i32)

declare { <64 x i32>, <256 x i1> } @llvm.hivm.vaddc.v64u32(<64 x i32>, <64 x i32>, <256 x i1>)

declare { <64 x i32>, <256 x i1> } @llvm.hivm.vaddcs.v64u32(<64 x i32>, <64 x i32>, <256 x i1>, <256 x i1>)

declare { <64 x i32>, <64 x i32> } @llvm.hivm.vmull.v64u32(<64 x i32>, <64 x i32>, <256 x i1>)

declare <64 x i32> @llvm.hivm.vxor.v64u32.x(<64 x i32>, <64 x i32>, <256 x i1>)

declare <64 x i32> @llvm.hivm.vadds.v64u32.x(<64 x i32>, i32, <256 x i1>)

declare { <64 x i32>, <64 x i32> } @llvm.hivm.vintlv.v64u32(<64 x i32>, <64 x i32>)

declare { <256 x i1>, i32 } @llvm.hivm.plt.b32.v300(i32)

declare void @llvm.hivm.vstsx1.v64u32(<64 x i32>, ptr addrspace(6), i32, i32, i32, <256 x i1>)

declare void @llvm.hivm.MOV.UB.TO.OUT.ALIGN.V2.DV(ptr addrspace(1), ptr addrspace(6), i64, i64)

define void @TRANDOM_int32_4x256(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2) #0 {
  %4 = getelementptr i32, ptr addrspace(1) %0, i64 0
  %5 = load i32, ptr addrspace(1) %4, align 4
  %6 = getelementptr i32, ptr addrspace(1) %0, i64 1
  %7 = load i32, ptr addrspace(1) %6, align 4
  %8 = getelementptr i32, ptr addrspace(1) %1, i64 0
  %9 = load i32, ptr addrspace(1) %8, align 4
  %10 = getelementptr i32, ptr addrspace(1) %1, i64 1
  %11 = load i32, ptr addrspace(1) %10, align 4
  %12 = getelementptr i32, ptr addrspace(1) %1, i64 2
  %13 = load i32, ptr addrspace(1) %12, align 4
  %14 = getelementptr i32, ptr addrspace(1) %1, i64 3
  %15 = load i32, ptr addrspace(1) %14, align 4
  br label %16

16:                                               ; preds = %aivscope.latch, %3
  %17 = phi i64 [ %122, %aivscope.latch ], [ 0, %3 ]
  %18 = icmp slt i64 %17, 1
  br i1 %18, label %19, label %123

19:                                               ; preds = %16
  %20 = call <256 x i1> @llvm.hivm.pset.b32(i32 0)
  %21 = call <64 x i32> @llvm.hivm.vbr.s32.v300(i32 %9)
  %22 = call <64 x i32> @llvm.hivm.vbr.s32.v300(i32 %11)
  %23 = call <64 x i32> @llvm.hivm.vbr.s32.v300(i32 %13)
  %24 = call <64 x i32> @llvm.hivm.vbr.s32.v300(i32 %15)
  %25 = call <64 x i32> @llvm.hivm.vbr.s32.v300(i32 %5)
  %26 = call <64 x i32> @llvm.hivm.vbr.s32.v300(i32 %7)
  %27 = call <64 x i32> @llvm.hivm.vbr.u32.v300(i32 0)
  %28 = call <64 x i32> @llvm.hivm.vbr.u32.v300(i32 -766435501)
  %29 = call <64 x i32> @llvm.hivm.vbr.u32.v300(i32 -845247145)
  %30 = call <64 x i32> @llvm.hivm.vci.v64s32(i32 0, i32 0)
  %31 = call { <64 x i32>, <256 x i1> } @llvm.hivm.vaddc.v64u32(<64 x i32> %21, <64 x i32> %30, <256 x i1> %20)
  %32 = extractvalue { <64 x i32>, <256 x i1> } %31, 0
  %33 = extractvalue { <64 x i32>, <256 x i1> } %31, 1
  %34 = call { <64 x i32>, <256 x i1> } @llvm.hivm.vaddcs.v64u32(<64 x i32> %22, <64 x i32> %27, <256 x i1> %33, <256 x i1> %20)
  %35 = extractvalue { <64 x i32>, <256 x i1> } %34, 0
  %36 = extractvalue { <64 x i32>, <256 x i1> } %34, 1
  %37 = call { <64 x i32>, <256 x i1> } @llvm.hivm.vaddcs.v64u32(<64 x i32> %23, <64 x i32> %27, <256 x i1> %36, <256 x i1> %20)
  %38 = extractvalue { <64 x i32>, <256 x i1> } %37, 0
  %39 = extractvalue { <64 x i32>, <256 x i1> } %37, 1
  %40 = call { <64 x i32>, <256 x i1> } @llvm.hivm.vaddcs.v64u32(<64 x i32> %24, <64 x i32> %27, <256 x i1> %39, <256 x i1> %20)
  %41 = extractvalue { <64 x i32>, <256 x i1> } %40, 0
  %42 = extractvalue { <64 x i32>, <256 x i1> } %40, 1
  br label %43

43:                                               ; preds = %74, %19
  %44 = phi i64 [ %120, %74 ], [ 0, %19 ]
  %45 = phi <64 x i32> [ %109, %74 ], [ %32, %19 ]
  %46 = phi <64 x i32> [ %112, %74 ], [ %35, %19 ]
  %47 = phi <64 x i32> [ %115, %74 ], [ %38, %19 ]
  %48 = phi <64 x i32> [ %118, %74 ], [ %41, %19 ]
  %49 = icmp slt i64 %44, 4
  br i1 %49, label %50, label %121

50:                                               ; preds = %43
  br label %51

51:                                               ; preds = %60, %50
  %52 = phi i64 [ %73, %60 ], [ 0, %50 ]
  %53 = phi <64 x i32> [ %68, %60 ], [ %45, %50 ]
  %54 = phi <64 x i32> [ %65, %60 ], [ %46, %50 ]
  %55 = phi <64 x i32> [ %70, %60 ], [ %47, %50 ]
  %56 = phi <64 x i32> [ %62, %60 ], [ %48, %50 ]
  %57 = phi <64 x i32> [ %71, %60 ], [ %25, %50 ]
  %58 = phi <64 x i32> [ %72, %60 ], [ %26, %50 ]
  %59 = icmp slt i64 %52, 10
  br i1 %59, label %60, label %74

60:                                               ; preds = %51
  %61 = call { <64 x i32>, <64 x i32> } @llvm.hivm.vmull.v64u32(<64 x i32> %53, <64 x i32> %28, <256 x i1> %20)
  %62 = extractvalue { <64 x i32>, <64 x i32> } %61, 0
  %63 = extractvalue { <64 x i32>, <64 x i32> } %61, 1
  %64 = call { <64 x i32>, <64 x i32> } @llvm.hivm.vmull.v64u32(<64 x i32> %55, <64 x i32> %29, <256 x i1> %20)
  %65 = extractvalue { <64 x i32>, <64 x i32> } %64, 0
  %66 = extractvalue { <64 x i32>, <64 x i32> } %64, 1
  ; %54的数据流是正确的
  %67 = call <64 x i32> @llvm.hivm.vxor.v64u32.x(<64 x i32> %66, <64 x i32> %54, <256 x i1> %20)
  %68 = call <64 x i32> @llvm.hivm.vxor.v64u32.x(<64 x i32> %67, <64 x i32> %57, <256 x i1> %20)
  ; %56的数据流是正确的
  %69 = call <64 x i32> @llvm.hivm.vxor.v64u32.x(<64 x i32> %63, <64 x i32> %56, <256 x i1> %20)
  %70 = call <64 x i32> @llvm.hivm.vxor.v64u32.x(<64 x i32> %69, <64 x i32> %58, <256 x i1> %20)
  %71 = call <64 x i32> @llvm.hivm.vadds.v64u32.x(<64 x i32> %57, i32 -1640531527, <256 x i1> %20)
  %72 = call <64 x i32> @llvm.hivm.vadds.v64u32.x(<64 x i32> %58, i32 -1150833019, <256 x i1> %20)
  %73 = add i64 %52, 1
  br label %51

74:                                               ; preds = %51
  %75 = call { <64 x i32>, <64 x i32> } @llvm.hivm.vintlv.v64u32(<64 x i32> %53, <64 x i32> %55)
  %76 = extractvalue { <64 x i32>, <64 x i32> } %75, 0
  %77 = extractvalue { <64 x i32>, <64 x i32> } %75, 1
  %78 = call { <64 x i32>, <64 x i32> } @llvm.hivm.vintlv.v64u32(<64 x i32> %54, <64 x i32> %56)
  %79 = extractvalue { <64 x i32>, <64 x i32> } %78, 0
  %80 = extractvalue { <64 x i32>, <64 x i32> } %78, 1
  %81 = call { <64 x i32>, <64 x i32> } @llvm.hivm.vintlv.v64u32(<64 x i32> %76, <64 x i32> %79)
  %82 = extractvalue { <64 x i32>, <64 x i32> } %81, 0
  %83 = extractvalue { <64 x i32>, <64 x i32> } %81, 1
  %84 = call { <64 x i32>, <64 x i32> } @llvm.hivm.vintlv.v64u32(<64 x i32> %77, <64 x i32> %80)
  %85 = extractvalue { <64 x i32>, <64 x i32> } %84, 0
  %86 = extractvalue { <64 x i32>, <64 x i32> } %84, 1
  %87 = call { <256 x i1>, i32 } @llvm.hivm.plt.b32.v300(i32 256)
  %88 = extractvalue { <256 x i1>, i32 } %87, 0
  %89 = extractvalue { <256 x i1>, i32 } %87, 1
  %90 = call { <256 x i1>, i32 } @llvm.hivm.plt.b32.v300(i32 %89)
  %91 = extractvalue { <256 x i1>, i32 } %90, 0
  %92 = extractvalue { <256 x i1>, i32 } %90, 1
  %93 = call { <256 x i1>, i32 } @llvm.hivm.plt.b32.v300(i32 %92)
  %94 = extractvalue { <256 x i1>, i32 } %93, 0
  %95 = extractvalue { <256 x i1>, i32 } %93, 1
  %96 = call { <256 x i1>, i32 } @llvm.hivm.plt.b32.v300(i32 %95)
  %97 = extractvalue { <256 x i1>, i32 } %96, 0
  %98 = extractvalue { <256 x i1>, i32 } %96, 1
  %99 = mul i64 %44, 256
  %100 = getelementptr i32, ptr addrspace(6) null, i64 %99
  call void @llvm.hivm.vstsx1.v64u32(<64 x i32> %82, ptr addrspace(6) %100, i32 0, i32 2, i32 0, <256 x i1> %88)
  %101 = add i64 %99, 64
  %102 = getelementptr i32, ptr addrspace(6) null, i64 %101
  call void @llvm.hivm.vstsx1.v64u32(<64 x i32> %83, ptr addrspace(6) %102, i32 0, i32 2, i32 0, <256 x i1> %91)
  %103 = add i64 %99, 128
  %104 = getelementptr i32, ptr addrspace(6) null, i64 %103
  call void @llvm.hivm.vstsx1.v64u32(<64 x i32> %85, ptr addrspace(6) %104, i32 0, i32 2, i32 0, <256 x i1> %94)
  %105 = add i64 %99, 192
  %106 = getelementptr i32, ptr addrspace(6) null, i64 %105
  call void @llvm.hivm.vstsx1.v64u32(<64 x i32> %86, ptr addrspace(6) %106, i32 0, i32 2, i32 0, <256 x i1> %97)
  %107 = call <64 x i32> @llvm.hivm.vbr.u32.v300(i32 64)
  %108 = call { <64 x i32>, <256 x i1> } @llvm.hivm.vaddc.v64u32(<64 x i32> %45, <64 x i32> %107, <256 x i1> %20)
  %109 = extractvalue { <64 x i32>, <256 x i1> } %108, 0
  %110 = extractvalue { <64 x i32>, <256 x i1> } %108, 1
  %111 = call { <64 x i32>, <256 x i1> } @llvm.hivm.vaddcs.v64u32(<64 x i32> %46, <64 x i32> %27, <256 x i1> %110, <256 x i1> %20)
  %112 = extractvalue { <64 x i32>, <256 x i1> } %111, 0
  %113 = extractvalue { <64 x i32>, <256 x i1> } %111, 1
  %114 = call { <64 x i32>, <256 x i1> } @llvm.hivm.vaddcs.v64u32(<64 x i32> %47, <64 x i32> %27, <256 x i1> %113, <256 x i1> %20)
  %115 = extractvalue { <64 x i32>, <256 x i1> } %114, 0
  %116 = extractvalue { <64 x i32>, <256 x i1> } %114, 1
  %117 = call { <64 x i32>, <256 x i1> } @llvm.hivm.vaddcs.v64u32(<64 x i32> %48, <64 x i32> %27, <256 x i1> %116, <256 x i1> %20)
  %118 = extractvalue { <64 x i32>, <256 x i1> } %117, 0
  %119 = extractvalue { <64 x i32>, <256 x i1> } %117, 1
  %120 = add i64 %44, 1
  br label %43

121:                                              ; preds = %43
  br label %aivscope.latch

aivscope.latch:                                   ; preds = %121
  %122 = add i64 %17, 1
  br label %16, !llvm.loop !3

123:                                              ; preds = %16
  %124 = getelementptr i32, ptr addrspace(1) %2, i64 0
  call void @llvm.hivm.MOV.UB.TO.OUT.ALIGN.V2.DV(ptr addrspace(1) %124, ptr addrspace(6) null, i64 34359738432, i64 1125899906843648)
  ret void
}

attributes #0 = { "target-cpu"="dav-c310-vec" "target-features"="+ATOMIC,+ArchV130,+AregRedefinable,+ArithmeticBf16,+AtomicForB8 ,+F8e4m3,+F8e5m2,+F8e8m0,+FFTSBlk,+Fp4e1m2x2,+Fp4e2m1x2,+LDExtRefine,+MOVX8,+MSTX,+SPR7bits,+SyncV,+dav-c310-vec" }

!llvm.module.flags = !{!0}
!hivm.annotations = !{!1, !2}

!0 = !{i32 2, !"Debug Info Version", i32 3}
!1 = !{ptr @TRANDOM_int32_4x256, !"kernel", i32 1}
!2 = !{ptr @TRANDOM_int32_4x256, !"kernel_with_simd", i32 1}
!3 = distinct !{!3, !4}
!4 = !{!"llvm.loop.aivector_scope"}

• 在黄区使用2.28的代码构建的ccec编译器，生成代码应该是正确的。

[Zhendong404]

Track in issue #219