Better reporting of unrecognized tokens

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(())
 }

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

ptx/src/lib.rs

@@ -28,6 +28,8 @@ pub mod ast;
 mod test;
 mod translate;
 
+use std::fmt;
+
 pub use crate::ptx::ModuleParser;
 pub use lalrpop_util::lexer::Token;
 pub use lalrpop_util::ParseError;
@@ -36,6 +38,86 @@ pub use translate::to_spirv_module;
 pub use translate::KernelInfo;
 pub use translate::TranslateError;
 
+pub trait ModuleParserExt {
+    fn parse_checked<'input>(
+        txt: &'input str,
+    ) -> Result<ast::Module<'input>, Vec<ParseError<usize, Token<'input>, 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<ParseError<usize, Token<'input>, ast::PtxError>>,
+    );
+}
+
+impl ModuleParserExt for ModuleParser {
+    fn parse_checked<'input>(
+        txt: &'input str,
+    ) -> Result<ast::Module<'input>, Vec<ParseError<usize, Token<'input>, 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<ParseError<usize, Token<'input>, 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<Loc, Tok, Err>);
+
+impl<'a, Loc, Tok, Err> fmt::Display for DisplayParseError<'a, Loc, Tok, Err>
+where
+    Loc: fmt::Display + Into<usize> + 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<T>(x: Vec<Option<T>>) -> Vec<T> {
     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"));
+    }
+}

ptx/src/ptx.lalrpop

@@ -526,8 +526,15 @@ Statement: Option<ast::Statement<ast::ParsedArgParams<'input>>> = {
     <p:PredAt?> <i:Instruction> ";" => Some(ast::Statement::Instruction(p, i)),
     PragmaStatement => None,
     "{" <s:Statement*> "}" => 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
     }

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<usize>,
         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<usize>,
         submodule_index: Option<usize>,
-        recoverable: &[ptx::ParseError<usize, Token<'input>, PtxError>],
-        unrecoverable: Option<ptx::ParseError<usize, Token<'input>, PtxError>>,
+        errors: &[ptx::ParseError<usize, Token<'input>, 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(())
     }