From 81e421b1845b59ba66dde1e6006664c2f33fde62 Mon Sep 17 00:00:00 2001 From: i3wanna2 <15910307812@163.com> Date: Mon, 16 Jun 2025 17:01:49 +0800 Subject: [PATCH 1/3] add .h --- include/CMakeLists.txt | 1 + include/flagtree/CMakeLists.txt | 3 + include/flagtree/Common/UnifiedHardwareBase.h | 27 ++++++++ .../gbuilder/aipu.bin | Bin 0 -> 92564 bytes .../gbuilder/aipubt_gb_add_kernel.py | 63 ++++++++++++++++++ .../gbuilder/aipugb.log | 42 ++++++++++++ .../gbuilder/compass_ir.bin | 0 .../gbuilder/compass_ir.txt | 46 +++++++++++++ .../gbuilder/op_lib/add_kernel.o | Bin 0 -> 4064 bytes .../runtime/dump/cluster0/core0/tec0.log | 2 + .../runtime/dump/cluster0/core0/tec1.log | 2 + .../runtime/dump/cluster0/core0/tec2.log | 2 + .../runtime/dump/cluster0/core0/tec3.log | 2 + .../runtime/dump/hts_aiff.log | 0 .../runtime/dump/tsm.log | 0 .../gbuilder/op_lib/add_kernel.cl | 35 ++++++++++ .../gbuilder/op_lib/add_kernel.log | 3 + .../gbuilder/op_lib/add_kernel.o | Bin 0 -> 4064 bytes third_party/aipu/unified_hardware_aipu.h | 24 +++++++ 19 files changed, 252 insertions(+) create mode 100644 include/flagtree/CMakeLists.txt create mode 100644 include/flagtree/Common/UnifiedHardwareBase.h create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipu.bin create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipubt_gb_add_kernel.py create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipugb.log create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.bin create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.txt create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/op_lib/add_kernel.o create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec0.log create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec1.log create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec2.log create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec3.log create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/hts_aiff.log create mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/tsm.log create mode 100644 third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.cl create mode 100644 third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.log create mode 100644 third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.o create mode 100644 third_party/aipu/unified_hardware_aipu.h diff --git a/include/CMakeLists.txt b/include/CMakeLists.txt index 109c292fe..93a0c83ea 100644 --- a/include/CMakeLists.txt +++ b/include/CMakeLists.txt @@ -1 +1,2 @@ add_subdirectory(triton) +add_subdirectory(flagtree) diff --git a/include/flagtree/CMakeLists.txt b/include/flagtree/CMakeLists.txt new file mode 100644 index 000000000..27c703b3c --- /dev/null +++ b/include/flagtree/CMakeLists.txt @@ -0,0 +1,3 @@ +add_subdirectory(Conversion) +add_subdirectory(Dialect) +add_subdirectory(Target) diff --git a/include/flagtree/Common/UnifiedHardwareBase.h b/include/flagtree/Common/UnifiedHardwareBase.h new file mode 100644 index 000000000..dc165c728 --- /dev/null +++ b/include/flagtree/Common/UnifiedHardwareBase.h @@ -0,0 +1,27 @@ +#ifndef UNIFIED_HARDWARE_BASE_H +#define UNIFIED_HARDWARE_BASE_H + +#include + +namespace mlir { + +//this is the unified hardware abstraction for hardware +//to determined if these abstraction is specified, using std::optional is needed +//using in passes: if(uh_flagtree->xxx()){...} + +class UnifiedHardware{ + +public: + virtual ~UnifiedHardware() = default; + + //DMA + virtual std::optional getAllocSpaceForDMATag() const { + return std::nullopt; + } + +}; + + +} // namespace mlir + +#endif // UNIFIED_HARDWARE_BASE_H diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipu.bin b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipu.bin new file mode 100644 index 0000000000000000000000000000000000000000..6af3938cf016148514948367d2b348b8eaf435df GIT binary patch literal 92564 zcmeI3e{dYdmB)K_SMpjb+hfT_GWf?zRth7GSV^`e|Hv^~N!XSR2piwIsHuj+m@5XL9Cp5k=Ki{<^FYKXSH1HlgZL z1(ox?o_@VE-IlC~v*edw` z>Mmhh#e&SkRug7+?)-c9DdA2gx?e7nE+Nat-`I;t41c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0-qlOn^?-YbZ?=BpY)$;dV{8Q zJ4@ZB_QmG;{0y zEMx1FvzW|1{+*WPJxnd{)wEC3d794G^ej#LH9cF?1)2_Mx=_<|G(A_-^E7R0dcLNM zG+nIe1)5%{>8mtdqUlmiFVghYn!ZNUi#5Ge(=h%aKAr}uijcsxJ!^4yMjW5+Ag;+1Lf${%#ryIj77ocR2jH>2inTd+P?KesPjs^;fH z%;$xLmNFm0=H>klk2?Jgn@lP1ExNosukY}UV!n8b)bf+YH=X!Hm+n>K@65kEere-a z19C_|cBx+}&)8N~YwH90S+5m~@s@hwGp{o^x5tBDn~Kj=>xY5-%UF{fTmGGk2Xf7m2 z%kP{l{=J!8`*x<7A(`{>D4neTuJWZrYI!|>b)LFi*7e)wcd5o}rsFwX{rY^8?PJT; z^}(oBwWm?l-tDUPb*MTosp|Y;RnK}v)&9e(p8c|_3r?##@Q$hrCsaM>V^z=fs`smT z#i}-!t9pK|s*4&`UA$e@3p!N2FsbURhJ}{vcgZ8FogG&7qL)>D^=VaK^Ny+)PpEpy z$EsfHEwa~twmfXFr^||k&W%6Yr`#Wml&U@KZpH3RsoKZxQS5msRp+z667}Q#X;w;T zS-)Saf3{YCfmVM&tG`gIe~woFT&@0jTK%S0|9q|fBCY;nt^NgC{R_4FuhQx-(dsYN z>R+VQf3;TsHCp|PwfdK6^)J=xU#8X1$H6euzmel%Ff&x<)}|S1aOv>F%Chs(NQn9Bf&K3}ueb(nNl(l|JG6)E3fv!ipgmllYZP`b zw1;aC?k{}M9@~FA`a^rTK7e))w8!>8jsDOc+aLFPK5Wm#o9Les_6oFnM$!K*^dA#C zEd7rQJ@I4oPYc~C{m0SY+hF$}_Rpezzqmj8vgkic*bOG;a|PNxrqJh#(LX5kko3p> z`9ul& z_ThfbGmicn(LW<}Iof?$^xuO1OsrQ;(%%&N+->L|6uMpdhlHMJY!vxr3GKUD@b$E! z|3+c=wh3L1c3%ShhtWTU{)^Fn6#XAS|1tDOyYD#qkDz~A=qBkuj{Xm$e@5tb>7Paa z!|2b7?fmd9!TL?1&mBSkpwQ)L_l1O>cn%&ADEeoF{wL|56`ENq-aP)2rwo6#8}PAHw3>}43CS-lw2$T4zehad=#PB78KJ{y_hr!^`4~*hk51`t;(RKW z{TF&j`iF#`!26wH3GK_r`diT-?|0rdq07`_X?C{R4PE8WTE<_S||f z5AR25VegdwcE%hrR_7>|kh;hT%?F>yYKH%KAlu({pfN1#gzt2N_B=t_d=M|s zG_!AL&x6ir&qKW0_;kS7$QV3+cmY$isXf2{0LS0Q#CUVxr(ygtM*3+jUY<7I2XMSo zINq-Dx9{irc(b6$9~nP$>0u>*CLO=u<9IjV{#4e_zk`;B>$0yp zx=nkY)u8E_KHr$GJb(Ui=bky%Vm8i7*nj`#BN)Fs``>L?4D)r#J7)G|N2~RV4$Jbf zGi=d>nf=xyV*G<%Gf`G-X0hK^`EoO{yw=g{8_mS~ji++Ik6wlOw|={0f3+BNhf6O9 zopkA1(8Dg>2>KCdPp!zuSGSv?Yd&`DuYrGwcYdxt--r0eT=4@V{=}SD{n_2u`R!aR zOr;sLJ>9aJSwA2CuwgcP!`26`W!4=>ORU-Kbz6TO^r4zu&ZiH3Bi|Z1=;ZVLZ-f19 zD;z63vw_FsJkC7m*d?ETKc2GVjT=5?zK$DU1SdpvSlJR+&rnREJf=TkP zWr?#sZCRq$w?$j8Hsx8?7Ngjzs5rA>KJa|&m2kM=hC<6%BjW$cyVGjUE3ukZ2z_jW z-+F$m#41XQ`Ms$b`v1+`GjF^9KR#7r9dPk@=gAVQVvWNm+mUH_@5@35JPG6H6B7-c zRgV4q_tbcM))ZO=hI79;eX7t3lsTI72vn=Oph4B%`9jad^~!s@V?PV=id^w|yk%Pa zA}!uBy*`Z3efg&V;};pD_ImWDap|w&y3L2%57v+7@`dY>pwI_{zM%j!>=^v|T=eJX z@Ui_bG-etq@$X^pKi?Q``Q|+9y?TdV3yh~MU#b0l#WP7fUiKK)kABujg=o+2>!!#* z9N*tH{~9s>9$Y2rdye_SyH^%kd0sI;d3!cO{>lEFU*NYcLjFm;sf=0oJ?`Yo2+t?T zFKOrBA3=T{a`7I11^91p__v$@{wp25XRVrVJjQ3uw<0ayF2H=^@!aFz0__Rle8P5H zfobwF@EP+l%|t$8J!x6rK^*_eGT%_1EB{`bqP)yM?%(qGyxe%$`U?5?k`|BOTlo6> zkZXK+jG-0IcyK)$5<2s6#JCUcQRls4K8?WoxYEhd^kL8~`z>p?u`e-V?8~;?@ZEOq zkDpwBcTcf@H2s&a$J|%Wk0xzB{$2AwjQQ{VSnT)y4c6Zl$bWZxmHdbG_rF~Ezj+R` zR=D#2&?N0Y3+u5v-UV%bT!i`Iw*Ls$Yqx&BT5V4Ytj{Z5>*HH5fY#Q-N-cjYwS2uE z*ULW4hwE{@#ColRIKE9^VncpcKK5w&Bg=o*d~DO==~~7=|LIQ|oTm%4^>nA^uj`w% z^4qlXy5{wr_x>U`9xa-Gi{^i_*0RT{0QP6F2g-iydS;!3Ifd=M?C;gGKRPcTuICl7 zUTt*EUtWHTmOm3({_H-g=11Tgg_b`4a=c~zH@WhMU%M-pgex$*hov$Jg0XPH=! z0{a8X{x0yT`hJA$e}N_LZ-JxY_gjHa=q;aIUjE{>ALi=g{ui!w=4;@E#&99b&)vAc zIE%l34!$7voBI0-pErGSyZZZ5@_7yR`*5C?`Uw0UaO-Dazwg#h!+zhbpYjTQ5dVIG zaqM^(H)a~PY`!ib%kOdJ_ns3?%J_}stM`KshArzNJY(2$&P+7D?fl-jpZP=Gr;4+@ zp7U0Tb-&*qiajaxne?yK@e6|8GDBS_gq{05v63n6e`15G4=vX8U#j|vpkMW0q1i*4 zyZJ_R65h?|(5154h*+OIV)!z$Z7}$1aKW zR^PuKo1#4AO9IB+=Fek6J|rMVvOM2UZnrGRp9GJ8a`_bbHiP`R0QuA6%AXcj{L6{Cy``@5qZs(yC6s*g0P`k7jx@1JBp$ou$a#!Q^GocqrO_J*!z6!+0j3ZRwL!$685XEd;tCl$bqaZ54n)w`DE|cUHR~p*g!Jg z-xsW|s;{c6XB|D!zV2Y|Vs%w@RVcV3ncTJR#v2D>-O+fc#!BAPaYJ`(U?4uoWmR;b zw_{hlFIv^n-@C46?W*ef8yhya?uhjCe5E(i8B4~y`+{Ayq1v@|^|i4z?Om&*vHIxR zP*;6jU92{=wywIpe(kCbFjUvpUdOh^`}V{J)&;k2ZNfJFL91uH>EXTco=)Bf$zcs( zsIIH4Ulm+Y+a3#bu35b*R9o8_s;jB4Ufo{TUQ^v!T^kM6)YWyi*GB6?b!)q>3r3Ui z&PXCQ(6z3*rgmj$eQ$JMBpUDPiX@}GiJn*_(I4*{jP&j8T_#EnRPYghPU6HPMSAQfKzb7WH^>;<~_QeO+b;i1)dwT}g zCkLY)dxX7ab=~^G=s$#5p66m&h zj;;osEphahK_^Nby_wrxdIz^Ja_mvisjD5`%gej$cY|gw`+XqCUG~2L-R81C4*IxD z9|fIu)&C0UjLUusbn7)v{XYV2&35#Ef*y16`4#9HBzjmwps<+NFwK+qUK4ViSG&$@T<5NOYh3l#xazHO)m!7Lx5ia(ja~1S$e-T~cM!WncICHiYi`vtQo3}-v;-=;X=<;p1wZh5-via7?&8@XFK%;}+u9Ir-WuTzf#Tb8?#|`yTQ-Zjc@J}fcQGe;A9I3t@-G@T z-x1lo?RLA(yazq)J-iM+pzVW^uDyL7$-V8}1JT5;NbaO2(wXduL_0epaFgtd^)TnO zssohpg>@dR(Vj?8yq#Z`F?*x&zFer>$z?Gka1R(5bcJRS`#u&Kh;{T2blNvD?GoIb zv@3(L4*g0TZY=Gw?s(s1hFD*xcKNO@xJM53_w;w~jWMXy#?q0vJJJ~)j7A3gBb}WC zQ(ceT6^#!n#=Bzi?p=f2t-E(h!9b_!%()^V3hn<;yz)lp7)J^u6>z_+-*zw>WUVbJ%KHE|F$dJ?4 z1C5pSa33CGIv?n&?GxklXKKH|1$x;|KJ}whyW2F0z`la5CI}U1c(3;AOb{y2oM1xKm>>Y5g-CYfCvx) zB0vO)01+SpMBs}}pa`1JUjyLp4RrMP!FzjsgS(>1UGO4qPYkca7b}2I6KeTpWfC=f zsRs|nKl0!&Gw=hx>&2P!FMmfN1Rv3c-bB17Hqe{w=2txV#=hon`|@({vnY<5m<^i8 z=W)7wp*LM|9zxDy;dxXEDO=WnTK$NvIuYcFyD literal 0 HcmV?d00001 diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipubt_gb_add_kernel.py b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipubt_gb_add_kernel.py new file mode 100644 index 000000000..2b8096089 --- /dev/null +++ b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipubt_gb_add_kernel.py @@ -0,0 +1,63 @@ +import numpy as np +from AIPUBuilder.core import BuilderOpPlugin, register_optype, BuilderParams, Tensor +from AIPUBuilder.plugin_loader import register_plugin, PluginType + + +op_type = register_optype("DSL_add_kernel") + +def _get_addr(base, offset_in_byte): + if offset_in_byte == 0: + return base + + ret = Tensor(base) + ret.memory_offset().set_base_offset(base.memory_offset()) + ret.memory_offset().relative_offset += offset_in_byte + return ret + +@register_plugin(PluginType.Builder, 0) +class DSL_add_kernelPlugin(BuilderOpPlugin): + def get_graph_pattern(self): + return ([("useless", op_type)], []) + + def get_score(self): + return 10 + + def set_target(self, target): + return True + + def check_params(self, nodes): + return True + + def setup(self, sgnode, nodes): + sgnode.attrs["keeping_layout"] = False + return True + + def generate_code_name(self, sgnode, nodes): + return "op_lib/add_kernel.o" + + def generate_descriptor(self, sgnode, nodes): + desc_base = sgnode.attrs["descriptorbase"] + ro = BuilderParams() + return ro + + def generate_params(self, sgnode, nodes): + desc_base = sgnode.attrs["descriptorbase"] + ro = BuilderParams() + ro.append(sgnode.inputs[0]) + ro.append(sgnode.inputs[1]) + ro.append(sgnode.outputs[0]) + value = np.int32(nodes[0].params["var_14"]) + ro.append(int(value.view("int32"))) + value = np.int32(nodes[0].params["var_33"]) + ro.append(int(value.view("int32"))) + value = np.int32(nodes[0].params["var_35"]) + ro.append(int(value.view("int32"))) + value = np.int32(nodes[0].params["var_37"]) + ro.append(int(value.view("int32"))) + value = np.int32(nodes[0].params["var_3"]) + ro.append(int(value.view("int32"))) + value = np.int32(nodes[0].params["var_39"]) + ro.append(int(value.view("int32"))) + value = np.int32(nodes[0].params["var_41"]) + ro.append(int(value.view("int32"))) + return ro diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipugb.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipugb.log new file mode 100644 index 000000000..ff4bae16f --- /dev/null +++ b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipugb.log @@ -0,0 +1,42 @@ +Current Working Directory: /root/flagtree/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder +Command Line: aipugb compass_ir.txt -w compass_ir.bin --target X2_1204 +Standard Output & Error: +[I] [IRChecker] Start to check IR: /root/flagtree/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.txt +[I] [IRChecker] model_name: unknown +[I] [IRChecker] IRChecker: All IR pass +[I] [tools.cpp : 352] BuildTool version: 6.1.17784. Build for target X2_1204 PID: 42160 +[I] [tools.cpp : 851] global cwd: /tmp/e77881b08dfe9a5bd2c506fa5be9126b7d478215d54cd681637723a3b0f97 +[I] [graph.cpp :1605] loading graph weight: compass_ir.bin size: 0x0 +[I] [actg.cpp : 471] new sgnode with actg: 0 +[I] [datalayout_schedule.cpp:2443] Layout loss: 0 +[I] [datalayout_schedule.cpp:2444] Layout scheduling ... +[I] [datalayout_schedule.cpp:2447] The layout loss for graph unknown: 0 +[I] [datalayout_schedule.cpp:1397] The graph unknown post optimized score:0 +[I] [datalayout_schedule.cpp:1401] layout schedule costs: 0.079781ms +[I] [IRChecker] Start to check IR: +[I] [IRChecker] model_name: cost_model +[I] [IRChecker] IRChecker: All IR pass +[I] [load_balancer.cpp:3032] enable multicore schedule optimization for load balance strategy 1 it may degrade performance on single core targets. +[I] [soms_scheduler.cpp: 195] [EVAL] init time t1: 47 ms +[I] [soms_scheduler.cpp: 201] [EVAL] unsafe check time t2: 10 ms +[I] [soms_scheduler.cpp: 947] get max in loop 0 +[I] [soms_scheduler.cpp: 214] [EVAL] mem assignment time t3: 26 ms +[I] [builder.cpp:1829] [EVAL] duration time of mem allocation: 94 ms +[W] [subgraphop.cpp: 613] empty kernel name, will trying to use the first kernel in elf for node:subgraph_2_dsl_add_kernel +[I] [builder.cpp:1948] The graph DDR Footprint requirement(estimation) of feature maps: +[I] [builder.cpp:1949] Read and Write:51.94KB +[I] [builder2.cpp:1666] memory statistics for this graph (unknown) +[I] [builder.cpp: 635] Text section: 0x00003640 Bytes ( 0.013MB) +[I] [builder.cpp: 635] Weight section: 0x00000000 Bytes ( 0.000MB) +[I] [builder.cpp: 635] Ro section: 0x00000500 Bytes ( 0.001MB) +[I] [builder.cpp: 635] Desc section: 0x00000000 Bytes ( 0.000MB) +[I] [builder.cpp: 635] Stack section: 0x00040400 Bytes ( 0.251MB) +[I] [builder.cpp: 635] Workspace section: 0x00000000 Bytes ( 0.000MB) +[I] [builder.cpp: 635] Activations section: 0x00010fc4 Bytes ( 0.066MB) +[I] [builder.cpp: 635] Total: 0x00054f04 Bytes ( 0.332MB) +[I] [tools.cpp :1480] - compile time: 0.242 s +[I] [tools.cpp :1382] With GM optimization, DDR Footprint stastic(estimation): +[I] [tools.cpp :1389] Read and Write:51.94KB +[I] [tools.cpp :1436] - draw graph time: 0 s +[I] [tools.cpp :1991] remove global cwd: /tmp/e77881b08dfe9a5bd2c506fa5be9126b7d478215d54cd681637723a3b0f97 +Total errors: 0, warnings: 1 diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.bin b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.bin new file mode 100644 index 000000000..e69de29bb diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.txt b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.txt new file mode 100644 index 000000000..241b02f9f --- /dev/null +++ b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.txt @@ -0,0 +1,46 @@ + +model_name=unknown +layer_number=3 +data_format=NHWC +precision=float32 +batch_size=1 +input_tensors=[input0,input1] +output_tensors=[temp_var_0] +compat_quantized_model=false + +layer_id=0 +layer_name=0_input +layer_type=Input +layer_bottom=[] +layer_bottom_shape=[] +layer_bottom_type=[] +layer_top=[input0] +layer_top_shape=[[4432]] +layer_top_type=[float32] + +layer_id=1 +layer_name=1_input +layer_type=Input +layer_bottom=[] +layer_bottom_shape=[] +layer_bottom_type=[] +layer_top=[input1] +layer_top_shape=[[4432]] +layer_top_type=[float32] + +layer_id=2 +layer_name=2_dsl_add_kernel +layer_type=DSL_add_kernel +layer_bottom=[input0, input1] +layer_bottom_shape=[[4432], [4432]] +layer_bottom_type=[float32, float32] +layer_top=[temp_var_0] +layer_top_shape=[[4432]] +layer_top_type=[float32] +var_14=4432 +var_33=5 +var_35=1 +var_37=1 +var_3=1 +var_39=0 +var_41=0 diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/op_lib/add_kernel.o b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/op_lib/add_kernel.o new file mode 100644 index 0000000000000000000000000000000000000000..06a5c3b0a5857417646b93388c392cd70556b395 GIT binary patch literal 4064 zcmcgvZ){Ul6hCimq4ag~HXw}=WP_QFF`B^;N0{NR9S%EUX3%K(Qb+ej*8OSIRh=IU zWdnolkR^n~7)WOHgCH3AfYD4yQ8UT#&j?@q=Hej4jXzZqH4;+KdGFo6``*xr7+-RG z@9&;-&pG#=-#zbj_m*u>NsH&o69HU!NSCbFx+<6tq7>lmBB!*YtC%f?3Gl(M zcW*iv+Y$S*aVg8gvdc>DY^0UliZJGcF?e`N$sIWaEWg{!${b3(e36o)F|`#dmH5g! z6K@JC@vFh}ce#Ajq10~LZpud&A&v-mCE`8-*C8Gha1ik!jHj+u^EbL(seReS0dU*_)7@0X~O)c4Bc{ z5;2WKr_J@As)BE_quD2&rj5Ek&phYs^~)5^C+W8=v-)ev@3!b~burc=En*cFV+|#E zo)^yq0-iNq=3J-S|Hbhl+v4)ErYapDZ*a2{BR*D|&1?JFx&`-tSl=U$@4NFpwpTch z%kTPF#Rl_ya*@f#D^KcJE{aPxrluO({ib~SQ>(qM4PNGvlv-~8i}PM4mz$W*Ltbme zo<=Knl<4>&NNM|bJZs9w(O#)wpX@ES*)O%(TW zKV0Ao>1WDxoGxQ2hVs}y1_jE&}lsX8WVSIwDOH?{9C@2+W0nu z_lfL@`43?{vT>h`xS2c$ALR$+qv0pZ{iTikSC=~nTyyY`)^PsOICOMjAs-sP5B|}+ zh3fBHLVn1`K$V#f!o#&Xo;(R=xc^5!@_5Ak(|G#M4}g8vGe-QGlzlZ7XH~Cml)Au-a-15h4 z{p*$23jS#QOr?2W<&(jH7w>1(_`Vp=AJpPapG;|dwdeo8#_)=o=_==rS|jEO8uOTGY{UcpgxcNaPpbV%@JC`jJ!)?f!r{1@=m;m&NG#DFPIgAR z!aLPnoxRZ3-X88!6TND;F!V>d!QUH8s{ZC@o?p|KY>RZ+EVOlYbl9d|?NGzt)p@()_8t?e=Lk0{*^aBH6YR``3D?w+7nY)21ilAN+oVI>ymx+P@bA{&O`S{_hr3 zr~T8Vo*qUGdBPv!A3dbckCbp3+WQ`RT33XPWVDPL>DFSSCuML>N4kvU&DeMhBpVid W66UgxaZL94LwdsqNY4FIjDG?FfP`!S literal 0 HcmV?d00001 diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec0.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec0.log new file mode 100644 index 000000000..02e7b0e2c --- /dev/null +++ b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec0.log @@ -0,0 +1,2 @@ +[WARN]: tec0's vfp_ctrl is disabled, then decode fp inst exception. +[WARN]: tec0 decode fp instruction exception jump from pc=0x00007280 to new pc=0x00005030. diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec1.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec1.log new file mode 100644 index 000000000..325212db9 --- /dev/null +++ b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec1.log @@ -0,0 +1,2 @@ +[WARN]: tec1's vfp_ctrl is disabled, then decode fp inst exception. +[WARN]: tec1 decode fp instruction exception jump from pc=0x00007280 to new pc=0x00005030. diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec2.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec2.log new file mode 100644 index 000000000..a30f06aba --- /dev/null +++ b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec2.log @@ -0,0 +1,2 @@ +[WARN]: tec2's vfp_ctrl is disabled, then decode fp inst exception. +[WARN]: tec2 decode fp instruction exception jump from pc=0x00007280 to new pc=0x00005030. diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec3.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec3.log new file mode 100644 index 000000000..207afcde4 --- /dev/null +++ b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec3.log @@ -0,0 +1,2 @@ +[WARN]: tec3's vfp_ctrl is disabled, then decode fp inst exception. +[WARN]: tec3 decode fp instruction exception jump from pc=0x00007280 to new pc=0x00005030. diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/hts_aiff.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/hts_aiff.log new file mode 100644 index 000000000..e69de29bb diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/tsm.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/tsm.log new file mode 100644 index 000000000..e69de29bb diff --git a/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.cl b/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.cl new file mode 100644 index 000000000..31aa8bec9 --- /dev/null +++ b/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.cl @@ -0,0 +1,35 @@ +#include + +GEN_DMA_DIRECT_EXT2INT(kGlobal, kLsram); +GEN_DMA_DIRECT_INT2EXT(kLsram, kGlobal); + +__kernel void add_kernel(__global float* var_11, __global float* var_20, __global float* var_31, int var_14, int var_33, int var_35, int var_37, int var_3, int var_39, int var_41); + +__kernel void add_kernel(__global float* var_11, __global float* var_20, __global float* var_31, int var_14, int var_33, int var_35, int var_37, int var_3, int var_39, int var_41) { + __lsram float buf[1024]; + __lsram float buf_1[1024]; + __lsram float buf_2[1024]; + int var_0 = get_local_size(0); + int var_1 = get_local_id(0); + int var_4 = (var_3 * var_0); + int var_5 = (var_4 + var_1); + int var_8 = (var_5 * 1024); + int var_9 = var_8; + int var_12 = (var_9 + 1024); + int var_15 = var_14; + int var_16 = min(var_12, var_15); + int var_17 = max(var_16, var_9); + int var_18 = (var_17 - var_9); + int cse_var_1 = (var_18 * 4); + DmaDirect_kGlobal_to_kLsram((int)buf, (int)(var_11 + var_9), cse_var_1, cse_var_1, cse_var_1, cse_var_1); + DmaDirect_kGlobal_to_kLsram((int)buf_1, (int)(var_20 + var_9), cse_var_1, cse_var_1, cse_var_1, cse_var_1); + for (int var_24 = 0; var_24 < (1024 - 0); var_24 += 8) { + float8 var_26 = __vload((__lsram float8*)(buf + var_24), ALL_TRUE_w); + float8 var_28 = __vload((__lsram float8*)(buf_1 + var_24), ALL_TRUE_w); + float8 var_29 = (var_26 + var_28); + __vstore(var_29, (__lsram float8*)(buf_2 + var_24), ALL_TRUE_w); + } + DmaDirect_kLsram_to_kGlobal((int)(var_31 + var_9), (int)buf_2, cse_var_1, cse_var_1, cse_var_1, cse_var_1); + barrier(CLK_LOCAL_MEM_FENCE);return; + barrier(CLK_LOCAL_MEM_FENCE); +} diff --git a/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.log b/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.log new file mode 100644 index 000000000..98c4aff04 --- /dev/null +++ b/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.log @@ -0,0 +1,3 @@ +Current Working Directory: /root/flagtree/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib +Command Line: aipuocc -c -target aipu -mcpu=X2_1204 -I /usr/local/lib/python3.10/dist-packages/tvm/aipu/include -o add_kernel.o add_kernel.cl +Standard Output & Error: diff --git a/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.o b/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.o new file mode 100644 index 0000000000000000000000000000000000000000..06a5c3b0a5857417646b93388c392cd70556b395 GIT binary patch literal 4064 zcmcgvZ){Ul6hCimq4ag~HXw}=WP_QFF`B^;N0{NR9S%EUX3%K(Qb+ej*8OSIRh=IU zWdnolkR^n~7)WOHgCH3AfYD4yQ8UT#&j?@q=Hej4jXzZqH4;+KdGFo6``*xr7+-RG z@9&;-&pG#=-#zbj_m*u>NsH&o69HU!NSCbFx+<6tq7>lmBB!*YtC%f?3Gl(M zcW*iv+Y$S*aVg8gvdc>DY^0UliZJGcF?e`N$sIWaEWg{!${b3(e36o)F|`#dmH5g! z6K@JC@vFh}ce#Ajq10~LZpud&A&v-mCE`8-*C8Gha1ik!jHj+u^EbL(seReS0dU*_)7@0X~O)c4Bc{ z5;2WKr_J@As)BE_quD2&rj5Ek&phYs^~)5^C+W8=v-)ev@3!b~burc=En*cFV+|#E zo)^yq0-iNq=3J-S|Hbhl+v4)ErYapDZ*a2{BR*D|&1?JFx&`-tSl=U$@4NFpwpTch z%kTPF#Rl_ya*@f#D^KcJE{aPxrluO({ib~SQ>(qM4PNGvlv-~8i}PM4mz$W*Ltbme zo<=Knl<4>&NNM|bJZs9w(O#)wpX@ES*)O%(TW zKV0Ao>1WDxoGxQ2hVs}y1_jE&}lsX8WVSIwDOH?{9C@2+W0nu z_lfL@`43?{vT>h`xS2c$ALR$+qv0pZ{iTikSC=~nTyyY`)^PsOICOMjAs-sP5B|}+ zh3fBHLVn1`K$V#f!o#&Xo;(R=xc^5!@_5Ak(|G#M4}g8vGe-QGlzlZ7XH~Cml)Au-a-15h4 z{p*$23jS#QOr?2W<&(jH7w>1(_`Vp=AJpPapG;|dwdeo8#_)=o=_==rS|jEO8uOTGY{UcpgxcNaPpbV%@JC`jJ!)?f!r{1@=m;m&NG#DFPIgAR z!aLPnoxRZ3-X88!6TND;F!V>d!QUH8s{ZC@o?p|KY>RZ+EVOlYbl9d|?NGzt)p@()_8t?e=Lk0{*^aBH6YR``3D?w+7nY)21ilAN+oVI>ymx+P@bA{&O`S{_hr3 zr~T8Vo*qUGdBPv!A3dbckCbp3+WQ`RT33XPWVDPL>DFSSCuML>N4kvU&DeMhBpVid W66UgxaZL94LwdsqNY4FIjDG?FfP`!S literal 0 HcmV?d00001 diff --git a/third_party/aipu/unified_hardware_aipu.h b/third_party/aipu/unified_hardware_aipu.h new file mode 100644 index 000000000..b15421d8f --- /dev/null +++ b/third_party/aipu/unified_hardware_aipu.h @@ -0,0 +1,24 @@ +#ifndef UNIFIED_HARDWARE_AIPU_H +#define UNIFIED_HARDWARE_AIPU_H + +#include + +#include "triton/flagtree/Common/UnifiedHardwareBase.h" + +namespace mlir { + +namespace aipu { + +class UnifiedHardwareAIPU final : public mlir::UnifiedHardware { + + //DMA + std::optional getAllocSpaceForDMATag() const override{ + return 11; + } +}; + +} // namespace aipu + +} // namespace mlir + +#endif // UNIFIED_HARDWARE_AIPU_H \ No newline at end of file From 811f7fdf074771bf280d4af8453a7b28b863c527 Mon Sep 17 00:00:00 2001 From: i3wanna2 <15910307812@163.com> Date: Mon, 16 Jun 2025 17:16:07 +0800 Subject: [PATCH 2/3] add .h --- .../gbuilder/aipu.bin | Bin 92564 -> 0 bytes .../gbuilder/aipubt_gb_add_kernel.py | 63 ------------------ .../gbuilder/aipugb.log | 42 ------------ .../gbuilder/compass_ir.bin | 0 .../gbuilder/compass_ir.txt | 46 ------------- .../gbuilder/op_lib/add_kernel.o | Bin 4064 -> 0 bytes .../runtime/dump/cluster0/core0/tec0.log | 2 - .../runtime/dump/cluster0/core0/tec1.log | 2 - .../runtime/dump/cluster0/core0/tec2.log | 2 - .../runtime/dump/cluster0/core0/tec3.log | 2 - .../runtime/dump/hts_aiff.log | 0 .../runtime/dump/tsm.log | 0 .../gbuilder/op_lib/add_kernel.cl | 35 ---------- .../gbuilder/op_lib/add_kernel.log | 3 - .../gbuilder/op_lib/add_kernel.o | Bin 4064 -> 0 bytes 15 files changed, 197 deletions(-) delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipu.bin delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipubt_gb_add_kernel.py delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipugb.log delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.bin delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.txt delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/op_lib/add_kernel.o delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec0.log delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec1.log delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec2.log delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec3.log delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/hts_aiff.log delete mode 100644 third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/tsm.log delete mode 100644 third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.cl delete mode 100644 third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.log delete mode 100644 third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.o diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipu.bin b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipu.bin deleted file mode 100644 index 6af3938cf016148514948367d2b348b8eaf435df..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 92564 zcmeI3e{dYdmB)K_SMpjb+hfT_GWf?zRth7GSV^`e|Hv^~N!XSR2piwIsHuj+m@5XL9Cp5k=Ki{<^FYKXSH1HlgZL z1(ox?o_@VE-IlC~v*edw` z>Mmhh#e&SkRug7+?)-c9DdA2gx?e7nE+Nat-`I;t41c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0z`la5CI}U1c(3;AOb{y2oM1x zKm>>Y5g-CYfCvx)B0vO)01+SpM1Tko0U|&IhyW2F0-qlOn^?-YbZ?=BpY)$;dV{8Q zJ4@ZB_QmG;{0y zEMx1FvzW|1{+*WPJxnd{)wEC3d794G^ej#LH9cF?1)2_Mx=_<|G(A_-^E7R0dcLNM zG+nIe1)5%{>8mtdqUlmiFVghYn!ZNUi#5Ge(=h%aKAr}uijcsxJ!^4yMjW5+Ag;+1Lf${%#ryIj77ocR2jH>2inTd+P?KesPjs^;fH z%;$xLmNFm0=H>klk2?Jgn@lP1ExNosukY}UV!n8b)bf+YH=X!Hm+n>K@65kEere-a z19C_|cBx+}&)8N~YwH90S+5m~@s@hwGp{o^x5tBDn~Kj=>xY5-%UF{fTmGGk2Xf7m2 z%kP{l{=J!8`*x<7A(`{>D4neTuJWZrYI!|>b)LFi*7e)wcd5o}rsFwX{rY^8?PJT; z^}(oBwWm?l-tDUPb*MTosp|Y;RnK}v)&9e(p8c|_3r?##@Q$hrCsaM>V^z=fs`smT z#i}-!t9pK|s*4&`UA$e@3p!N2FsbURhJ}{vcgZ8FogG&7qL)>D^=VaK^Ny+)PpEpy z$EsfHEwa~twmfXFr^||k&W%6Yr`#Wml&U@KZpH3RsoKZxQS5msRp+z667}Q#X;w;T zS-)Saf3{YCfmVM&tG`gIe~woFT&@0jTK%S0|9q|fBCY;nt^NgC{R_4FuhQx-(dsYN z>R+VQf3;TsHCp|PwfdK6^)J=xU#8X1$H6euzmel%Ff&x<)}|S1aOv>F%Chs(NQn9Bf&K3}ueb(nNl(l|JG6)E3fv!ipgmllYZP`b zw1;aC?k{}M9@~FA`a^rTK7e))w8!>8jsDOc+aLFPK5Wm#o9Les_6oFnM$!K*^dA#C zEd7rQJ@I4oPYc~C{m0SY+hF$}_Rpezzqmj8vgkic*bOG;a|PNxrqJh#(LX5kko3p> z`9ul& z_ThfbGmicn(LW<}Iof?$^xuO1OsrQ;(%%&N+->L|6uMpdhlHMJY!vxr3GKUD@b$E! z|3+c=wh3L1c3%ShhtWTU{)^Fn6#XAS|1tDOyYD#qkDz~A=qBkuj{Xm$e@5tb>7Paa z!|2b7?fmd9!TL?1&mBSkpwQ)L_l1O>cn%&ADEeoF{wL|56`ENq-aP)2rwo6#8}PAHw3>}43CS-lw2$T4zehad=#PB78KJ{y_hr!^`4~*hk51`t;(RKW z{TF&j`iF#`!26wH3GK_r`diT-?|0rdq07`_X?C{R4PE8WTE<_S||f z5AR25VegdwcE%hrR_7>|kh;hT%?F>yYKH%KAlu({pfN1#gzt2N_B=t_d=M|s zG_!AL&x6ir&qKW0_;kS7$QV3+cmY$isXf2{0LS0Q#CUVxr(ygtM*3+jUY<7I2XMSo zINq-Dx9{irc(b6$9~nP$>0u>*CLO=u<9IjV{#4e_zk`;B>$0yp zx=nkY)u8E_KHr$GJb(Ui=bky%Vm8i7*nj`#BN)Fs``>L?4D)r#J7)G|N2~RV4$Jbf zGi=d>nf=xyV*G<%Gf`G-X0hK^`EoO{yw=g{8_mS~ji++Ik6wlOw|={0f3+BNhf6O9 zopkA1(8Dg>2>KCdPp!zuSGSv?Yd&`DuYrGwcYdxt--r0eT=4@V{=}SD{n_2u`R!aR zOr;sLJ>9aJSwA2CuwgcP!`26`W!4=>ORU-Kbz6TO^r4zu&ZiH3Bi|Z1=;ZVLZ-f19 zD;z63vw_FsJkC7m*d?ETKc2GVjT=5?zK$DU1SdpvSlJR+&rnREJf=TkP zWr?#sZCRq$w?$j8Hsx8?7Ngjzs5rA>KJa|&m2kM=hC<6%BjW$cyVGjUE3ukZ2z_jW z-+F$m#41XQ`Ms$b`v1+`GjF^9KR#7r9dPk@=gAVQVvWNm+mUH_@5@35JPG6H6B7-c zRgV4q_tbcM))ZO=hI79;eX7t3lsTI72vn=Oph4B%`9jad^~!s@V?PV=id^w|yk%Pa zA}!uBy*`Z3efg&V;};pD_ImWDap|w&y3L2%57v+7@`dY>pwI_{zM%j!>=^v|T=eJX z@Ui_bG-etq@$X^pKi?Q``Q|+9y?TdV3yh~MU#b0l#WP7fUiKK)kABujg=o+2>!!#* z9N*tH{~9s>9$Y2rdye_SyH^%kd0sI;d3!cO{>lEFU*NYcLjFm;sf=0oJ?`Yo2+t?T zFKOrBA3=T{a`7I11^91p__v$@{wp25XRVrVJjQ3uw<0ayF2H=^@!aFz0__Rle8P5H zfobwF@EP+l%|t$8J!x6rK^*_eGT%_1EB{`bqP)yM?%(qGyxe%$`U?5?k`|BOTlo6> zkZXK+jG-0IcyK)$5<2s6#JCUcQRls4K8?WoxYEhd^kL8~`z>p?u`e-V?8~;?@ZEOq zkDpwBcTcf@H2s&a$J|%Wk0xzB{$2AwjQQ{VSnT)y4c6Zl$bWZxmHdbG_rF~Ezj+R` zR=D#2&?N0Y3+u5v-UV%bT!i`Iw*Ls$Yqx&BT5V4Ytj{Z5>*HH5fY#Q-N-cjYwS2uE z*ULW4hwE{@#ColRIKE9^VncpcKK5w&Bg=o*d~DO==~~7=|LIQ|oTm%4^>nA^uj`w% z^4qlXy5{wr_x>U`9xa-Gi{^i_*0RT{0QP6F2g-iydS;!3Ifd=M?C;gGKRPcTuICl7 zUTt*EUtWHTmOm3({_H-g=11Tgg_b`4a=c~zH@WhMU%M-pgex$*hov$Jg0XPH=! z0{a8X{x0yT`hJA$e}N_LZ-JxY_gjHa=q;aIUjE{>ALi=g{ui!w=4;@E#&99b&)vAc zIE%l34!$7voBI0-pErGSyZZZ5@_7yR`*5C?`Uw0UaO-Dazwg#h!+zhbpYjTQ5dVIG zaqM^(H)a~PY`!ib%kOdJ_ns3?%J_}stM`KshArzNJY(2$&P+7D?fl-jpZP=Gr;4+@ zp7U0Tb-&*qiajaxne?yK@e6|8GDBS_gq{05v63n6e`15G4=vX8U#j|vpkMW0q1i*4 zyZJ_R65h?|(5154h*+OIV)!z$Z7}$1aKW zR^PuKo1#4AO9IB+=Fek6J|rMVvOM2UZnrGRp9GJ8a`_bbHiP`R0QuA6%AXcj{L6{Cy``@5qZs(yC6s*g0P`k7jx@1JBp$ou$a#!Q^GocqrO_J*!z6!+0j3ZRwL!$685XEd;tCl$bqaZ54n)w`DE|cUHR~p*g!Jg z-xsW|s;{c6XB|D!zV2Y|Vs%w@RVcV3ncTJR#v2D>-O+fc#!BAPaYJ`(U?4uoWmR;b zw_{hlFIv^n-@C46?W*ef8yhya?uhjCe5E(i8B4~y`+{Ayq1v@|^|i4z?Om&*vHIxR zP*;6jU92{=wywIpe(kCbFjUvpUdOh^`}V{J)&;k2ZNfJFL91uH>EXTco=)Bf$zcs( zsIIH4Ulm+Y+a3#bu35b*R9o8_s;jB4Ufo{TUQ^v!T^kM6)YWyi*GB6?b!)q>3r3Ui z&PXCQ(6z3*rgmj$eQ$JMBpUDPiX@}GiJn*_(I4*{jP&j8T_#EnRPYghPU6HPMSAQfKzb7WH^>;<~_QeO+b;i1)dwT}g zCkLY)dxX7ab=~^G=s$#5p66m&h zj;;osEphahK_^Nby_wrxdIz^Ja_mvisjD5`%gej$cY|gw`+XqCUG~2L-R81C4*IxD z9|fIu)&C0UjLUusbn7)v{XYV2&35#Ef*y16`4#9HBzjmwps<+NFwK+qUK4ViSG&$@T<5NOYh3l#xazHO)m!7Lx5ia(ja~1S$e-T~cM!WncICHiYi`vtQo3}-v;-=;X=<;p1wZh5-via7?&8@XFK%;}+u9Ir-WuTzf#Tb8?#|`yTQ-Zjc@J}fcQGe;A9I3t@-G@T z-x1lo?RLA(yazq)J-iM+pzVW^uDyL7$-V8}1JT5;NbaO2(wXduL_0epaFgtd^)TnO zssohpg>@dR(Vj?8yq#Z`F?*x&zFer>$z?Gka1R(5bcJRS`#u&Kh;{T2blNvD?GoIb zv@3(L4*g0TZY=Gw?s(s1hFD*xcKNO@xJM53_w;w~jWMXy#?q0vJJJ~)j7A3gBb}WC zQ(ceT6^#!n#=Bzi?p=f2t-E(h!9b_!%()^V3hn<;yz)lp7)J^u6>z_+-*zw>WUVbJ%KHE|F$dJ?4 z1C5pSa33CGIv?n&?GxklXKKH|1$x;|KJ}whyW2F0z`la5CI}U1c(3;AOb{y2oM1xKm>>Y5g-CYfCvx) zB0vO)01+SpMBs}}pa`1JUjyLp4RrMP!FzjsgS(>1UGO4qPYkca7b}2I6KeTpWfC=f zsRs|nKl0!&Gw=hx>&2P!FMmfN1Rv3c-bB17Hqe{w=2txV#=hon`|@({vnY<5m<^i8 z=W)7wp*LM|9zxDy;dxXEDO=WnTK$NvIuYcFyD diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipubt_gb_add_kernel.py b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipubt_gb_add_kernel.py deleted file mode 100644 index 2b8096089..000000000 --- a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipubt_gb_add_kernel.py +++ /dev/null @@ -1,63 +0,0 @@ -import numpy as np -from AIPUBuilder.core import BuilderOpPlugin, register_optype, BuilderParams, Tensor -from AIPUBuilder.plugin_loader import register_plugin, PluginType - - -op_type = register_optype("DSL_add_kernel") - -def _get_addr(base, offset_in_byte): - if offset_in_byte == 0: - return base - - ret = Tensor(base) - ret.memory_offset().set_base_offset(base.memory_offset()) - ret.memory_offset().relative_offset += offset_in_byte - return ret - -@register_plugin(PluginType.Builder, 0) -class DSL_add_kernelPlugin(BuilderOpPlugin): - def get_graph_pattern(self): - return ([("useless", op_type)], []) - - def get_score(self): - return 10 - - def set_target(self, target): - return True - - def check_params(self, nodes): - return True - - def setup(self, sgnode, nodes): - sgnode.attrs["keeping_layout"] = False - return True - - def generate_code_name(self, sgnode, nodes): - return "op_lib/add_kernel.o" - - def generate_descriptor(self, sgnode, nodes): - desc_base = sgnode.attrs["descriptorbase"] - ro = BuilderParams() - return ro - - def generate_params(self, sgnode, nodes): - desc_base = sgnode.attrs["descriptorbase"] - ro = BuilderParams() - ro.append(sgnode.inputs[0]) - ro.append(sgnode.inputs[1]) - ro.append(sgnode.outputs[0]) - value = np.int32(nodes[0].params["var_14"]) - ro.append(int(value.view("int32"))) - value = np.int32(nodes[0].params["var_33"]) - ro.append(int(value.view("int32"))) - value = np.int32(nodes[0].params["var_35"]) - ro.append(int(value.view("int32"))) - value = np.int32(nodes[0].params["var_37"]) - ro.append(int(value.view("int32"))) - value = np.int32(nodes[0].params["var_3"]) - ro.append(int(value.view("int32"))) - value = np.int32(nodes[0].params["var_39"]) - ro.append(int(value.view("int32"))) - value = np.int32(nodes[0].params["var_41"]) - ro.append(int(value.view("int32"))) - return ro diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipugb.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipugb.log deleted file mode 100644 index ff4bae16f..000000000 --- a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/aipugb.log +++ /dev/null @@ -1,42 +0,0 @@ -Current Working Directory: /root/flagtree/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder -Command Line: aipugb compass_ir.txt -w compass_ir.bin --target X2_1204 -Standard Output & Error: -[I] [IRChecker] Start to check IR: /root/flagtree/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.txt -[I] [IRChecker] model_name: unknown -[I] [IRChecker] IRChecker: All IR pass -[I] [tools.cpp : 352] BuildTool version: 6.1.17784. Build for target X2_1204 PID: 42160 -[I] [tools.cpp : 851] global cwd: /tmp/e77881b08dfe9a5bd2c506fa5be9126b7d478215d54cd681637723a3b0f97 -[I] [graph.cpp :1605] loading graph weight: compass_ir.bin size: 0x0 -[I] [actg.cpp : 471] new sgnode with actg: 0 -[I] [datalayout_schedule.cpp:2443] Layout loss: 0 -[I] [datalayout_schedule.cpp:2444] Layout scheduling ... -[I] [datalayout_schedule.cpp:2447] The layout loss for graph unknown: 0 -[I] [datalayout_schedule.cpp:1397] The graph unknown post optimized score:0 -[I] [datalayout_schedule.cpp:1401] layout schedule costs: 0.079781ms -[I] [IRChecker] Start to check IR: -[I] [IRChecker] model_name: cost_model -[I] [IRChecker] IRChecker: All IR pass -[I] [load_balancer.cpp:3032] enable multicore schedule optimization for load balance strategy 1 it may degrade performance on single core targets. -[I] [soms_scheduler.cpp: 195] [EVAL] init time t1: 47 ms -[I] [soms_scheduler.cpp: 201] [EVAL] unsafe check time t2: 10 ms -[I] [soms_scheduler.cpp: 947] get max in loop 0 -[I] [soms_scheduler.cpp: 214] [EVAL] mem assignment time t3: 26 ms -[I] [builder.cpp:1829] [EVAL] duration time of mem allocation: 94 ms -[W] [subgraphop.cpp: 613] empty kernel name, will trying to use the first kernel in elf for node:subgraph_2_dsl_add_kernel -[I] [builder.cpp:1948] The graph DDR Footprint requirement(estimation) of feature maps: -[I] [builder.cpp:1949] Read and Write:51.94KB -[I] [builder2.cpp:1666] memory statistics for this graph (unknown) -[I] [builder.cpp: 635] Text section: 0x00003640 Bytes ( 0.013MB) -[I] [builder.cpp: 635] Weight section: 0x00000000 Bytes ( 0.000MB) -[I] [builder.cpp: 635] Ro section: 0x00000500 Bytes ( 0.001MB) -[I] [builder.cpp: 635] Desc section: 0x00000000 Bytes ( 0.000MB) -[I] [builder.cpp: 635] Stack section: 0x00040400 Bytes ( 0.251MB) -[I] [builder.cpp: 635] Workspace section: 0x00000000 Bytes ( 0.000MB) -[I] [builder.cpp: 635] Activations section: 0x00010fc4 Bytes ( 0.066MB) -[I] [builder.cpp: 635] Total: 0x00054f04 Bytes ( 0.332MB) -[I] [tools.cpp :1480] - compile time: 0.242 s -[I] [tools.cpp :1382] With GM optimization, DDR Footprint stastic(estimation): -[I] [tools.cpp :1389] Read and Write:51.94KB -[I] [tools.cpp :1436] - draw graph time: 0 s -[I] [tools.cpp :1991] remove global cwd: /tmp/e77881b08dfe9a5bd2c506fa5be9126b7d478215d54cd681637723a3b0f97 -Total errors: 0, warnings: 1 diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.bin b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.bin deleted file mode 100644 index e69de29bb..000000000 diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.txt b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.txt deleted file mode 100644 index 241b02f9f..000000000 --- a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/compass_ir.txt +++ /dev/null @@ -1,46 +0,0 @@ - -model_name=unknown -layer_number=3 -data_format=NHWC -precision=float32 -batch_size=1 -input_tensors=[input0,input1] -output_tensors=[temp_var_0] -compat_quantized_model=false - -layer_id=0 -layer_name=0_input -layer_type=Input -layer_bottom=[] -layer_bottom_shape=[] -layer_bottom_type=[] -layer_top=[input0] -layer_top_shape=[[4432]] -layer_top_type=[float32] - -layer_id=1 -layer_name=1_input -layer_type=Input -layer_bottom=[] -layer_bottom_shape=[] -layer_bottom_type=[] -layer_top=[input1] -layer_top_shape=[[4432]] -layer_top_type=[float32] - -layer_id=2 -layer_name=2_dsl_add_kernel -layer_type=DSL_add_kernel -layer_bottom=[input0, input1] -layer_bottom_shape=[[4432], [4432]] -layer_bottom_type=[float32, float32] -layer_top=[temp_var_0] -layer_top_shape=[[4432]] -layer_top_type=[float32] -var_14=4432 -var_33=5 -var_35=1 -var_37=1 -var_3=1 -var_39=0 -var_41=0 diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/op_lib/add_kernel.o b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/gbuilder/op_lib/add_kernel.o deleted file mode 100644 index 06a5c3b0a5857417646b93388c392cd70556b395..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 4064 zcmcgvZ){Ul6hCimq4ag~HXw}=WP_QFF`B^;N0{NR9S%EUX3%K(Qb+ej*8OSIRh=IU zWdnolkR^n~7)WOHgCH3AfYD4yQ8UT#&j?@q=Hej4jXzZqH4;+KdGFo6``*xr7+-RG z@9&;-&pG#=-#zbj_m*u>NsH&o69HU!NSCbFx+<6tq7>lmBB!*YtC%f?3Gl(M zcW*iv+Y$S*aVg8gvdc>DY^0UliZJGcF?e`N$sIWaEWg{!${b3(e36o)F|`#dmH5g! z6K@JC@vFh}ce#Ajq10~LZpud&A&v-mCE`8-*C8Gha1ik!jHj+u^EbL(seReS0dU*_)7@0X~O)c4Bc{ z5;2WKr_J@As)BE_quD2&rj5Ek&phYs^~)5^C+W8=v-)ev@3!b~burc=En*cFV+|#E zo)^yq0-iNq=3J-S|Hbhl+v4)ErYapDZ*a2{BR*D|&1?JFx&`-tSl=U$@4NFpwpTch z%kTPF#Rl_ya*@f#D^KcJE{aPxrluO({ib~SQ>(qM4PNGvlv-~8i}PM4mz$W*Ltbme zo<=Knl<4>&NNM|bJZs9w(O#)wpX@ES*)O%(TW zKV0Ao>1WDxoGxQ2hVs}y1_jE&}lsX8WVSIwDOH?{9C@2+W0nu z_lfL@`43?{vT>h`xS2c$ALR$+qv0pZ{iTikSC=~nTyyY`)^PsOICOMjAs-sP5B|}+ zh3fBHLVn1`K$V#f!o#&Xo;(R=xc^5!@_5Ak(|G#M4}g8vGe-QGlzlZ7XH~Cml)Au-a-15h4 z{p*$23jS#QOr?2W<&(jH7w>1(_`Vp=AJpPapG;|dwdeo8#_)=o=_==rS|jEO8uOTGY{UcpgxcNaPpbV%@JC`jJ!)?f!r{1@=m;m&NG#DFPIgAR z!aLPnoxRZ3-X88!6TND;F!V>d!QUH8s{ZC@o?p|KY>RZ+EVOlYbl9d|?NGzt)p@()_8t?e=Lk0{*^aBH6YR``3D?w+7nY)21ilAN+oVI>ymx+P@bA{&O`S{_hr3 zr~T8Vo*qUGdBPv!A3dbckCbp3+WQ`RT33XPWVDPL>DFSSCuML>N4kvU&DeMhBpVid W66UgxaZL94LwdsqNY4FIjDG?FfP`!S diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec0.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec0.log deleted file mode 100644 index 02e7b0e2c..000000000 --- a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec0.log +++ /dev/null @@ -1,2 +0,0 @@ -[WARN]: tec0's vfp_ctrl is disabled, then decode fp inst exception. -[WARN]: tec0 decode fp instruction exception jump from pc=0x00007280 to new pc=0x00005030. diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec1.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec1.log deleted file mode 100644 index 325212db9..000000000 --- a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec1.log +++ /dev/null @@ -1,2 +0,0 @@ -[WARN]: tec1's vfp_ctrl is disabled, then decode fp inst exception. -[WARN]: tec1 decode fp instruction exception jump from pc=0x00007280 to new pc=0x00005030. diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec2.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec2.log deleted file mode 100644 index a30f06aba..000000000 --- a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec2.log +++ /dev/null @@ -1,2 +0,0 @@ -[WARN]: tec2's vfp_ctrl is disabled, then decode fp inst exception. -[WARN]: tec2 decode fp instruction exception jump from pc=0x00007280 to new pc=0x00005030. diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec3.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec3.log deleted file mode 100644 index 207afcde4..000000000 --- a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/cluster0/core0/tec3.log +++ /dev/null @@ -1,2 +0,0 @@ -[WARN]: tec3's vfp_ctrl is disabled, then decode fp inst exception. -[WARN]: tec3 decode fp instruction exception jump from pc=0x00007280 to new pc=0x00005030. diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/hts_aiff.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/hts_aiff.log deleted file mode 100644 index e69de29bb..000000000 diff --git a/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/tsm.log b/third_party/aipu/python/test/compass_dsl4763d5ba84114530a0c7188c26d9954a/runtime/dump/tsm.log deleted file mode 100644 index e69de29bb..000000000 diff --git a/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.cl b/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.cl deleted file mode 100644 index 31aa8bec9..000000000 --- a/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.cl +++ /dev/null @@ -1,35 +0,0 @@ -#include - -GEN_DMA_DIRECT_EXT2INT(kGlobal, kLsram); -GEN_DMA_DIRECT_INT2EXT(kLsram, kGlobal); - -__kernel void add_kernel(__global float* var_11, __global float* var_20, __global float* var_31, int var_14, int var_33, int var_35, int var_37, int var_3, int var_39, int var_41); - -__kernel void add_kernel(__global float* var_11, __global float* var_20, __global float* var_31, int var_14, int var_33, int var_35, int var_37, int var_3, int var_39, int var_41) { - __lsram float buf[1024]; - __lsram float buf_1[1024]; - __lsram float buf_2[1024]; - int var_0 = get_local_size(0); - int var_1 = get_local_id(0); - int var_4 = (var_3 * var_0); - int var_5 = (var_4 + var_1); - int var_8 = (var_5 * 1024); - int var_9 = var_8; - int var_12 = (var_9 + 1024); - int var_15 = var_14; - int var_16 = min(var_12, var_15); - int var_17 = max(var_16, var_9); - int var_18 = (var_17 - var_9); - int cse_var_1 = (var_18 * 4); - DmaDirect_kGlobal_to_kLsram((int)buf, (int)(var_11 + var_9), cse_var_1, cse_var_1, cse_var_1, cse_var_1); - DmaDirect_kGlobal_to_kLsram((int)buf_1, (int)(var_20 + var_9), cse_var_1, cse_var_1, cse_var_1, cse_var_1); - for (int var_24 = 0; var_24 < (1024 - 0); var_24 += 8) { - float8 var_26 = __vload((__lsram float8*)(buf + var_24), ALL_TRUE_w); - float8 var_28 = __vload((__lsram float8*)(buf_1 + var_24), ALL_TRUE_w); - float8 var_29 = (var_26 + var_28); - __vstore(var_29, (__lsram float8*)(buf_2 + var_24), ALL_TRUE_w); - } - DmaDirect_kLsram_to_kGlobal((int)(var_31 + var_9), (int)buf_2, cse_var_1, cse_var_1, cse_var_1, cse_var_1); - barrier(CLK_LOCAL_MEM_FENCE);return; - barrier(CLK_LOCAL_MEM_FENCE); -} diff --git a/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.log b/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.log deleted file mode 100644 index 98c4aff04..000000000 --- a/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.log +++ /dev/null @@ -1,3 +0,0 @@ -Current Working Directory: /root/flagtree/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib -Command Line: aipuocc -c -target aipu -mcpu=X2_1204 -I /usr/local/lib/python3.10/dist-packages/tvm/aipu/include -o add_kernel.o add_kernel.cl -Standard Output & Error: diff --git a/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.o b/third_party/aipu/python/test/compass_dsl_add_kernel_8bb556285cb04e918905aca7ffc0cb65/gbuilder/op_lib/add_kernel.o deleted file mode 100644 index 06a5c3b0a5857417646b93388c392cd70556b395..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 4064 zcmcgvZ){Ul6hCimq4ag~HXw}=WP_QFF`B^;N0{NR9S%EUX3%K(Qb+ej*8OSIRh=IU zWdnolkR^n~7)WOHgCH3AfYD4yQ8UT#&j?@q=Hej4jXzZqH4;+KdGFo6``*xr7+-RG z@9&;-&pG#=-#zbj_m*u>NsH&o69HU!NSCbFx+<6tq7>lmBB!*YtC%f?3Gl(M zcW*iv+Y$S*aVg8gvdc>DY^0UliZJGcF?e`N$sIWaEWg{!${b3(e36o)F|`#dmH5g! z6K@JC@vFh}ce#Ajq10~LZpud&A&v-mCE`8-*C8Gha1ik!jHj+u^EbL(seReS0dU*_)7@0X~O)c4Bc{ z5;2WKr_J@As)BE_quD2&rj5Ek&phYs^~)5^C+W8=v-)ev@3!b~burc=En*cFV+|#E zo)^yq0-iNq=3J-S|Hbhl+v4)ErYapDZ*a2{BR*D|&1?JFx&`-tSl=U$@4NFpwpTch z%kTPF#Rl_ya*@f#D^KcJE{aPxrluO({ib~SQ>(qM4PNGvlv-~8i}PM4mz$W*Ltbme zo<=Knl<4>&NNM|bJZs9w(O#)wpX@ES*)O%(TW zKV0Ao>1WDxoGxQ2hVs}y1_jE&}lsX8WVSIwDOH?{9C@2+W0nu z_lfL@`43?{vT>h`xS2c$ALR$+qv0pZ{iTikSC=~nTyyY`)^PsOICOMjAs-sP5B|}+ zh3fBHLVn1`K$V#f!o#&Xo;(R=xc^5!@_5Ak(|G#M4}g8vGe-QGlzlZ7XH~Cml)Au-a-15h4 z{p*$23jS#QOr?2W<&(jH7w>1(_`Vp=AJpPapG;|dwdeo8#_)=o=_==rS|jEO8uOTGY{UcpgxcNaPpbV%@JC`jJ!)?f!r{1@=m;m&NG#DFPIgAR z!aLPnoxRZ3-X88!6TND;F!V>d!QUH8s{ZC@o?p|KY>RZ+EVOlYbl9d|?NGzt)p@()_8t?e=Lk0{*^aBH6YR``3D?w+7nY)21ilAN+oVI>ymx+P@bA{&O`S{_hr3 zr~T8Vo*qUGdBPv!A3dbckCbp3+WQ`RT33XPWVDPL>DFSSCuML>N4kvU&DeMhBpVid W66UgxaZL94LwdsqNY4FIjDG?FfP`!S From 1e73080296f0e56f80c732b2c37ae822a735e947 Mon Sep 17 00:00:00 2001 From: i3wanna2 <15910307812@163.com> Date: Thu, 19 Jun 2025 17:54:08 +0800 Subject: [PATCH 3/3] add uh, needed to fix mod.set_attr(uh_aipu_ptr, ptr_attr) --- include/CMakeLists.txt | 3 +-- include/flagtree/CMakeLists.txt | 3 --- include/flagtree/Common/UnifiedHardwareBase.h | 4 ++-- .../triton/Dialect/Triton/IR/TritonAttrDefs.td | 1 + third_party/aipu/backend/compiler.py | 17 +++++++++++++++++ third_party/aipu/triton_aipu.cc | 8 ++++++++ third_party/aipu/unified_hardware_aipu.h | 8 +++----- 7 files changed, 32 insertions(+), 12 deletions(-) delete mode 100644 include/flagtree/CMakeLists.txt diff --git a/include/CMakeLists.txt b/include/CMakeLists.txt index 93a0c83ea..72181b98f 100644 --- a/include/CMakeLists.txt +++ b/include/CMakeLists.txt @@ -1,2 +1 @@ -add_subdirectory(triton) -add_subdirectory(flagtree) +add_subdirectory(triton) \ No newline at end of file diff --git a/include/flagtree/CMakeLists.txt b/include/flagtree/CMakeLists.txt deleted file mode 100644 index 27c703b3c..000000000 --- a/include/flagtree/CMakeLists.txt +++ /dev/null @@ -1,3 +0,0 @@ -add_subdirectory(Conversion) -add_subdirectory(Dialect) -add_subdirectory(Target) diff --git a/include/flagtree/Common/UnifiedHardwareBase.h b/include/flagtree/Common/UnifiedHardwareBase.h index dc165c728..03e8aba38 100644 --- a/include/flagtree/Common/UnifiedHardwareBase.h +++ b/include/flagtree/Common/UnifiedHardwareBase.h @@ -4,7 +4,7 @@ #include namespace mlir { - +namespace flagtree { //this is the unified hardware abstraction for hardware //to determined if these abstraction is specified, using std::optional is needed //using in passes: if(uh_flagtree->xxx()){...} @@ -21,7 +21,7 @@ class UnifiedHardware{ }; - +} // namespace flagtree } // namespace mlir #endif // UNIFIED_HARDWARE_BASE_H diff --git a/include/triton/Dialect/Triton/IR/TritonAttrDefs.td b/include/triton/Dialect/Triton/IR/TritonAttrDefs.td index 571d2b55b..898ca1559 100644 --- a/include/triton/Dialect/Triton/IR/TritonAttrDefs.td +++ b/include/triton/Dialect/Triton/IR/TritonAttrDefs.td @@ -4,6 +4,7 @@ include "mlir/IR/EnumAttr.td" include "mlir/IR/AttrTypeBase.td" + // Attributes for LoadOp and StoreOp def TT_CacheModifierAttr : I32EnumAttr< "CacheModifier", "", diff --git a/third_party/aipu/backend/compiler.py b/third_party/aipu/backend/compiler.py index f3554128a..b1e255480 100644 --- a/third_party/aipu/backend/compiler.py +++ b/third_party/aipu/backend/compiler.py @@ -7,6 +7,8 @@ from mlir.passmanager import PassManager from mlir.ir import Context, Module +from triton._C.libtriton.aipu import UnifiedHardwareAIPU + from dataclasses import dataclass import functools import hashlib @@ -14,6 +16,7 @@ from types import ModuleType + def min_dot_size(target: GPUTarget): return lambda lhsType, rhsType: (1, 1, 1) @@ -80,6 +83,7 @@ def load_dialects(self, ctx): @staticmethod def make_ttir(mod, metadata, opt): + pm = ir.pass_manager(mod.context) pm.enable_debug() passes.common.add_inliner(pm) @@ -96,6 +100,19 @@ def make_ttir(mod, metadata, opt): @staticmethod def make_linalg(mod, metadata, opt): + + import ctypes + uh_aipu = UnifiedHardwareAIPU() + obj_ptr = id(uh_aipu) + uh_aipu_ptr = ctypes.cast(obj_ptr, ctypes.c_void_p) + + + from mlir.ir import IntegerAttr, IntegerType + ctx = Context() + ptr_type = IntegerType.get_unsigned(64, ctx) + ptr_attr = IntegerAttr.get(ptr_type, uh_aipu_ptr.value) + # mod.set_attr("uh_aipu_ptr", ptr_attr) #报错 + pm = ir.pass_manager(mod.context) pm.enable_debug() # add pass here diff --git a/third_party/aipu/triton_aipu.cc b/third_party/aipu/triton_aipu.cc index b617232df..1df4d481f 100644 --- a/third_party/aipu/triton_aipu.cc +++ b/third_party/aipu/triton_aipu.cc @@ -22,6 +22,8 @@ #include "triton-shared/Conversion/TritonToLinalg/TritonToLinalg.h" #include "triton-shared/Conversion/TritonToLinalgExperimental/TritonToLinalgExperimental.h" +#include "unified_hardware_aipu.h" + #include #include #include @@ -69,6 +71,7 @@ void init_triton_aipu(py::module &&m) { init_triton_aipu_common(m.def_submodule("common")); auto passes = m.def_submodule("passes"); init_triton_aipu_passes_convert(passes.def_submodule("convert")); + // load dialects m.def("load_dialects", [](mlir::MLIRContext &context) { using namespace mlir; @@ -88,4 +91,9 @@ void init_triton_aipu(py::module &&m) { context.loadAllAvailableDialects(); }); // register passes here + + //flagtree + py::class_(m, "UnifiedHardwareAIPU") + .def(py::init<>()); + } diff --git a/third_party/aipu/unified_hardware_aipu.h b/third_party/aipu/unified_hardware_aipu.h index b15421d8f..f8f272ecc 100644 --- a/third_party/aipu/unified_hardware_aipu.h +++ b/third_party/aipu/unified_hardware_aipu.h @@ -3,22 +3,20 @@ #include -#include "triton/flagtree/Common/UnifiedHardwareBase.h" +#include "flagtree/Common/UnifiedHardwareBase.h" namespace mlir { - namespace aipu { -class UnifiedHardwareAIPU final : public mlir::UnifiedHardware { +class UnifiedHardwareAIPU final : public mlir::flagtree::UnifiedHardware { //DMA std::optional getAllocSpaceForDMATag() const override{ - return 11; + return std::optional(11); } }; } // namespace aipu - } // namespace mlir #endif // UNIFIED_HARDWARE_AIPU_H \ No newline at end of file