Better reporting of unrecognized tokens

6 files changed, 461 insertions, 27 deletions

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 0000000..1d7df96
--- /dev/null
+++ b/hip_runtime-sys/lib/amdhip64.lib
diff --git a/ptx/src/lib.rs b/ptx/src/lib.rs
index 4ade4e8..db9fc23 100644
--- a/ptx/src/lib.rs
+++ b/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"));
+    }
+}
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<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
     }
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<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(())

     }