From 0ca14d740fcf76579d17f5b573f3f003f04592bc Mon Sep 17 00:00:00 2001 From: Andrzej Janik Date: Mon, 13 Dec 2021 22:25:26 +0100 Subject: [PATCH] Better reporting of unrecognized tokens --- hip_runtime-sys/build.rs | 18 +- hip_runtime-sys/lib/amdhip64.def | 308 +++++++++++++++++++++++++++++++ hip_runtime-sys/lib/amdhip64.lib | Bin 0 -> 71480 bytes ptx/src/lib.rs | 112 +++++++++++ ptx/src/ptx.lalrpop | 11 +- zluda_dump/src/trace.rs | 39 ++-- 6 files changed, 461 insertions(+), 27 deletions(-) create mode 100644 hip_runtime-sys/lib/amdhip64.def create mode 100644 hip_runtime-sys/lib/amdhip64.lib diff --git a/hip_runtime-sys/build.rs b/hip_runtime-sys/build.rs index 9c3d9b4..3bc1250 100644 --- a/hip_runtime-sys/build.rs +++ b/hip_runtime-sys/build.rs @@ -1,9 +1,21 @@ use std::env::VarError; +use std::{env, path::PathBuf}; fn main() -> Result<(), VarError> { println!("cargo:rustc-link-lib=dylib=amdhip64"); - //println!("cargo:rustc-link-search=native=/opt/rocm/lib/"); - println!("cargo:rustc-link-search=native=/home/ubuntu/hipamd/build/lib"); - println!("cargo:rustc-link-search=native=/home/vosen/hipamd/build/lib"); + if cfg!(windows) { + let env = env::var("CARGO_CFG_TARGET_ENV")?; + if env == "msvc" { + let mut path = PathBuf::from(env::var("CARGO_MANIFEST_DIR")?); + path.push("lib"); + println!("cargo:rustc-link-search=native={}", path.display()); + } else { + println!("cargo:rustc-link-search=native=C:\\Windows\\System32"); + }; + } else { + //println!("cargo:rustc-link-search=native=/opt/rocm/lib/"); + println!("cargo:rustc-link-search=native=/home/ubuntu/hipamd/build/lib"); + println!("cargo:rustc-link-search=native=/home/vosen/hipamd/build/lib"); + } Ok(()) } diff --git a/hip_runtime-sys/lib/amdhip64.def b/hip_runtime-sys/lib/amdhip64.def new file mode 100644 index 0000000..515d5c6 --- /dev/null +++ b/hip_runtime-sys/lib/amdhip64.def @@ -0,0 +1,308 @@ +LIBRARY AMDHIP64 +EXPORTS +__hipPopCallConfiguration +__hipPushCallConfiguration +__hipRegisterFatBinary +__hipRegisterFunction +__hipRegisterManagedVar +__hipRegisterSurface +__hipRegisterTexture +__hipRegisterVar +__hipUnregisterFatBinary +hipApiName +hipArray3DCreate +hipArrayCreate +hipArrayDestroy +hipBindTexture +hipBindTexture2D +hipBindTextureToArray +hipBindTextureToMipmappedArray +hipChooseDevice +hipConfigureCall +hipCreateChannelDesc +hipCreateSurfaceObject +hipCreateTextureObject +hipCtxCreate +hipCtxDestroy +hipCtxDisablePeerAccess +hipCtxEnablePeerAccess +hipCtxGetApiVersion +hipCtxGetCacheConfig +hipCtxGetCurrent +hipCtxGetDevice +hipCtxGetFlags +hipCtxGetSharedMemConfig +hipCtxPopCurrent +hipCtxPushCurrent +hipCtxSetCacheConfig +hipCtxSetCurrent +hipCtxSetSharedMemConfig +hipCtxSynchronize +hipDestroyExternalMemory +hipDestroyExternalSemaphore +hipDestroySurfaceObject +hipDestroyTextureObject +hipDeviceCanAccessPeer +hipDeviceComputeCapability +hipDeviceDisablePeerAccess +hipDeviceEnablePeerAccess +hipDeviceGet +hipDeviceGetAttribute +hipDeviceGetByPCIBusId +hipDeviceGetCacheConfig +hipDeviceGetLimit +hipDeviceGetName +hipDeviceGetP2PAttribute +hipDeviceGetPCIBusId +hipDeviceGetSharedMemConfig +hipDeviceGetStreamPriorityRange +hipDevicePrimaryCtxGetState +hipDevicePrimaryCtxRelease +hipDevicePrimaryCtxReset +hipDevicePrimaryCtxRetain +hipDevicePrimaryCtxSetFlags +hipDeviceReset +hipDeviceSetCacheConfig +hipDeviceSetSharedMemConfig +hipDeviceSynchronize +hipDeviceTotalMem +hipDriverGetVersion +hipDrvMemcpy2DUnaligned +hipDrvMemcpy3D +hipDrvMemcpy3DAsync +hipEnableActivityCallback +hipEventCreate +hipEventCreateWithFlags +hipEventDestroy +hipEventElapsedTime +hipEventQuery +hipEventRecord +hipEventSynchronize +hipExtGetLinkTypeAndHopCount +hipExtLaunchKernel +hipExtLaunchMultiKernelMultiDevice +hipExtMallocWithFlags +hipExtModuleLaunchKernel +hipExtStreamCreateWithCUMask +hipExtStreamGetCUMask +hipExternalMemoryGetMappedBuffer +hipFree +hipFreeArray +hipFreeHost +hipFreeMipmappedArray +hipFuncGetAttribute +hipFuncGetAttributes +hipFuncSetAttribute +hipFuncSetCacheConfig +hipFuncSetSharedMemConfig +hipGLGetDevices +hipGetChannelDesc +hipGetCmdName +hipGetDevice +hipGetDeviceCount +hipGetDeviceFlags +hipGetDeviceProperties +hipGetErrorName +hipGetErrorString +hipGetLastError +hipGetMipmappedArrayLevel +hipGetSymbolAddress +hipGetSymbolSize +hipGetTextureAlignmentOffset +hipGetTextureObjectResourceDesc +hipGetTextureObjectResourceViewDesc +hipGetTextureObjectTextureDesc +hipGetTextureReference +hipGraphAddDependencies +hipGraphAddEmptyNode +hipGraphAddKernelNode +hipGraphAddMemcpyNode +hipGraphAddMemcpyNode1D +hipGraphAddMemsetNode +hipGraphCreate +hipGraphDestroy +hipGraphExecDestroy +hipGraphExecKernelNodeSetParams +hipGraphGetNodes +hipGraphGetRootNodes +hipGraphInstantiate +hipGraphKernelNodeGetParams +hipGraphKernelNodeSetParams +hipGraphLaunch +hipGraphMemcpyNodeGetParams +hipGraphMemcpyNodeSetParams +hipGraphMemsetNodeGetParams +hipGraphMemsetNodeSetParams +hipGraphicsGLRegisterBuffer +hipGraphicsMapResources +hipGraphicsResourceGetMappedPointer +hipGraphicsUnmapResources +hipGraphicsUnregisterResource +hipHccModuleLaunchKernel +hipHostAlloc +hipHostFree +hipHostGetDevicePointer +hipHostGetFlags +hipHostMalloc +hipHostRegister +hipHostUnregister +hipImportExternalMemory +hipImportExternalSemaphore +hipInit +hipInitActivityCallback +hipIpcCloseMemHandle +hipIpcGetEventHandle +hipIpcGetMemHandle +hipIpcOpenEventHandle +hipIpcOpenMemHandle +hipKernelNameRef +hipLaunchByPtr +hipLaunchCooperativeKernel +hipLaunchCooperativeKernelMultiDevice +hipLaunchKernel +hipMalloc +hipMalloc3D +hipMalloc3DArray +hipMallocArray +hipMallocHost +hipMallocManaged +hipMallocMipmappedArray +hipMallocPitch +hipMemAdvise +hipMemAllocHost +hipMemAllocPitch +hipMemGetAddressRange +hipMemGetInfo +hipMemPrefetchAsync +hipMemPtrGetInfo +hipMemRangeGetAttribute +hipMemRangeGetAttributes +hipMemcpy +hipMemcpy2D +hipMemcpy2DAsync +hipMemcpy2DFromArray +hipMemcpy2DFromArrayAsync +hipMemcpy2DToArray +hipMemcpy2DToArrayAsync +hipMemcpy3D +hipMemcpy3DAsync +hipMemcpyAsync +hipMemcpyAtoH +hipMemcpyDtoD +hipMemcpyDtoDAsync +hipMemcpyDtoH +hipMemcpyDtoHAsync +hipMemcpyFromArray +hipMemcpyFromSymbol +hipMemcpyFromSymbolAsync +hipMemcpyHtoA +hipMemcpyHtoD +hipMemcpyHtoDAsync +hipMemcpyParam2D +hipMemcpyParam2DAsync +hipMemcpyPeer +hipMemcpyPeerAsync +hipMemcpyToArray +hipMemcpyToSymbol +hipMemcpyToSymbolAsync +hipMemcpyWithStream +hipMemset +hipMemset2D +hipMemset2DAsync +hipMemset3D +hipMemset3DAsync +hipMemsetAsync +hipMemsetD16 +hipMemsetD16Async +hipMemsetD32 +hipMemsetD32Async +hipMemsetD8 +hipMemsetD8Async +hipMipmappedArrayCreate +hipMipmappedArrayDestroy +hipMipmappedArrayGetLevel +hipModuleGetFunction +hipModuleGetGlobal +hipModuleGetTexRef +hipModuleLaunchKernel +hipModuleLaunchKernelExt +hipModuleLoad +hipModuleLoadData +hipModuleLoadDataEx +hipModuleOccupancyMaxActiveBlocksPerMultiprocessor +hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags +hipModuleOccupancyMaxPotentialBlockSize +hipModuleOccupancyMaxPotentialBlockSizeWithFlags +hipModuleUnload +hipOccupancyMaxActiveBlocksPerMultiprocessor +hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags +hipOccupancyMaxPotentialBlockSize +hipPeekAtLastError +hipPointerGetAttributes +hipProfilerStart +hipProfilerStop +hipRegisterActivityCallback +hipRegisterApiCallback +hipRemoveActivityCallback +hipRemoveApiCallback +hipRuntimeGetVersion +hipSetDevice +hipSetDeviceFlags +hipSetupArgument +hipSignalExternalSemaphoresAsync +hipStreamAddCallback +hipStreamAttachMemAsync +hipStreamBeginCapture +hipStreamCreate +hipStreamCreateWithFlags +hipStreamCreateWithPriority +hipStreamDestroy +hipStreamEndCapture +hipStreamGetFlags +hipStreamGetPriority +hipStreamIsCapturing +hipStreamQuery +hipStreamSynchronize +hipStreamWaitEvent +hipTexObjectCreate +hipTexObjectDestroy +hipTexObjectGetResourceDesc +hipTexObjectGetResourceViewDesc +hipTexObjectGetTextureDesc +hipTexRefGetAddress +hipTexRefGetAddressMode +hipTexRefGetArray +hipTexRefGetBorderColor +hipTexRefGetFilterMode +hipTexRefGetFlags +hipTexRefGetFormat +hipTexRefGetMaxAnisotropy +hipTexRefGetMipmapFilterMode +hipTexRefGetMipmapLevelBias +hipTexRefGetMipmapLevelClamp +hipTexRefGetMipmappedArray +hipTexRefSetAddress +hipTexRefSetAddress2D +hipTexRefSetAddressMode +hipTexRefSetArray +hipTexRefSetBorderColor +hipTexRefSetFilterMode +hipTexRefSetFlags +hipTexRefSetFormat +hipTexRefSetMaxAnisotropy +hipTexRefSetMipmapFilterMode +hipTexRefSetMipmapLevelBias +hipTexRefSetMipmapLevelClamp +hipTexRefSetMipmappedArray +hipUnbindTexture +hipWaitExternalSemaphoresAsync +hiprtcAddNameExpression +hiprtcCompileProgram +hiprtcCreateProgram +hiprtcDestroyProgram +hiprtcGetCode +hiprtcGetCodeSize +hiprtcGetErrorString +hiprtcGetLoweredName +hiprtcGetProgramLog +hiprtcGetProgramLogSize diff --git a/hip_runtime-sys/lib/amdhip64.lib b/hip_runtime-sys/lib/amdhip64.lib new file mode 100644 index 0000000000000000000000000000000000000000..1d7df961e35d5ae012398e42d2e2b286facdfc58 GIT binary patch literal 71480 zcmeHQf4o&iwf}BJRK#D1W@biaW@ZEzMD+Q3TrLRI1@8kipRX4j?%i-t-Fwb?=blTK zp09d6AFm%~RAgmlWK?EkW~N0(WM)QWW=3RXMn-04W@Kij@3UrR&#d{ehkfoj_}9Ci z&w?{+)|&6k?3vl~Ypwmlb4vBz{#TBF#bo<$+RW)^&FGjpZR!l0KUD>%ojK#ond0|~ z@xYkXz&j=a@0tx0bU#3J^@WOV1&BT|Q_-~m(Wc3Y)&qi`28ce{P_zdi`qTqBVe^2LYnZuT!)gAo|=nif#mmuA8jra)9XbV->9d z1nmZhu3xNZ8$k4hvlU$n5Z!Q^qN@O+FHTZ)89;R77)5Ud1nmKczSN^=3qW+!6h#{W zqMIiwx&#pP1VHrVnxgvvqOUAgv;`pAa+DjERUk+@WsAwG^=qZ5c8x=*n0is(kP_zRe`sQp!n*pNRj#qRkK=iHg zibeoI4*^8CzgE#sfau#jif#dj?wG0QdVuIVrzzS55N$tE(FTC%yT>ZJ3?SMuUeQ|t zK@R~$-&>|=J3w^jEJd3DqVG>s^fo}y69Ca&HAOoCq91fC+6oZeJyp?0far&l6s-dU z?FESLS)phbK=h;6D%uVZ?d(u=B|!A!@ru>}f_4K$yLuGe4iNpML(!E0(Y@mpjR1lk z28e!op`z^o(a&Znx(*=v`N@hd1q3|~5Z%|OXeU7Qi^Ynz0z~)EP;?bQ^vel~MgTz% z07Scc6m0{Des#K{D*>VhCMbF{AZRZ@^y?Lhb^t_s&QWwNK=hjv6}=M>^dvy^V85bW z0MT!|6>SBG9y&$QdO*XWNQjQ@QFJSkV{cu)eBq*=^Oh`IK6l>It|be5mYug`dFP_J z3l{dwoPp)b7oXqV?MvV-nTQlvw&485uUrFYhncdaAXq(!fAhwX&+bdbnO{ zl&cjiUtS)pEsw>M=I~H|Ts23azg%0At}YKX()zij#+-7cRA0*>Ly=r(xYFz9QmdFa zrm?70DXmWXE-2MQFjg$nS~^@`Rq9PcDkh4_FH1)n!}XxNiDL3O*^%u2a_#&|JyuKk z(f#FGXRW-rG?;?@l+TnH^?GUT^toO2wA7Fi7Q>cvq*TzBnVSwZ>eaOZjFQU2T53+Y z(w9|1f7uF#WIE#mN? zHad+F^Sn)f$7j!$3fIsfOQJkfS~-yRq-njgw>KRc%49u0*P2%eDtTZ={jdE2V)&>0q@khKF!Op5D@QuvF`>)&-WMYT{gN zyC3$KGfkWCYS}%}XxV(GMI%sGsUo`;buy{svt6C)U~RaOc9m+SmF0nQL)DAW^P(JG zzw3&OlCB@+=z2X^t9mhx=KNIjJlUBj;Kz418ujwZVfh-(m{-ujM{#uIa|P+V+4gtAaBf`DN9H7_SY6<`~I(c}-fMpElSV-{UdqHH*^0-rCxZ zx#w3(1Lf6~RCf6g-`AQh=Ac@5I-7wqX6Mk_ikQ54Je{H(;ksfKxPaECk${Uugt)-onz6I4tgw+v8_puu z;aHM`ec42Zq*RDYn4YA;1M#|pYaz+;Q_`cMmXsV{vZr3HrS(Qx6*en6V{@H(^?J3= z%3{Kv)Y3-1ToH@S`Kgps94^^i8j^AJWN1LsVJ>Ai(e8AO9I&VZS$i5XhHEZeJGins z(An2lSF3TFg;_jbdZ}7!X;H!8a%zRrNw3zyw9+_l)hab(nxE!H2K43D#FBKVI$ZBf zv&|5bL?C{jsa#M_-?*=ov`lVls(^V3@#sJrOVU+oJ*~*?B2yg0HahiEt-rIcZ*E#k zD}8CDH*1C=zNa;Bu+~_+xY{QM%S^>EJ$d=&aupcY(~#3BSA%gQ8q>ryD$AD9na1W* zEOsbu1aLC0jbmm;A^{b&WwJR9^((5)ENn~78%cY!m`sxMd}&i%OVdVAsa_hCqaiHWxgz_QPT`hZb`K&YQ>T@Oj~|oWvEfAG|FntByz-gwzSFTe6Jj#xIAsY7L2xu zv$f?+Mpw8B+A^lW`1uRPbHfHM}6U7ni4P(r4I8+r-)0<=&zB-PxIj znL3H58H)3?7nN$+kZPJux!2}lFx zQx;_^&K!lfabkZ9k!%Q&1G_oqa_y()~fY}JxwS1AU{XTY?FG5Cd%km zBZ4T4P%B4%3l$Q>PxI0lDmX zwY1_V?s8oXlNME9O~@ObmR5soX^^HBBx*CXoE%%{tnF!t2hW3!OjngA!lg!eO)Ar( zRHeKhwJ%*~d7P6BmvT~cv!KFS%Jd7jPsp&`#L>SUz3FrHN-`61R+EzLNCjo6mr_+p z4H0tC)G4zO${4PkHJNAPHs8_eDK~n>yHsXC28~QQ*x9$HEO$apz>25APnRMcG+|H5 zU@0*ljh4)ExuN%e7o{GTr#wGxEUc`mit-9c5#ka(^>kI*=8;a*&A!b3xm{a$DGBgz-3n+R|B+ieFBQuJECh zetWSzPeR1c%~0CVyKLObjxT7v zlU>oHC+abM`S}CYl_fb?Wzve}%Pvbt)J#~*DwZQF=b~nQDa9>rL}T8Foa}igG8EzJ zRZDWxCxs~#keFL)luU(?l;Gq@&KvP%=Z(11=k@ju*GiS%wTnt4V!f2k8L0NYVW=mq zi#d0#UZv|h_|sL_p}_f#7x?SBo3!e-mpK2kugsZ>p7bUWsi z;3sbs9Gvqj16CVRVJ*dz-RE;GF5G|U@Jr!lm4 zt;{bHgxE}Hu(~FA|6oZt#?^GOL~@L2T+h3Hb5)98U{5 zdt^(~#&E5(zIvF>p;T&&VNxuwu9Rqy7AMqtBWg&$W8(rQ9dX5}v#*ao1)u54HyWj0 z+ITi_G98=mXw6Aimn&VRnz&j{6(w{muA^b*08-XH8R1;hId#%f+(ZOZr|9g8S4dUI zIojETLZf92M{ZuF&nq+IInsO=RhCY3d`X=(od??ASJ%~C`GrF=8|n~OCeE2o`uYNl zlyq*uNG89^qVvF+PMWVNmF51Vkv1$XyDS}%N2TWRyDVQetNs465zctN^!(I2r_~T{ zTvS`1y>eWZf{qWS!Vg}Z1b7tJU6QUcTT)8N=H*H+Qb((D9%)h${D2l2K-dU^<4ew| z*89?WS9L&6z;ybY?`xe~9%!U>FH?*i(8(^IYYb=nH2i9Puq4`<%kopDS!Ja>RBhC& zHM_)iOh8jCSy-0kQodC9q=xnz?OuX?6i8S{~N0>^7^_bK9&f;k4OMWM1xd z!?X%+?H?H3^D8TDt1X+MF~qAVe?}FH)@jrmy`6pHcG$d;8eLLS_B=s8*CNY>@&Mh_ zT3s)R&8&LEWm@vW#>x@5cwAb&32kwQ$>r%dOWp@;dmnK6`++Oo4~*LYG&TSa zZUFi|0NninVD{y}4VMF_d=R+egTUAi0ksbSdp-nQ^kHDfhk;jK0c^ekIQ}ERx{m-) zeFS(tS?ym5bbl1s@=;*wRlvrpfMY)ftp6A=ZXx>q`#%m`cr|eA)xgY8 z0M~v3n7j#CzX^DH6VUi1u;-J&icbOeeG0hn)4*My2F|?(*m?~x>odTn&j80?3v9R+ z824FV&1ZoJKMTBWGq8O#aL(s|8$Sn3z7DwjI$-SQfi<58c7GmNd_A!3df@CY0M~v2 zIPC`DsvCewUj#1uA~5Dg;H@_Tdu{}Jz65Of5-{Z^VB<}|#G8RjZU&yX8K`|3xbMrr z;;#T(z5<-Q1=z3!82eRV%In@uobv`D=_Knz~x^DCT;`P zZ3CX#22{QQ?EVID!L7iKTY=f%1U7#YIQ}-^(%XRX-vUOy1w8aE;I+2{J8uViz75>+ zZD8gd!1Z?kr+o+5^c~>D?ZAfZz_H&2F8eMpeh2W@9l%37fMwqUwto+pbtkatPGI8q zfwz4hc;fp&?Ji*FT|oB_fUQ3Orrr%~yc?MGLtxzxfxSNjR@?*Zx(9gekAUqz0y=gA zSMCJH{}@>FV_^4>fu3E!?Yn@Ep8!|>1Q>rWFmf;O@V&r=KLxh`6qxlh;JTjyC;uF{ z^yk3iKL`5m19sjAEdB+s^%ua5`+=+O2PXUy82Kgez%PNG-N3fp!0EpNuKX1+;Q`>y z4*+`~09O1O*zs%NoISv`dw>&v1HAJ$z>~iL`X2;#JqUFF7TEe*;FO1e^$!70KLqsu z4%qoSV9vw9=7)h3eh;ksJ@EMNffbJccRm8l{R6P&4?xGGz@|rmNq+<``6KYeAA!na zz`c(Fulf^k%b$Smy}+Kmz`&n@C;treKMp+hII!+7z_=%XHBSH!KLK3ySKyAn0&|`O zZh8`!@;6}P-++lv0T(|7Jp2@J(cghxe+ORmG_dVyV8%1R)z3gUARy-hoKB1doqRP&+WeBL^l}H7>!faS0Y5lwj*Y31%Fe;F^OI z9Ct{9w;z%K;}g7Te1g5>6I2dOu;~kJTk$`qY~VERDuN)6Wlg2!P!SA*nD(?6DK8DKPkaelM)OZlir8o^Q!|8ZArr;HrifQP;bj-jRn29qn3uobM`~%Lx zKjM}6C(K4C=Aa96F%RcrJ{DjhUWHeq8;h_Q=b;Dxj3ro#WjG%f;5E1q|ALF~uXru~ z4X?v;tUw7X(ThH$ScTQ-M;Wij8!&)DR8U0?|BgC_(7-U(;Efo;TD%GWfs65HvBg(!T(?#{ugh@JMd1t3zy>ExD4;Xdb}6!!~3xTAHe1KAU=c-;|hEPSK_0% z3LnEpd>mKf6WD}L;#2rEuEA$;Ek29Q_#Cdo=W#v0fE(~d+=wsXCftlK<15&Lui_Sb z4O{VbY{NHjE53=_@GabqZ{rSp2ix&o?7;VMC%%un@B`e9AL1VT2s`m(?7~lQFMf)j z;peyyzrg+YC3fRicmTh~9{dIm;$M7fY#h>vw{(>j)S3HTo z;VJwbPvaSC)ybITfMjfPU@|T_C^pCzFz6l4FzS zB*!JsO^#2Vmzi=6*dt1~27@a4QQTh{MTjSha2xn(UiUJ|q4&S$550-@5IkHnI z>WX8ZK2&lXp=c*7S(bVE{76$uFU~5f7pqLr)s)_+-ZmwWlX*S!r0!03M(#ewOhoWht7n|Kev2xqH;|BrJkyB<&wi_YTp~2 z<<$x~-Du8{22YN*FOA0xCnjjlUc~Q(la8~wsI0gh#}-_tfV&DV?%2S&QN`R&SX#_j zS@TuRXImS}mqcE)K1=c>a@vAQKBqBAOE}xj1A`u=(*`GnBK{W~6@GRah-1zbe8ph% zk_8n;yyPvx*DB%)!PjcSxq+`NPie9fA$ft-kurw)!l3~f+vG|V^fVCB7Y+q-d)gF5XCT_<+vKi!TWhUb$b#BPM=hken0i$RuLnNg zc^4ZaC1dc+mG*7P%%zLlDRV7#YgnCx`7!oZxUrmd4UM^`^%~8r!i3@OKBfG#W&2qk>#fDsbSZbCZKKZ`on}Qj~Pcmv~#4|jU<{J_@axqIG zS2W)g$dxQ)?PHXzyFA_m&WM zMPU2DPa$T*P((TxWc0=&x}pkKQ6LjsH&({xj52iAj1dJHUj*2NY%8>0HZ2UBFEniq zVbb+1A9R#9SqQM4$*gU~U{Z)0Ef_J+Cc%v=Zf4+Uy30d%AIL3-$cfn|%63xSWWDi8 zDi*lr@8qxG{e4GkpLg>e`2u(DwT8YH@0C% z-$S>8Nw2`YbzWeZX8t8u-y}n4F??TFcp>6<%zZTZMvG!fk=)C*+I^moD0t`F$%wgN zogc#|Ci?2K+5`(dfP$B>{hZ`ihJR3?GZKCPbhkM60QBACDCRyMfKDX;0ca6@fms)O z6qk~Dkm(UnezlA2&ZHyTRMZz7j&c-TNqr z8u_!2hGe6e^9>I>=zYB-Tiw}Wp-*vGr||S)0PlMm-ZwHxN-*tnR?W6{h6P)nv&%0J zurdYD)}@qpRw~&vOq%m_&ErR%hG~5uLGN@K0r?u$^6rVbj9_QwiorhNAUV$YxM5lu z2%PmPNp>18JbsIhyEy0IeuOyzH*wxGQI+jI6m^J^^H~G$^90%vxqlTxM}?l4gmJ$s zU=Dh;vi(H>&UPNHPSN-FpE-!tB9@nao~Ayn=FGt3*D1+~QZ}!WTS;0sa5f!?%TA?( zQOfCu#53tB2lfFK{c6#1j$?YiP0GZ>tx(0e4beJfCZwQ4_ z(nno^%d{e=y~^iy_P?BHDyG=HRln{13c+>HbCoM=8i%vvT44|9o#84$za1*cKzFlF zY{_?L(A%};)mk=OB`gEJyJ|ho>vsSg&9F~QRDG+VlXdls1!mcYZmr^nId8?F` zylfg^4m{m6Rw-=_*VgH3BwjCB0y2ETVnv&?RAPtxA2Yr=KV?jF;>Wq>xGv1~+!U*W zIW28EzjZd60rOK)&NI0&$AnqDJ_%(E4cE9J^?Z$L zH=&M~QUCN%OL|e<6R4IlZ#MbPq8Rzh_U@_QA~MnmX0H_)nVkmer0mht@GV!DmIuv{ zyG$-^HeU4fjGgfju__ZUGZB_b);B_aq!5cwCg}m zoUP=t6co!5)h^w7Yx=&P&W9swyI%C&EIQ)8=R`+9XWt;QQ#PFo!*ldK*M*LoLPtDv zG^i7{PS6b3b0g@a=>ddd6W;n90^E@}& z@rdU}TW;=5q;+Ct-m*5I&8h7OeTGGu&W@;PMkl@x%b5W$a&q^}NLgo9T9%E16eWFBy+)o!X2nxMfvsXDkKN2eYaIwX_Ws$gdT0~xeg3hq?w8B0cBKhd4 zkajB+_EBY@1ZiQ*sg&lLPU)@!P44ned~_nL2TNgg%ApaCfbK&Z8?9EAAMw)XLUy*I_`;mHUUq`W1l>jg76vwVIeTKsI zgH=ZRSH>F+7k8?R1RWnq-YE%J&#y9ls1c|#J1)=T3$HT!9HhS-^}Etv4xfEUe%Pmq zq;xw&+B0vrtaH>EFKIlTJhr7DOy04~h}>TfoV zoHM1|S89|_?#OtP0Ne>+w_+RsU^b?1Q*tvixx?v4V&v#=HYRTy1K?;JBEl1RpZ-=%Q;m94LNq z7+`j9b+9(2+;_onl05T_)RQ{&NBJ9rtkkmsk>mc1Bg#i82V|TXBtx27<_yE3zZZ&s zQMEyTq1MtTV$2_FEz>hDwU$q10rf|% zW&gB-B2V`DqP7>G-`%}@;i8`NmMmL7ciz&jB@26&owo$zc}pqM%Q|NHz3jey)LW#N zosrcM^-}6({BJg9XT5Aj#wWR*SwQ`fT>I;{NZp67!VV_$1kfU?eoxx1=_K)%a8T7ruMNrZ3)O7| zah!{&Ce#*)6)1$5({ia_9^@jDKLlPRD^Lh68~U6YJq0w&P&I0{huhR*XXMoCaR3g% zSejX7_dr@JU8F9@D;(>RNlR!ma_aH~0EaR(wR&w0D`#A#HzWCQa<9i&U`_Y-Sji5P zXw0RB4aSiRwaCabJKV$Sm}>!8fkJ4JvGWKIE!z;~Woo`QXXdp2(*YbAz_j-lZE&8G zBIS*>o1-MQb>oSr-+W19&vK$dv+h$_a6C|?Y_YNWXocoo*5V*mpb%o@Z9j=4+P9AO z8|E<_(>fw=4Rc0LIgbZ$tisfH*jg&xjGTGJ^8h?YBW0Um7F5%tT5Jw;oWb&3+q8(9 zQ#-!~$b90bmjG}IH#yFa(J&z;u%WD=S4-_h6 zKA)qh>o*oy(=9dj{ymYS$*n;Pj0cL8EH)#20mpK$CvhMv zPzW{h^d@UmcGI*qv}~iqnf1Ov?DUA)Q&(ezUt_i4ntpz<-f)t}b1s5ekUY?8eZI&- zY)5a1wV9JST6^0}Ps!u$ONZ()W}=;v+6|)Q4xH2ySdI@Z}sgKKW_-*~VtCQ#dZ&!smsu0)I` z=?#077WP>nR-g!C?0q)VL5xlp zntq-$b9(2?0i5X~i|ZrVMGtvj(yHj@(IfBiSsdLMw`fO>Eh6W%V|ve=6-LfXaV%gS zDDtFZE&A*{T)HU8fvi9w)SR7BI&u1kJXGP%(t>U4^F1ewt4)+zLyoNu|B)k011JlO z2MX10WaNJ($BHdx+FG~&dj z>U@Q3+}he;9J$r8XMu-h?(1@#=KDu%r*9!cZPT7~q~*Mdp>dbLEtW5*H`0##s~Mua z&DtVX&K#FcNV^%9GC#Gzc%V>giPYpGhLt;TH~*IE2=B6=24Jy5%=M_ui)RH2l{E)> zGJx|mGB+ErAgn+UyqP)AaVmfw2QSCwx3%)-%;f3D%Rg&;X;{^+aUeFEU1BhObLEyw z8f#NaHL6^#$iq!4ud;=12{~5B%N%67VJn?$TaeB7pqyQ}DFDvbm=QB>3!VpxwCWi- zGp}O+Tp-bf=_OCM`7*|Kqh4cBm9I_!fevu~_E8}YoA{iBAt_CcOQPgTqo1u~Q-z>y7 z?(0QX?60#h+q<@l&Fq#-T+81Ae|n>+*__dZ)^#fsn(fV@HN4oou4M3(uZR}0a>i}C zAF@(og^t2GZlN9#sl{FgNxHe<2`1*3e>OSy`jmFGA@VrA6GxYUlMD*Y>;8 zB4W<2GM$dR!9f&eQ5GBz6ltTedGLTnQzyc%%NP!31qvbOyb(^_q*~xSP^9i-qhZa(UXrd#>uIIenF2M9 znX?~4w;ujoVTQZ_ah&F#b8K!|w^5CcAq$cR3e{~+yQZ1mkd0`Z`&cNWC}jk*VPhI^ zKdoX$p675L<}_ymxaocyE9n{!HFNF6pJDTuvE8*dGR*yLe=x$Z_W!QvT874bUbR-2 zu~Gj`4x;HE#xYfE=&^U}e=v0IugiktfkGoe&InI;buYHjj59C}Vg(8zMrJGj$q>1x zF$=Ki{un!be=|cg&(b&`+keP!2_mLruZ^0` z`aP$Bt>qxLqI|oZ1IOiHYq&FW#^{#;c!!HyPgj?R8mVz;*cx{1tmT~&JM1peg5`lC zV_0k#>s<;}`(osYTY*AIIrCuJnZMK^wc)fO)@$BvkUTe%R`HsC+s5`>F0)YEv1b3oktlgEMkv1IEFR%YSfB3f*(?7arf*k`wj7g;O4PvQlggSOUMV)ql@A41JBaNxNp z)C+TVBI&$ugTk$po$H8A_k~!Ue!!r#`Cd+JPJg*V^}hYcE6@rQsm&PT2Q{MO@U=C> z$gPJDaYW-Iwl%z*QJZ#5KCJOvSG5*o^Q|Z6Elif1S6G-ig=`)%vakCQiKrGy>VutD z0ePTEtBI{qu9RqUc_MxI^5m^R5v%VdtUwXmGjqlna*ql4E7wu&qYV14a5>?9KZr3(fNQYvDBCQew~U(;QX0 z=H`X60!2y~Td`cDu+r=c!L1?XoI{bnna>y`?Sk9_Yin=2)<$bXEymvKpVer_32Hm6 zfnvRGGsj$5SygSdHgiT}nnQd}WA)V2RcWKQzb&`XBmMk3gJ}H9Tg`9!I% z>$uTIGfw+Ch!rT*I&(%&vPb(;4q~97&DU#Wb##-B8+2>M3tgz3Ir9iQH@`WE*N#%g z+Ub`$>VBO8e#J%#6s+0WjLjmpaI|({0e;oOq7T&bk~H5cV_3IvEb-Be*05qd z;A<|{ey+!@4qmJE7%SV?U92{gErz#^*keF=;+#J+;70 zKj3KlbtZeajns~E<-AvEFY|{StsPi7J&o*8?%`PU;V?c(S%E?$R%8tMk%JT{So5tu zQnH;E-t-QuFjk<|lKt30YD>u?{dt##H_IwZ^K}?|%lyPf3KWbtmqPWJv#O%?@VyGn z+IJ3J%Wb-y#AdNS4WR}v=r&!}$R5YfLa6#9c^qjJrH$<7{9NIBpWw3~d7#i!jm&HB za}d?{{WydbXcg}l4xajWU~70ey^&_N_j9~_X98_)v%k#4r0*WKRMOaddbfp2A62xV znlEAGZSgA?Yi_AgY85Xs`+FdSH*cgx#928#aW;Tob432-#Q5hXtt#h%BK_v9oElgN zV9%)X(=I;}XXfOLJ{$C#yqpETG0AJB`4%}VXB>MafCr1`BBxS#-ikb}*y-YLd2ZB4 zK6(1hm;KC~87Y0y>miBuEVa^EIdg9J~NPM%%?;9);cZEu;eH|*~v zzWZ%et2yvMk+v6`r#+(ZLLYABIL+5`Y%cl-g{r>-TdKGEuFIL9(dqr8E?TwLV%cJL z5r33u+5Omd-+IhhC6V>~V-7O=0Iu~Yd7x1H$*Eg9-}_UB8T=S8M;b*7%6a3G)z)4M zw{0bjApTh)dfySY3dsXS>Nd9OeOw}%j}-I7tw0f^*bMD24pKYrtjGHQ6B=*0)>&UY zJeXEm?Ef*pN`F;o(T_K`R=1IM^Q1x-?soZ?p<6@Dd56*3`)?Ma{n=>?wCOizZ0+>-AXcCV zVr;+rK#mx2sL6qgV8u!|j$`FIxojFSXWdAiCJ*9>=EL?J!U_~BVFd4BiKjnl9~$B~ zcpF_Qw7TW?6Bz@KcJZq9!BUH>h1iOAl7qLXG}2iq z4^ zx}V4V=R0`M&X|9qgW9eUB3836a54Aum_OOY+s|YE3mv>?Ys^2%#eLSs{1*k0pN%p9 zp{Tg-))gHzDy&kHF*%S*)qm%h@7qw$!YigU9G*`#q!( + txt: &'input str, + ) -> Result, Vec, ast::PtxError>>>; + + // Returned AST might be malformed. Some users, like logger, want to look at + // malformed AST to record information - list of kernels or such + fn parse_unchecked<'input>( + txt: &'input str, + ) -> ( + ast::Module<'input>, + Vec, ast::PtxError>>, + ); +} + +impl ModuleParserExt for ModuleParser { + fn parse_checked<'input>( + txt: &'input str, + ) -> Result, Vec, ast::PtxError>>> { + let mut errors = Vec::new(); + let maybe_ast = ptx::ModuleParser::new().parse(&mut errors, txt); + match (&*errors, maybe_ast) { + (&[], Ok(ast)) => Ok(ast), + (_, Err(unrecoverable)) => { + errors.push(unrecoverable); + Err(errors) + } + (_, Ok(_)) => Err(errors), + } + } + + fn parse_unchecked<'input>( + txt: &'input str, + ) -> ( + ast::Module<'input>, + Vec, ast::PtxError>>, + ) { + let mut errors = Vec::new(); + let maybe_ast = ptx::ModuleParser::new().parse(&mut errors, txt); + let ast = match maybe_ast { + Ok(ast) => ast, + Err(unrecoverable_err) => { + errors.push(unrecoverable_err); + ast::Module { + version: (0, 0), + directives: Vec::new(), + } + } + }; + (ast, errors) + } +} + +pub struct DisplayParseError<'a, Loc, Tok, Err>(pub &'a str, pub &'a ParseError); + +impl<'a, Loc, Tok, Err> fmt::Display for DisplayParseError<'a, Loc, Tok, Err> +where + Loc: fmt::Display + Into + Copy, + Tok: fmt::Display, + Err: fmt::Display, +{ + fn fmt(&self, f: &mut fmt::Formatter) -> fmt::Result { + match self.1 { + ParseError::UnrecognizedToken { + token: (start, token, end), + .. + } => { + let full_instruction = + unsafe { self.0.get_unchecked((*start).into()..(*end).into()) }; + write!( + f, + "`{}` unrecognized token `{}` found at {}:{}", + full_instruction, token, start, end + ) + } + _ => self.fmt(f), + } + } +} + pub(crate) fn without_none(x: Vec>) -> Vec { x.into_iter().filter_map(|x| x).collect() } @@ -53,3 +135,33 @@ pub(crate) fn vector_index<'input>( }), } } + +#[cfg(test)] +mod tests { + use crate::{DisplayParseError, ModuleParser, ModuleParserExt}; + + #[test] + fn error_report_unknown_instructions() { + let module = r#" + .version 6.5 + .target sm_30 + .address_size 64 + + .visible .entry add( + .param .u64 input, + ) + { + .reg .u64 x; + does_not_exist.u64 x, x; + ret; + }"#; + let errors = match ModuleParser::parse_checked(module) { + Err(e) => e, + Ok(_) => panic!(), + }; + assert_eq!(errors.len(), 1); + let reporter = DisplayParseError(module, &errors[0]); + let build_log_string = format!("{}", reporter); + assert!(build_log_string.contains("does_not_exist")); + } +} diff --git a/ptx/src/ptx.lalrpop b/ptx/src/ptx.lalrpop index 181a727..f6fd9cd 100644 --- a/ptx/src/ptx.lalrpop +++ b/ptx/src/ptx.lalrpop @@ -526,8 +526,15 @@ Statement: Option>> = { ";" => Some(ast::Statement::Instruction(p, i)), PragmaStatement => None, "{" "}" => Some(ast::Statement::Block(without_none(s))), - ! ";" => { - let (err, _) = (<>); + @L ! ";" @R => { + let (start, mut err, _, end) = (<>); + // TODO: report this error more generally, perhaps in user error? + err.error = match err.error { + ParseError::UnrecognizedToken { token: (_, token, _), expected } => { + ParseError::UnrecognizedToken { token: (start, token, end), expected } + } + e => e + }; errors.push(err.error); None } diff --git a/zluda_dump/src/trace.rs b/zluda_dump/src/trace.rs index 3bdf807..eac6bbd 100644 --- a/zluda_dump/src/trace.rs +++ b/zluda_dump/src/trace.rs @@ -1,4 +1,5 @@ use ptx::{ast::PtxError, Token}; +use ptx::{DisplayParseError, ModuleParserExt}; use crate::{cuda::CUmodule, dark_api, log, Settings}; use std::{ @@ -170,24 +171,18 @@ impl StateTracker { submodule_index: Option, module_text: &str, ) { - let mut errors = Vec::new(); - let ast = ptx::ModuleParser::new().parse(&mut errors, module_text); - let ast = match (&*errors, ast) { - (&[], Ok(ast)) => ast, - (err_vec, res) => { - fn_logger.log(log::LogEntry::ModuleParsingError( - DumpWriter::get_file_name(module_index, version, submodule_index, "log"), - )); - fn_logger.log_io_error(self.writer.save_module_error_log( - module_index, - version, - submodule_index, - err_vec, - res.err(), - )); - return; - } - }; + let (ast, errors) = ptx::ModuleParser::parse_unchecked(module_text); + if !errors.is_empty() { + fn_logger.log(log::LogEntry::ModuleParsingError( + DumpWriter::get_file_name(module_index, version, submodule_index, "log"), + )); + fn_logger.log_io_error(self.writer.save_module_error_log( + module_index, + version, + submodule_index, + &*errors, + )); + } } pub(crate) fn module_exists(&self, hmod: CUmodule) -> bool { @@ -238,8 +233,7 @@ impl DumpWriter { module_index: usize, version: Option, submodule_index: Option, - recoverable: &[ptx::ParseError, PtxError>], - unrecoverable: Option, PtxError>>, + errors: &[ptx::ParseError, PtxError>], ) -> io::Result<()> { let mut log_file = match &self.dump_dir { None => return Ok(()), @@ -252,8 +246,9 @@ impl DumpWriter { "log", )); let mut file = File::create(log_file)?; - for error in unrecoverable.iter().chain(recoverable.iter()) { - writeln!(file, "{}", error)?; + for error in errors { + let pretty_print_error = DisplayParseError("", error); + writeln!(file, "{}", pretty_print_error)?; } Ok(()) }